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

Question

（关于设备（全局）内存阵列也有类似的问题，例如my own 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

即。编译器展开循环并发出一串 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;

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