3

私は現在CUDAを学んでおり、いくつかの演習に取り組んでいます。その 1 つは、要素ごとに 1 スレッド、行ごとに 1 スレッド、列ごとに 1 スレッドの 3 つの異なる方法で行列を追加するカーネルを実装することです。行列は正方形で、1D ベクトルとして実装されています。

A[N*row + col]

直感的には、最初のオプションはスレッドのオーバーヘッドが原因で最も遅く、2 番目のオプションは単一のスレッドが隣接するデータを処理するため、最も高速であると予想しました。

CPU で、8000 x 8000 の密行列を使用すると、次のようになります。

Adding on CPU - Adding down columns
Compute Time Taken: 2.21e+00 s
Adding on CPU - Adding across rows
Compute Time Taken: 2.52e-01 s

したがって、より多くのキャッシュヒットにより、約1桁の速度が向上します。同じ行列を持つ GPU では、次のようになります。

Adding one element per thread 
Compute Time Taken: 7.42e-05 s
Adding one row per thread 
Compute Time Taken: 2.52e-05 s
Adding one column per thread 
Compute Time Taken: 1.57e-05 s

これは私には直感的ではありません。最後のケースの 30 ～ 40% の高速化は、約 1000 x 1000 行列以上で一貫しています。これらのタイミングはカーネル実行のみであり、ホストとデバイス間のデータ転送は含まれていないことに注意してください。以下は、比較のための私の 2 つのカーネルです。

__global__
void matAddKernel2(float* A, float* B, float* C, int N)
{
        int row = threadIdx.x + blockDim.x * blockIdx.x;
        if (row < N)
        {
                int j;
                for (j = 0; j < N; j++)
                {
                        C[N*row + j] = A[N*row + j] + B[N*row + j];
                }
        }
}



__global__
void matAddKernel3(float* A, float* B, float* C, int N)
{
        int col = threadIdx.x + blockDim.x * blockIdx.x;
        int j;

        if (col < N)
        {
                for (j = 0; j < N; j++)
                {
                        C[col + N*j] = A[col + N*j] + B[col + N*j];
                }
        }
}

私の質問は、GPU スレッドが隣接するデータを操作することで利益を得ないように思われるのはなぜですか?

4

1 に答える 1

5

GPU スレッドは、隣接するデータを操作することで恩恵を受けます。欠けているのは、GPU スレッドは CPU スレッドのような独立したスレッドではなく、warp と呼ばれるグループで動作することです。ワープは 32 のスレッドをグループ化し、幅 32 の SIMD 命令を実行する単一の CPU スレッドと同様の方法で動作します。

したがって、実際には、列ごとに 1 つのスレッドを使用するコードが最も効果的です。これは、ワープ内の隣接するスレッドがメモリから隣接するデータ位置にアクセスし、それがグローバル メモリにアクセスする最も効果的な方法であるためです。

詳細については、CUDA のドキュメントを参照してください。

于 2013-05-31T20:39:58.490 に答える