カーネルの定数メモリにアクセスするにはどうすればいいですか？

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

何らかの理由で、私はresults_hに "0"しか得られません。私は能力1.1のカードを使ってCUDA 4.0を走らせています。

アイデア？ありがとう！

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

は動作します。しかし、一定のメモリアドレスを引数としてカーネルに渡すことは、定数メモリを使用する間違った方法です。つまり、定数メモリキャッシュを介してメモリにアクセスする命令をコンパイラが生成するのを忘れてしまいます。

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

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

後者の場合には、constBuf_dがld.const.s32を介してアクセスされていることを

.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 

注意を生成します。このカーネルで

.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_ 

：1.2 PTXは、あなたのカーネル用に生成された計算機能の比較、ld.global.s32ではなく、定数メモリキャッシュが使用されます。

Answer 2

優れた答え@タロン。しかし、私は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にコピーする必要のあるホストメモリへのポインタです。