如何最好地将本地内存数组初始化为 0？

Question

（关于设备（全局）内存阵列也有类似的问题，例如我自己的问题。）

假设我有一个 CUDA 内核代码，如下所示：

my_arr[MyCompileTimeConstant];

/* ... */

for(unsigned i = 0; i < foo(); i++) {
   my_arr[bar(i)] += some_value;
}

现在，我想my_arr在开始添加其条目之前初始化为全零。我能做得比简单的循环更好吗

for(unsigned i = 0; i < MyCompileTimeConstant; i++) {
   my_arr[i] = 0;
}

？

注意：我特意在编译时设置了循环范围和数组大小常量。如果它们在运行时通过，问题会略有不同。当然，它可能不会像改变 CPU 上运行的代码那样改变 CUDA 的答案

Answer 1

一个简单的循环应该是“最佳”方法（但请参阅下面的最终评论）。使用以下内核作为示例：

template<int version>
__global__
void tkernel(int *A, int *B, int *C, int n)
{
    int biglocal[100];

    switch(version) {
        case 1:
            for(int i=0; i<100; i++) {
                biglocal[i] = 0;
            };

            break;

        case 2:
            memset(&biglocal[0], 0, 100*sizeof(int));
            break;


        case 3:
            const int4 zero = {0, 0, 0, 0};
            int4 *p = reinterpret_cast<int4*>(&biglocal[0]);
#pragma unroll
            for(int i=0; i<100/4; i++) {
                p[i] = zero;
            }

            break;
    }

    if (n>0) {
        for(int i=0; i<100; i++) {
            biglocal[A[threadIdx.x*i]] += B[threadIdx.x*i];
        }
        C[threadIdx.x] = biglocal[n];
    }
}

template __global__ void tkernel<1>(int *, int *, int *, int);
template __global__ void tkernel<2>(int *, int *, int *, int);
template __global__ void tkernel<3>(int *, int *, int *, int);

这里我们有三种不同的方法来将大型本地内存数组清零，再加上一些代码来说服编译器整个初始化序列和本地数组不应该被优化掉。

查看使用 CUDA 6 版本编译器为计算 2.1 目标发出的 PTX，版本 1 和 3 均如下所示：

.local .align 4 .b8     __local_depot0[400];
.reg .b64   %SP;
.reg .b64   %SPL;
.reg .pred  %p<3>;
.reg .s32   %r<67>;
.reg .s64   %rd<73>;


mov.u64     %SPL, __local_depot0;
ld.param.u64    %rd4, [_Z7tkernelILi1EEvPiS0_S0_i_param_0];
ld.param.u64    %rd5, [_Z7tkernelILi1EEvPiS0_S0_i_param_1];
ld.param.u64    %rd6, [_Z7tkernelILi1EEvPiS0_S0_i_param_2];
ld.param.u32    %r21, [_Z7tkernelILi1EEvPiS0_S0_i_param_3];
add.u64     %rd7, %SPL, 0;
mov.u32     %r66, 0;
st.local.u32    [%rd7], %r66;
st.local.u32    [%rd7+4], %r66;
st.local.u32    [%rd7+8], %r66;
st.local.u32    [%rd7+12], %r66;
st.local.u32    [%rd7+16], %r66;
st.local.u32    [%rd7+20], %r66; 

    // etc

IE。编译器展开循环并发出一串 32 位存储指令。版本 3 中的技巧int4产生了与简单循环相同的代码，这有点令人惊讶。然而，版本 2 得到了这个：

.local .align 4 .b8     __local_depot1[400];
.reg .b64   %SP;
.reg .b64   %SPL;
.reg .pred  %p<4>;
.reg .s16   %rs<2>;
.reg .s32   %r<66>;
.reg .s64   %rd<79>;


mov.u64     %SPL, __local_depot1;
ld.param.u64    %rd7, [_Z7tkernelILi2EEvPiS0_S0_i_param_0];
ld.param.u64    %rd8, [_Z7tkernelILi2EEvPiS0_S0_i_param_1];
ld.param.u64    %rd9, [_Z7tkernelILi2EEvPiS0_S0_i_param_2];
ld.param.u32    %r21, [_Z7tkernelILi2EEvPiS0_S0_i_param_3];
add.u64     %rd11, %SPL, 0;
mov.u64     %rd78, 0;

BB1_1:
add.s64     %rd12, %rd11, %rd78;
mov.u16     %rs1, 0;
st.local.u8     [%rd12], %rs1;
add.s64     %rd78, %rd78, 1;
setp.lt.u64 %p1, %rd78, 400;
@%p1 bra    BB1_1;

IE。执行 8 位写入的循环（注释表明简单的列表初始化也将产生这种类型的复制循环）。后者会比前者慢很多。除了存储的大小差异之外，展开的写入流是完全独立的，并且可以以任何顺序发出，以保持指令管道满，并且应该导致更高的指令吞吐量。我不相信在展开的情况下能够击败编译器，并且一个简单的循环看起来会产生与简单的矢量化尝试相同的代码。如果您真的很热衷，我想您可以尝试内联 PTX 来生成更广泛的存储。我不知道这样做是否会有任何性能优势。