6

CUDA プログラミングでは、共有メモリを使用する場合、データをグローバル メモリから共有メモリに移動する必要があります。このようなデータの転送にはスレッドが使用されます。

グローバルメモリから共有メモリにデータをコピーするためにブロック内のすべてのスレッドを関与させない方がよいことを（オンラインリソースの）どこかで読みました。このような考え方は、すべてのスレッドが一緒に実行されるわけではないという点で理にかなっています。ワープ内のスレッドは一緒に実行されます。しかし、私の懸念は、すべてのワープが順番に実行されていないことです。たとえば、スレッドを含むブロックは 3 つのワープに分割されます: war p0 (0 ～ 31 スレッド)、warp 1 (32 ～ 63 スレッド)、warp 2 (64 ～ 95 スレッド)。warp 0 が最初に実行されるとは限りません (そうですか?)。

では、データをグローバル メモリから共有メモリにコピーするには、どのスレッドを使用すればよいでしょうか。

4

2 に答える 2

7

単一のワープを使用して共有メモリ配列をロードするには、次のようにします。

__global__
void kernel(float *in_data)
{
    __shared__ float buffer[1024];

    if (threadIdx.x < warpSize) {
        for(int i = threadIdx; i  <1024; i += warpSize) {
            buffer[i] = in_data[i];
        }
    }
    __syncthreads();

    // rest of kernel follows
}

[免責事項: ブラウザで書かれ、テストされていません。自己責任で使用してください]

ここでの重要なポイントは__syncthreads()、ブロック内のすべてのスレッドが、共有メモリへのロードを実行する warp がロードを完了するまで待機することを保証するために を使用することです。私が投稿したコードは最初のワープを使用しましたが、ブロック内のスレッド インデックスを warpSize で割ることでワープ数を計算できます。また、1 次元ブロックを想定しました。2D または 3D ブロックでスレッド インデックスを計算するのは簡単なので、それは読者の課題として残します。

于 2013-03-18T07:03:32.447 に答える
0

ブロックがマルチプロセッサに割り当てられると、そのブロック内のすべてのスレッドが終了するまでそこに存在し、この間、ワープ スケジューラは準備完了オペランドを持つワープ間で混合します。したがって、マルチプロセッサに 3 つのワープを持つ 1 つのブロックがあり、1 つのワープだけがグローバルから共有メモリにデータをフェッチし、他の 2 つのワープがアイドル状態のままでおそらく__syncthreads()バリアを待機している場合、失うものは何もなく、グローバル メモリのレイテンシだけによって制限されます。とにかくいたでしょう。フェッチが完了するとすぐに、ワープは作業を進めることができます。

したがって、warp0 が最初に実行されるという保証は必要なく、任意のスレッドを使用できます。心に留めておくべき唯一の 2 つのことは、グローバル メモリへの結合されたアクセスを可能な限り確保することと、バンクの競合を回避することです。

于 2013-03-18T00:54:39.183 に答える