Как мне лучше всего инициализировать массив локальной памяти до 0?

Простой цикл должен быть «лучшим» подходом (но см. окончательный комментарий ниже). В качестве примера используем следующее ядро:

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);

Здесь у нас есть три различных способа обнуления большого массива локальной памяти, а также некоторый код, чтобы убедить компилятор в том, что вся последовательность инициализации и локальный массив не должны быть оптимизированы.

Глядя на PTX, выдаваемые для целей Compute 2.1 с помощью компилятора выпуска CUDA 6, обе версии 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-битных инструкций сохранения. Уловка int4 в версии 3 давала такой же код, как и простой цикл, что немного удивительно. Однако версия 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-битную запись (комментарии указывают, что простая инициализация списка также дает этот тип цикла копирования). Последний будет намного медленнее, чем первый. Помимо разницы в размерах хранилищ, развернутый поток операций записи полностью независим и может выполняться в любом порядке, что обеспечит заполнение конвейера инструкций и должно привести к более высокой пропускной способности инструкций. Я не верю, что в развернутом случае удастся превзойти компилятор, и простой цикл, похоже, даст тот же код, что и простая попытка векторизации. Если вы действительно заинтересованы, я думаю, вы могли бы попробовать встроенный PTX для создания более широких хранилищ. Я не знаю, будет ли при этом какое-либо преимущество в производительности.