CUDAで有効な値を取得できません

Question

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 

#include <stdio.h> 
#include <stdlib.h> 

#define NUMBEROFMX 256*64 

__global__ void reduce0(int *g_idata, int *g_odata) 
{ 
    extern __shared__ int sdata[]; 
    unsigned int tid = threadIdx.x; 
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; 
    sdata[tid] = g_idata[i]; 
    __syncthreads(); 


    for (unsigned int s = 1; s < blockDim.x; s *= 2) 
    { 
     if (tid % (2 * s) == 0) 
     { 
      sdata[tid] += sdata[tid + s]; 
     } 
     __syncthreads(); 
    } 
    if (tid == 0) 
    { 
     g_odata[blockIdx.x] = sdata[0]; 
    } 
} 

int main() 
{ 
    int *A; 
    int *B; 
    int *dev_A; 
    int *dev_B; 


    A = (int*)malloc(sizeof(int) * NUMBEROFMX); 
    B = (int*)malloc(sizeof(int) * NUMBEROFMX); 

    cudaMalloc((void**)&dev_A, sizeof(int)*NUMBEROFMX); 
    cudaMalloc((void**)&dev_B, sizeof(int)*NUMBEROFMX); 

    for (int i = 0; i < NUMBEROFMX; i++) 
    { 
     A[i] = 1; 
    } 

    cudaMemcpy(dev_A, A, sizeof(int)*NUMBEROFMX, cudaMemcpyHostToDevice); 
    reduce0 << <256, 64 >> >(dev_A, dev_B); 
    cudaMemcpy(B, dev_B, sizeof(int)*NUMBEROFMX, cudaMemcpyDeviceToHost); 
    printf("%d\n", B[0]); 
} 

私はCUDAプログラミングについて学びます。このコードは完了していませんが、私は木の削減と配列の合計のコードを作成したいです。私はB [0]に64を得ることを期待していますが、B [0]の値は有効ではありません。 NSIGHTを使用してこのコードをデバッグすると、g_odata [0]は64ですが、B [0]は無効です。どうしてか分かりません。

Answer 1

proper CUDA error checkingを使用した場合、またはcuda-memcheckでコードを実行した場合は、共有メモリ（範囲外の共有アドレスまたはローカルアドレス）が不正に使用されている可能性があります。

使用しようとしている動的な共有メモリ、について、あなたはカーネル呼び出しでブロックごとの共有メモリのサイズを指定する必要があります。これにより

reduce0<<<256, 64, 64*sizeof(int)>>>(dev_A, dev_B); 

期待通りにカーネルの動作を変更します。

共有メモリの詳細については、ブログ記事Using Shared Memory in CUDA C/C++を参照してください。