いくつかの要素は、ここでCUDA

Question

上のベクトル加算で計算されていなかったコードは次のとおりです。

#include "common/book.h" 

#define N 36 

__global__ void add(int *a, int *b, int *c) { 
    int tid = blockIdx.x * gridDim.y * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x; 
    if(tid < N) { 
     c[tid] = a[tid] + b[tid]; 
    } 
} 

int main() { 
    int a[N], b[N], c[N]; 
    int *dev_a, *dev_b, *dev_c; 
    cudaMalloc((void**) &dev_a, N * sizeof(int)); 
    cudaMalloc((void**) &dev_b, N * sizeof(int)); 
    cudaMalloc((void**) &dev_c, N * sizeof(int)); 
    for (int i = 0; i < N; i++) { 
     a[i] = -1; 
     b[i] = i * i; 
    } 

    cudaMemcpy(
       dev_a, 
       a, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    cudaMemcpy(
       dev_b, 
       b, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    dim3 grid_dim(3, 2); 
    dim3 block_dim(3, 2); 
    add<<<grid_dim, block_dim>>>(dev_a, dev_b, dev_c); 
    cudaMemcpy(
       c, 
       dev_c, 
       N * sizeof(int), 
       cudaMemcpyDeviceToHost 
        ); 
    for (int i = 0; i < N; i++) { 
     printf("%d + %d = %d\n", a[i], b[i], c[i]); 
    } 
    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c); 
} 

基本的に、私は、3×2のレイアウトとグリッド上、二つのベクトル要素単位を追加することで、各ブロックをしようとしていましたグリッドはスレッドの3x2レイアウトを有する。

-1 + 0 = -1 
-1 + 1 = 0 
-1 + 4 = 3 
-1 + 9 = 8 
-1 + 16 = 15 
-1 + 25 = 24 
-1 + 36 = 0 
-1 + 49 = 0 
-1 + 64 = 0 
-1 + 81 = 0 
-1 + 100 = 0 
-1 + 121 = 0 
-1 + 144 = 143 
-1 + 169 = 168 
-1 + 196 = 195 
-1 + 225 = 224 
-1 + 256 = 255 
-1 + 289 = 288 
-1 + 324 = 0 
-1 + 361 = 0 
-1 + 400 = 0 
-1 + 441 = 0 
-1 + 484 = 0 
-1 + 529 = 0 
-1 + 576 = 575 
-1 + 625 = 624 
-1 + 676 = 675 
-1 + 729 = 728 
-1 + 784 = 783 
-1 + 841 = 840 
-1 + 900 = 0 
-1 + 961 = 0 
-1 + 1024 = 0 
-1 + 1089 = 0 
-1 + 1156 = 0 
-1 + 1225 = 0 

はどうやらいくつかのブロックは単に無視されています。私は、コンパイルされたバイナリを実行したときにここで

は結果です。また、tidがカーネル関数addでどのように計算されているかを試してみましたが、いくつかのブロックが欠落しています。

提案がありますか？

Answer 1

唯一の問題は、あなたがすでに推測している通り、tidの計算によるものです。

マッピングを実行したり、算術演算を作成する方法はたくさんあります。一般的な2Dグリッドでは、xとyに2Dインデックスを作成し、yインデックスとxインデックスを掛けたグリッド幅（x）を使用すると便利です（つまり、覚えやすい方法）スレッド固有の1-Dのインデックスを作成する：

int idy = threadIdx.y+blockDim.y*blockIdx.y; // y-index 
int idx = threadIdx.x+blockDim.x*blockIdx.x; // x-index 
int tid = gridDim.x*blockDim.x*idy + idx;  // thread-unique 1D index 

gridDim.x*blockDim.x

はxのグリッド幅であり、スレッドの単位で表されます。

私たちはあなたのコードでは、この汎用2Dインデックススキームを使用する場合、私のために正常に動作するようです：

$ cat t10.cu 
#include <stdio.h> 

#define N 36 

__global__ void add(int *a, int *b, int *c) { 
    int idy = threadIdx.y+blockDim.y*blockIdx.y; 
    int idx = threadIdx.x+blockDim.x*blockIdx.x; 
    int tid = gridDim.x*blockDim.x*idy + idx; 
    if(tid < N) { 
     c[tid] = a[tid] + b[tid]; 
    } 
} 

int main() { 
    int a[N], b[N], c[N]; 
    int *dev_a, *dev_b, *dev_c; 
    cudaMalloc((void**) &dev_a, N * sizeof(int)); 
    cudaMalloc((void**) &dev_b, N * sizeof(int)); 
    cudaMalloc((void**) &dev_c, N * sizeof(int)); 
    for (int i = 0; i < N; i++) { 
     a[i] = -1; 
     b[i] = i * i; 
    } 

    cudaMemcpy(
       dev_a, 
       a, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    cudaMemcpy(
       dev_b, 
       b, 
       N * sizeof(int), 
       cudaMemcpyHostToDevice 
        ); 
    dim3 grid_dim(3, 2); 
    dim3 block_dim(3, 2); 
    add<<<grid_dim, block_dim>>>(dev_a, dev_b, dev_c); 
    cudaMemcpy(
       c, 
       dev_c, 
       N * sizeof(int), 
       cudaMemcpyDeviceToHost 
        ); 
    for (int i = 0; i < N; i++) { 
     printf("%d + %d = %d\n", a[i], b[i], c[i]); 
    } 
    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c); 
} 
$ nvcc -arch=sm_35 -o t10 t10.cu 
$ cuda-memcheck ./t10 
========= CUDA-MEMCHECK 
-1 + 0 = -1 
-1 + 1 = 0 
-1 + 4 = 3 
-1 + 9 = 8 
-1 + 16 = 15 
-1 + 25 = 24 
-1 + 36 = 35 
-1 + 49 = 48 
-1 + 64 = 63 
-1 + 81 = 80 
-1 + 100 = 99 
-1 + 121 = 120 
-1 + 144 = 143 
-1 + 169 = 168 
-1 + 196 = 195 
-1 + 225 = 224 
-1 + 256 = 255 
-1 + 289 = 288 
-1 + 324 = 323 
-1 + 361 = 360 
-1 + 400 = 399 
-1 + 441 = 440 
-1 + 484 = 483 
-1 + 529 = 528 
-1 + 576 = 575 
-1 + 625 = 624 
-1 + 676 = 675 
-1 + 729 = 728 
-1 + 784 = 783 
-1 + 841 = 840 
-1 + 900 = 899 
-1 + 961 = 960 
-1 + 1024 = 1023 
-1 + 1089 = 1088 
-1 + 1156 = 1155 
-1 + 1225 = 1224 
========= ERROR SUMMARY: 0 errors 
$ 

上記は正しい結果を提供する必要があります。パフォーマンスに関して、これはこのおもちゃの問題の最も効率的なマッピングではないかもしれません。この問題は32の倍数ではないスレッドブロックサイズを持っています。これは一般的に効率的なCUDAプログラミングには推奨されません。この場合の最適なマッピング（パフォーマンス/効率の点で）を試すのではなく、スレッドブロックを再編成してブロックあたり32スレッドの倍数以上を提供することが推奨されます。 16または32のスレッドをブロックのx次元に配置することで、索引付けの理解を容易にし、ほぼ最適なメモリー・アクセス性能を得ることができます。