在CUDA中，有什么高效的方法可以交换两个寄存器变量？

Question

在CUDA中，有什么高效的方法可以交换两个寄存器变量？

7

我开始编写一些CUDA代码，我想在内核中为两个变量做相当于std::swap()的操作；它们位于寄存器文件中（无溢出，不在某个缓冲区中等）。假设我有以下设备代码：

__device__ foo(/* some args here */) {

    /* etc. */

    int x = /* value v1 */;
    int y = /* value v2 */;

    /* etc. */

    swap(x,y);

    /* etc. */
}

现在，我可以直接写：

template <typename T> void swap ( T& a, T& b )
{
  T c(a); a=b; b=c;
}

但我想知道 - 是否有一些内置的CUDA功能可以实现这个功能呢？

注：

是的，我希望它可以在所有线程上运行。
不要担心我是否有足够的寄存器。假设我有足够的寄存器。

- einpoklum

1

我对你的问题不是很清楚。CUDA使用单指令多数据（SIMD）处理方案，这意味着您在内核函数中编写的指令将由计算网格中的所有线程执行。根据网格大小，您可能有足够的寄存器空间或没有。 - Vitality

2个回答

3

据我了解，这些都是完全无关的。 x和y并不是“真实”的对象：它们只存在于C++标准描述的抽象机器中。特别地，它们不对应于寄存器。

你可能想象编译器在创建程序时会将它们分配给寄存器，但事实上它们并不是这样工作的。存储在寄存器中的东西可以被重排、复制、变成其他东西，甚至完全消失。

特别地，无条件交换存储在寄存器中的两个变量通常不会生成任何代码 - 它的唯一效果是让编译器调整其内部表格，在那个时间点上存储了哪些对象在哪些寄存器中。

（即使是有条件的交换，让编译器自己做还是比自己写更好）。

- user1084944

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Vitality · Accepted Answer

我已经考虑了以下测试程序。

template <typename T> __device__ void inline swap_test_device1(T& a, T& b)
{
    T c(a); a=b; b=c;
}

template <typename T> __device__ void inline swap_test_device2(T a, T b)
{
    T c(a); a=b; b=c;
}

__global__ void swap_test_global(const int* __restrict__ input1, const int* __restrict__ input2, int* output1, int* output2) {

    int tx = threadIdx.x + blockIdx.x * blockDim.x;

    int x = input1[tx]*input1[tx];
    int y = input2[tx]*input2[tx];

    //swap_test_device2(x,y);
    swap_test_device1(x,y);

    output1[tx] = x;
    output2[tx] = y;

}

我已经对其进行了反汇编。当使用swap_test_device1和swap_test_device2时，结果相同。共同的反汇编代码如下：

MOV R1, c[0x1][0x100];
S2R R0, SR_CTAID.X;
S2R R2, SR_TID.X;
MOV32I R9, 0x4;  
IMAD R3, R0, c[0x0][0x8], R2;
IMAD R6.CC, R3, R9, c[0x0][0x28];
IMAD.HI.X R7, R3, R9, c[0x0][0x2c];
IMAD R10.CC, R3, R9, c[0x0][0x20];
LD.E R2, [R6];                         loads input1[tx] and stores it in R2
IMAD.HI.X R11, R3, R9, c[0x0][0x24];
IMAD R4.CC, R3, R9, c[0x0][0x30];
LD.E R0, [R10];                        loads input2[tx] and stores it in R0
IMAD.HI.X R5, R3, R9, c[0x0][0x34];
IMAD R8.CC, R3, R9, c[0x0][0x38];
IMAD.HI.X R9, R3, R9, c[0x0][0x3c];
IMUL R2, R2, R2;                       R2 = R2 * R2
ST.E [R4], R2;                         stores input1[tx]*input1[tx] in global memory
IMUL R0, R0, R0;                       R0 = R0 * R0
ST.E [R8], R0;                         stores input2[tx]*input2[tx] in global memory
EXIT ;

在反汇编的代码中似乎没有明确的交换。换句话说，编译器对于这个简单的例子，能够直接优化代码，将 x 和 y 写入正确的全局内存位置。

编辑

现在我考虑了以下更复杂的测试用例

__global__ void swap_test_global(const char* __restrict__ input1, const char* __restrict__ input2, char* output1, char* output2) {

    int tx = threadIdx.x + blockIdx.x * blockDim.x;

    char x = input1[tx];
    char y = input2[tx];

    //swap_test_device2(x,y);
    swap_test_device1(x,y);

    output1[tx] = (x >> 3) & y;
    output2[tx] = (y >> 5) & x;

 }

使用与上述__device__函数相同的函数。反汇编代码为：

MOV R1, c[0x1][0x100];              
S2R R0, SR_CTAID.X;                 
S2R R2, SR_TID.X;           
IMAD R0, R0, c[0x0][0x8], R2;       R0 = threadIdx.x + blockIdx.x * blockDim.x
BFE R7, R0, 0x11f;
IADD R8.CC, R0, c[0x0][0x28];
IADD.X R9, R7, c[0x0][0x2c];
IADD R10.CC, R0, c[0x0][0x20];
LD.E.S8 R4, [R8];                   R4 = x = input1[tx]
IADD.X R11, R7, c[0x0][0x24];
IADD R2.CC, R0, c[0x0][0x30];
LD.E.S8 R5, [R10];                  R5 = y = input2[tx]
IADD.X R3, R7, c[0x0][0x34];
IADD R12.CC, R0, c[0x0][0x38];
IADD.X R13, R7, c[0x0][0x3c];
SHR.U32 R0, R4, 0x3;                R0 = x >> 3
SHR.U32 R6, R5, 0x5;                R6 = y >> 5
LOP.AND R5, R0, R5;                 R5 = (x >> 3) & y
LOP.AND R0, R6, R4;                 R0 = (y >> 5) & x
ST.E.U8 [R2], R5;                   global memory store
ST.E.U8 [R12], R0;                  global memory store
EXIT ;

正如你所看到的，仍然没有明显的寄存器交换。