parallel-processing - CUDA カーネルはどのように起動されますか?

Question

2 つの行列を追加する単純な CUDA アプリケーションを作成しました。それはうまくコンパイルされています。すべてのスレッドによってカーネルがどのように起動されるのか、CUDA 内でどのようなフローが行われるのかを知りたいですか? つまり、すべてのスレッドが行列の各要素をどのように実行するかということです。

これが非常に基本的な概念であることは知っていますが、私はこれを知りません。流れに戸惑います。

Answer 1

ブロックのグリッドを起動します。

ブロックはマルチプロセッサに分割できないように割り当てられます (マルチプロセッサのブロック数によって、使用可能な共有メモリの量が決まります)。

ブロックはさらにワープに分割されます。同じ命令を実行するか、非アクティブな 32 個のスレッドである Fermi GPU の場合 (たとえば、同じワープ内の隣接するループよりも早くループを終了したり、処理ifを行わなかったりして分岐したため)。Fermi GPU では、一度に 1 つのマルチプロセッサで最大 2 つのワープが実行されます。

レイテンシ (つまり、メモリ アクセスまたはデータの依存関係が完了するための実行停止) が発生するたびに、別のワープが実行されます (同じブロックまたは異なるブロックの 1 つのマルチプロセッサに収まるワープの数は、それぞれが使用するレジスタの数によって決まります)。ブロックによって使用されるスレッドと共有メモリの量)。

このスケジューリングは透過的に行われます。つまり、あまり深く考える必要はありません。ただし、定義済みの整数ベクトルthreadIdx(ブロック内のスレッドの位置は?)、blockDim(1 つのブロックの大きさは?)、blockIdx(グリッド内のブロックの位置は?)、およびgridDim(グリッドの大きさは?)を使用することもできます。スレッド間で作業 (読み取り: 入力と出力) を分割します。また、さまざまなタイプのメモリに効果的にアクセスする方法を読みたいと思うかもしれません (単一のトランザクション内で複数のスレッドを処理できるようにするため) - しかし、それはトピックから外れています。

NSight は、専門用語のジャングルを通り抜けると、デバイスで何が起こっているかを把握できるグラフィカル デバッガーを提供します。デバッガーに表示されないもの (ストールの理由やメモリ不足など) に関しては、プロファイラーにも同じことが言えます。

別のカーネルを起動することで、グリッド内のすべてのスレッド (存在するすべて) を同期できます。重複しない順次カーネル実行の場合、それ以上の同期は必要ありません。

1 つのグリッド内 (または 1 つのカーネル実行 - ただし、呼び方は自由) 内のスレッドは、アトミック操作 (算術演算用) または適切なメモリ フェンス (ロードまたはストア アクセス用) を使用して、グローバル メモリ経由で通信できます。

1 つのブロック内のすべてのスレッドを組み込み命令と同期できます__syncthreads()(すべてのスレッドは後でアクティブになりますが、通常どおり、Fermi GPU で実行できるワープは最大で 2 つです)。1 つのブロック内のスレッドは、アトミック操作 (算術演算用) または適切なメモリ フェンス (ロードまたはストア アクセス用) を使用して、共有メモリまたはグローバル メモリを介して通信できます。

前述のように、ワープ内のすべてのスレッドは常に「同期」されていますが、一部のスレッドは非アクティブな場合があります。それらは、共有メモリまたはグローバル メモリ (またはコンピューティング機能 3 を備えた今後のハードウェアでは「レーン スワッピング」) を介して通信できます。アトミック操作 (算術演算用) および volatile 修飾された共有変数またはグローバル変数 (同じワープ内で順次発生するロードまたはストア アクセス) を使用できます。volatile 修飾子は、常にメモリにアクセスし、他のスレッドから状態を確認できないレジスターには決してアクセスしないようにコンパイラーに指示します。

さらに、分岐の決定や整数 (プレフィックス) の合計の計算に役立つワープ全体の投票関数があります。

OK、基本的にはそれだけです。それが役立つことを願っています。良い流れの書き込みがありました:-)。

Answer 2

4*4行列の加算の例を見てみましょう。次元4*4の2つの行列AとBがあります。

int main()
{
 int *a, *b, *c;            //To store your matrix A & B in RAM. Result will be stored in matrix C
 int *ad, *bd, *cd;         // To store matrices into GPU's RAM. 
 int N =4;                 //No of rows and columns.

 size_t size=sizeof(float)* N * N;

 a=(float*)malloc(size);     //Allocate space of RAM for matrix A
 b=(float*)malloc(size);     //Allocate space of RAM for matrix B

//allocate memory on device
  cudaMalloc(&ad,size);
  cudaMalloc(&bd,size);
  cudaMalloc(&cd,size);

//initialize host memory with its own indices
    for(i=0;i<N;i++)
      {
    for(j=0;j<N;j++)
         {
            a[i * N + j]=(float)(i * N + j);
            b[i * N + j]= -(float)(i * N + j);
         }
      }

//copy data from host memory to device memory
     cudaMemcpy(ad, a, size, cudaMemcpyHostToDevice);
     cudaMemcpy(bd, b, size, cudaMemcpyHostToDevice);

//calculate execution configuration 
   dim3 grid (1, 1, 1); 
   dim3 block (16, 1, 1);

//each block contains N * N threads, each thread calculates 1 data element

    add_matrices<<<grid, block>>>(ad, bd, cd, N);

   cudaMemcpy(c,cd,size,cudaMemcpyDeviceToHost);  
   printf("Matrix A was---\n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            printf("%f ",a[i*N+j]);
        printf("\n");
    }

   printf("\nMatrix B was---\n");
   for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            printf("%f ",b[i*N+j]);
        printf("\n");
    }

    printf("\nAddition of A and B gives C----\n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            printf("%f ",c[i*N+j]);   //if correctly evaluated, all values will be 0
        printf("\n");
    }



    //deallocate host and device memories
    cudaFree(ad); 
    cudaFree(bd); 
    cudaFree (cd);

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

    getch();
    return 1;
}

/////Kernel Part

__global__ void add_matrices(float *ad,float *bd,float *cd,int N)
{
  int index;
  index = blockIDx.x * blockDim.x + threadIDx.x            

  cd[index] = ad[index] + bd[index];
}

16*16行列の加算の例を見てみましょう。次元16*16の2つの行列AとBがあります。

まず、スレッド構成を決定する必要があります。GPUデバイスで実行される行列の加法の並列計算を実行するカーネル関数を起動することを想定しています。

これで、1つのカーネル関数で1つのグリッドが起動されます。グリッドには、3次元の方法で配置できるブロックを最大65,535個持つことができます。（65535 * 65535 * 65535）。

グリッド内のすべてのブロックには、最大1024のスレッドを含めることができます。これらのスレッドは、3次元の方法で配置することもできます（1024 * 1024 * 64）

ここでの問題は、16*16行列の加算です。

A | 1  2  3  4 |        B | 1  2  3  4 |      C| 1  2  3  4 |
  | 5  6  7  8 |   +      | 5  6  7  8 |   =   | 5  6  7  8 | 
  | 9 10 11 12 |          | 9 10 11 12 |       | 9 10 11 12 |  
  | 13 14 15 16|          | 13 14 15 16|       | 13 14 15 16|

計算を実行するには16のスレッドが必要です。

i.e. A(1,1) + B (1,1) = C(1,1)
     A(1,2) + B (1,2) = C(1,2) 
     .        .          .
     .        .          . 
     A(4,4) + B (4,4) = C(4,4) 

これらのスレッドはすべて同時に実行されます。したがって、16スレッドのブロックが必要です。便宜上、1つのブロックに（16 * 1 * 1）の方法でスレッドを配置します。16のスレッドはないため、これらの16のスレッドを格納するために必要なブロックは1つだけです。

したがって、グリッド構成は次のようになります。dim3 Grid(1,1,1)つまり、グリッドには1つのブロックのみが含まれ、ブロック構成は次のようになりますdim3 block(16,1,1)。つまり、ブロックには16個のスレッドが列方向に配置されます。

次のプログラムは、その実行について明確なアイデアを提供します。インデックス作成部分（つまり、threadIDs、blockDim、blockID）を理解することは重要な部分です。CUDAの文献を読む必要があります。インデックス作成について明確なアイデアが得られたら、ハーフバトルに勝ちます！だから、もちろん、cudaの本、さまざまなアルゴリズム、紙の鉛筆で時間を過ごしてください！

Answer 3

CUDA デバッガーである'Cuda-gdb'を試してください。