CUDAカーネルが配列をスキップする

Question

私はCUDAを使うことを学んでおり、実際には2つの大きな配列（〜512百万）の追加を計算したいので、カーネルにforループを追加しようとしました。なんらかの理由で、いつも特定のエントリをスキップします。

この例では、エラーは常に155であり、Nに10を掛けても、常に155のままです。私は別の数字を試してきましたが、常にエラーがあります。私はCPU上で同様のループ構造を実行し、エラーは発生しません。また、アルゴリズムがすべてのインデックスをカバーすることを確認するためにステップバイステップで歩いていて、何がうまくいかないのか理解できません。

GPUassert: no error D:/Projects/_Test/cpp/CudaTest4/CudaTest4/kernel.cu 49 
3474374 4194304: 49 
GPUassert: no error D:/Projects/_Test/cpp/CudaTest4/CudaTest4/kernel.cu 51 
3461574 4194304: 51 
GPUassert: no error D:/Projects/_Test/cpp/CudaTest4/CudaTest4/kernel.cu 53 
3448774 4194304: 53 
n: 3276800 
n/THREAD:3200.00 
n/THREAD/CHUNK :100.00 
GPUassert: no error D:/Projects/_Test/cpp/CudaTest4/CudaTest4/kernel.cu 68 
GPU took 2.000000 msec to Calculate. 
GPUassert: no error D:/Projects/_Test/cpp/CudaTest4/CudaTest4/kernel.cu 71 
GPU took 17.000000 msec to Copy Memory. 
Number of Errors: 155 
CUDA: Kernel Finished 
CUDA: Memory Freed 
CPU took 3.000000 msec. 
Number of Errors: 0 

コードは以下である：以下

#define _N (512*2*32*10) 
#define THREAD_PER_BLOCK (512*2) 
#define CHUNK 32 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) 
{ 
     fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
} 

__global__ void add(const int *a, const int *b, int *c, int n, int chunk) 
{ 
    int index = (threadIdx.x + blockIdx.x * blockDim.x) * chunk; 
    #pragma unroll 
    for (int i = 0; i < chunk; i++){ 
     index += i; 
     if (index < n) { 
      c[index] = a[index] + b[index]; 
     } 
    } 
} 

void random_ints(int* a, int Size) 
{ 
    for (int i = 0; i < Size; i++){ 
     a[i] = rand(); 
    } 
} 

void _run_kernel(int *a, int *b, int *c, int n) 
{ 
    FILE *file = fopen("error.log", "wb+"); 

    cudaDeviceReset(); 
    size_t f; 
    size_t t; 
    CUresult result; 

    int *d_a, *d_b, *d_c; 
    int size = n * sizeof(int); 

    gpuErrchk(cudaMalloc((void **)&d_a, size)); 
    result = cuMemGetInfo(&f, &t); fprintf(stdout, "%d %d: %d\n", f/1024, t/1024); 
    gpuErrchk(cudaMalloc((void **)&d_b, size)); 
    result = cuMemGetInfo(&f, &t); fprintf(stdout, "%d %d: %d\n", f/1024, t/1024); 
    gpuErrchk(cudaMalloc((void **)&d_c, size)); 
    result = cuMemGetInfo(&f, &t); fprintf(stdout, "%d %d: %d\n", f/1024, t/1024); 

    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice); 

    printf("n: %i\n", n); 
    float block = n; 
    block /= THREAD_PER_BLOCK; 
    printf("n/THREAD:%.2f\n", block); 
    block /= CHUNK; 
    printf("n/THREAD/CHUNK :%.2f\n", block); 
    clock_t t0, t1, t2; 
    t0 = clock(); 
    add <<<block, THREAD_PER_BLOCK >>>(d_a, d_b, d_c, n, CHUNK); 
    gpuErrchk(cudaPeekAtLastError()); 
    t1 = clock(); 
    printf("GPU took %f msec to Calculate.\n", (double)((t1 - t0)/1)); 
    gpuErrchk(cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost)); 
    t1 = clock(); 
    printf("GPU took %f msec to Copy Memory.\n", (double)((t1 - t0)/1)); 

    int err = 0; 
    for (int i = 0; i < n; i++){ 
     if (a[i] + b[i] != c[i]){ 
      fprintf(file,"%i: %i + %i = %i", i, a[i], b[i], c[i]); 
      fprintf(file,a[i] + b[i] == c[i] ? " CORRECT\n" : " WRONG\n"); 
      err += (a[i] + b[i] == c[i]) ? 0 : 1; 
     } 
     else { 
      fprintf(file, "__%i: %i + %i = %i", i, a[i], b[i], c[i]); 
      fprintf(file, a[i] + b[i] == c[i] ? " CORRECT\n" : " WRONG\n"); 
     } 
    } 
    printf("Number of Errors: %i\n", err); 

    printf("CUDA: Kernel Finished\n"); 

    //free(a); free(b); free(c); 
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); 

    printf("CUDA: Memory Freed\n"); 
} 

int main(void){ 
    int *a, *b, *c; 
    int size = _N * sizeof(int); 

    a = (int *)malloc(size); random_ints(a, _N); 
    b = (int *)malloc(size); random_ints(b, _N); 
    c = (int *)malloc(size); random_ints(c, _N); 

    _run_kernel(a, b, c, _N); 

    free(a); free(b); free(c); 

    return 0; 
} 

エラーログの出力のビットは、ここでのコードのために出力されます。あなたが見ることができるように、エラーはかなり無作為に分布しており、index = 464の後にエラーはそれ以上ありません。

__0: 41 + 18155 = 18196 CORRECT 
__1: 18467 + 31709 = 50176 CORRECT 
2: 6334 + 26000 = 0 WRONG 
__3: 26500 + 18702 = 45202 CORRECT 
4: 19169 + 20035 = 0 WRONG 
5: 15724 + 15174 = 0 WRONG 
__6: 11478 + 6984 = 18462 CORRECT 
7: 29358 + 28504 = 0 WRONG 
8: 26962 + 6996 = 0 WRONG 
9: 24464 + 20874 = 0 WRONG 
__10: 5705 + 10331 = 16036 CORRECT 
11: 28145 + 32699 = 0 WRONG 
12: 23281 + 19859 = 0 WRONG 
13: 16827 + 29141 = 0 WRONG 
14: 9961 + 32411 = 0 WRONG 
__15: 491 + 17109 = 17600 CORRECT 
16: 2995 + 28917 = 0 WRONG 
17: 11942 + 1760 = 0 WRONG 
18: 4827 + 29229 = 0 WRONG 
19: 5436 + 16198 = 0 WRONG 
20: 32391 + 10345 = 0 WRONG 
__21: 14604 + 17254 = 31858 CORRECT 
22: 3902 + 22576 = 0 WRONG 
23: 153 + 28099 = 0 WRONG 
24: 292 + 24998 = 0 WRONG 
25: 12382 + 5848 = 0 WRONG 
26: 17421 + 30879 = 0 WRONG 
27: 18716 + 29542 = 0 WRONG 
__28: 19718 + 30248 = 49966 CORRECT 
29: 19895 + 23121 = 0 WRONG 
30: 5447 + 22393 = 0 WRONG 
31: 21726 + 22088 = 0 WRONG 
__32: 14771 + 23925 = 38696 CORRECT 
__33: 11538 + 18892 = 30430 CORRECT 
34: 1869 + 6431 = 0 WRONG 
__35: 19912 + 24658 = 44570 CORRECT 
__36: 25667 + 21669 = 47336 CORRECT 
37: 26299 + 7573 = 0 WRONG 
__38: 17035 + 1533 = 18568 CORRECT 
39: 9894 + 9296 = 0 WRONG 
40: 28703 + 7969 = 0 WRONG 
41: 23811 + 22525 = 0 WRONG 
__42: 31322 + 29036 = 60358 CORRECT 
43: 30333 + 22855 = 0 WRONG 
44: 17673 + 17579 = 0 WRONG 
__45: 4664 + 6994 = 11658 CORRECT 
46: 15141 + 2007 = 0 WRONG 
__47: 7711 + 9 = 7720 CORRECT 
48: 28253 + 28831 = 0 WRONG 
49: 6868 + 7182 = 0 WRONG 
50: 25547 + 18461 = 0 WRONG 
51: 27644 + 1798 = 0 WRONG 
52: 32662 + 22136 = 0 WRONG 
__53: 32757 + 847 = 33604 CORRECT 
54: 20037 + 11319 = 0 WRONG 
__55: 12859 + 24677 = 37536 CORRECT 
56: 8723 + 16277 = 0 WRONG 
57: 9741 + 21895 = 0 WRONG 
58: 27529 + 17691 = 0 WRONG 
59: 778 + 15956 = 0 WRONG 
__60: 12316 + 9790 = 22106 CORRECT 
61: 3035 + 12661 = 0 WRONG 
62: 22190 + 19936 = 0 WRONG 
63: 1842 + 27036 = 0 WRONG 
__64: 288 + 482 = 770 CORRECT 
__65: 30106 + 31572 = 61678 CORRECT 
__66: 9040 + 31266 = 40306 CORRECT 
__67: 8942 + 26232 = 35174 CORRECT 

すべてのヘルプは本当にいただければ幸い、私はこの上で一日無駄にし、本当に行うにはシンプルなものでなければなりませんように思えますか？

EDIT：私のCPUコード

int index = (threadIdx.x + blockIdx.x * blockDim.x) * chunk; 
for (int i = 0; i < n/THREAD_PER_BLOCK/CHUNK; i++){ 
    for (int j = 0; j < THREAD_PER_BLOCK; j++){ 
     int index = j + i * (n/THREAD_PER_BLOCK); 
     index *= chunk; 
     for (int k = 0; k < CHUNK; k++){ 
      index += k; 
      if (index < n) { 
       c[index] = a[index] + b[index]; 
      } 
     } 
    } 
} 

Answer 1

カーネルであなたのインデックスロジックが台無しにされます。これに

__global__ void add(const int *a, const int *b, int *c, int n, int chunk) 
{ 
    int index = (threadIdx.x + blockIdx.x * blockDim.x) * chunk; 
    #pragma unroll 
    for (int i = 0; i < chunk; i++){ 
     index += i; 
     if (index < n) { 
      c[index] = a[index] + b[index]; 
     } 
    } 
} 

：あなたはこれを変更した場合

__global__ void add(const int *a, const int *b, int *c, int n, int chunk) 
{ 
    int index = (threadIdx.x + blockIdx.x * blockDim.x) * chunk; 
    #pragma unroll 
    for (int i = 0; i < chunk; i++){ 
     if (index < n) { 
      c[index] = a[index] + b[index]; 
     } 
     index++; 
    } 
} 

あなたのコードはエラーなしで実行されます。これはCUDAとはまったく関係がありませんが、基本的なCプログラミングの問題です。上記インデックスで

iteration (i): 0 1 2 3 4 ... 
index:  +0 +1 +3 +6 +10 ... 

ギャップがスキップされたエントリを引き起こす：

index += i; 

は、インデックスが連続する反復を通じて、このようになりますindex変数を作成すること：インデックスのこの種のは、あなたが望むものではありません。 indexには、反復ごとに1ずつ増やすだけです。あなたのCPUループは、各繰り返しでインデックスを1だけ大きくします。

for (int i = 0; i < n; i++){