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

Question

Мне не удается получить доступ к данным в моей постоянной памяти, и я не знаю почему.Вот фрагмент моего кода:

#define N 10
__constant__ int constBuf_d[N];

__global__ void foo( int *results, int *constBuf )
{
    int tdx = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tdx;

    if( idx < N )
    {
         results[idx] = constBuf[idx];
    }
}

// main routine that executes on the host
int main(int argc, char* argv[])
{
    int *results_h = new int[N];
    int *results_d = NULL;

    cudaMalloc((void **)&results_d, N*sizeof(int));

    int arr[10] = { 16, 2, 77, 40, 12, 3, 5, 3, 6, 6 };

    int *cpnt;
    cudaError_t err = cudaGetSymbolAddress((void **)&cpnt, "constBuf_d");

    if( err )
        cout << "error!";

    cudaMemcpyToSymbol((void**)&cpnt, arr, N*sizeof(int), 0, cudaMemcpyHostToDevice);

    foo <<< 1, 256 >>> ( results_d, cpnt );

    cudaMemcpy(results_h, results_d, N*sizeof(int), cudaMemcpyDeviceToHost);

    for( int i=0; i < N; ++i )
        printf("%i ", results_h[i] );
}

По какой-то причине я получаю только «0» в results_h.Я использую CUDA 4.0 с картой с возможностью 1.1.

Есть идеи?Спасибо!

Answer 1

Если вы добавите правильную проверку ошибок в свой код, вы обнаружите, что cudaMemcpyToSymbol завершается с ошибкой с недопустимым символом устройства. Вам нужно либо передать символ по имени, либо использовать вместо него cudaMemcpy. Итак, это:

cudaGetSymbolAddress((void **)&cpnt, "constBuf_d");
cudaMemcpy(cpnt, arr, N*sizeof(int), cudaMemcpyHostToDevice); 

или

cudaMemcpyToSymbol("constBuf_d", arr, N*sizeof(int), 0, cudaMemcpyHostToDevice);

или

cudaMemcpyToSymbol(constBuf_d, arr, N*sizeof(int), 0, cudaMemcpyHostToDevice);

будет работать. Сказав это, передача адреса постоянной памяти в качестве аргумента ядру является неправильным способом использования постоянной памяти - это лишает компилятор возможности генерировать инструкции для доступа к памяти через кэш постоянной памяти. Сравните вычислительные возможности 1.2 PTX, сгенерированные для вашего ядра:

    .entry _Z3fooPiS_ (
        .param .u32 __cudaparm__Z3fooPiS__results,
        .param .u32 __cudaparm__Z3fooPiS__constBuf)
    {
    .reg .u16 %rh<4>;
    .reg .u32 %r<12>;
    .reg .pred %p<3>;
    .loc    16  7   0
$LDWbegin__Z3fooPiS_:
    mov.u16     %rh1, %ctaid.x;
    mov.u16     %rh2, %ntid.x;
    mul.wide.u16    %r1, %rh1, %rh2;
    cvt.s32.u16     %r2, %tid.x;
    add.u32     %r3, %r2, %r1;
    mov.u32     %r4, 9;
    setp.gt.s32     %p1, %r3, %r4;
    @%p1 bra    $Lt_0_1026;
    .loc    16  14  0
    mul.lo.u32  %r5, %r3, 4;
    ld.param.u32    %r6, [__cudaparm__Z3fooPiS__constBuf];
    add.u32     %r7, %r6, %r5;
    ld.global.s32   %r8, [%r7+0];
    ld.param.u32    %r9, [__cudaparm__Z3fooPiS__results];
    add.u32     %r10, %r9, %r5;
    st.global.s32   [%r10+0], %r8;
$Lt_0_1026:
    .loc    16  16  0
    exit;
$LDWend__Z3fooPiS_:
    } // _Z3fooPiS_

с этим ядром:

__global__ void foo2( int *results )
{
    int tdx = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tdx;

    if( idx < N )
    {
         results[idx] = constBuf_d[idx];
    }
}

, который производит

    .entry _Z4foo2Pi (
        .param .u32 __cudaparm__Z4foo2Pi_results)
    {
    .reg .u16 %rh<4>;
    .reg .u32 %r<12>;
    .reg .pred %p<3>;
    .loc    16  18  0
$LDWbegin__Z4foo2Pi:
    mov.u16     %rh1, %ctaid.x;
    mov.u16     %rh2, %ntid.x;
    mul.wide.u16    %r1, %rh1, %rh2;
    cvt.s32.u16     %r2, %tid.x;
    add.u32     %r3, %r2, %r1;
    mov.u32     %r4, 9;
    setp.gt.s32     %p1, %r3, %r4;
    @%p1 bra    $Lt_1_1026;
    .loc    16  25  0
    mul.lo.u32  %r5, %r3, 4;
    mov.u32     %r6, constBuf_d;
    add.u32     %r7, %r5, %r6;
    ld.const.s32    %r8, [%r7+0];
    ld.param.u32    %r9, [__cudaparm__Z4foo2Pi_results];
    add.u32     %r10, %r9, %r5;
    st.global.s32   [%r10+0], %r8;
$Lt_1_1026:
    .loc    16  27  0
    exit;
$LDWend__Z4foo2Pi:
    } // _Z4foo2Pi

Обратите внимание, что во втором случае к constBuf_d обращаются через ld.const.s32, а не ld.global.s32, так что используется постоянный кэш памяти.

Answer 2

Отличный ответ @talonmies.Но я хотел бы отметить, что в cuda 5 произошли изменения. В функции MemcpyToSymbol () аргумент char * больше не поддерживается.

Примечания к выпуску CUDA 5 гласят:

** The use of a character string to indicate a device symbol, which was possible with certain API functions, is no longer supported. Instead, the symbol should be used directly.

Вместо этого необходимо сделать копию в постоянную память следующим образом:

cudaMemcpyToSymbol( dev_x, x, N * sizeof(float) );

В этом случае "dev_x"- указатель на постоянную память, а" x "- указатель на память хоста, которую необходимо скопировать в dev_x.