55

...または現在の経糸またはブロックの糸だけですか？

また、特定のブロックのスレッドが（カーネル内で）次の行に遭遇した場合

__shared__  float srdMem[128];

彼らはこのスペースを（ブロックごとに）一度だけ宣言しますか？

それらはすべて明らかに非同期で動作するため、ブロック22のスレッド23がこの行に到達する最初のスレッドであり、ブロック22のスレッド69がこの行に到達する最後のスレッドである場合、スレッド69はすでに宣言されていることを認識しますか？

4

5 に答える 5

74

この__syncthreads()コマンドは、ブロックレベルの同期バリアです。つまり、ブロック内のすべてのスレッドがバリアに到達したときに安全に使用できます。条件付きコードで使用することもでき__syncthreads()ますが、すべてのスレッドがそのようなコードを同じように評価する場合にのみ、実行がハングしたり、意図しない副作用が発生したりする可能性があります[4]。

使用例__syncthreads():(ソース）

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

グリッド内のすべてのスレッドを同期するために、現在ネイティブAPI呼び出しはありません。グリッドレベルでスレッドを同期する1つの方法は、その時点ですべてのスレッドが終了し、同じポイントから再開するため、連続したカーネル呼び出しを使用することです。これは、一般にCPU同期または暗黙的同期とも呼ばれます。したがって、それらはすべて同期されます。

この手法の使用例（ソース）：

CPU同期

2番目の質問について。はい、ブロックごとに指定された共有メモリの量を宣言します。使用可能な共有メモリの量はSMごとに測定されることを考慮に入れてください。したがって、起動構成とともに共有メモリがどのように使用されるかについては、十分に注意する必要があります。

于 2013-03-06T08:27:36.697 に答える
19

ここでのすべての回答に同意しますが、最初の質問で重要な点が1つ欠けていると思います。上記の回答で完全に回答されたため、2番目の回答には回答していません。

GPUでの実行は、ワープの単位で行われます。ワープは32スレッドのグループであり、一度に特定のワープの各スレッドが同じ命令を実行します。ブロックに128スレッドを割り当てると、GPUに（128/32 =）4ワープします。

ここで、質問は「すべてのスレッドが同じ命令を実行している場合、なぜ同期が必要なのですか？」になります。答えは、 SAMEブロックに属するワープを同期する必要があるということです。__syncthreadsは、ワープ内のスレッドを同期しません。それらはすでに同期されています。同じブロックに属するワープを同期します。

そのため、質問に対する答えは次のとおりです。__syncthreadsはグリッド内のすべてのスレッドを同期しませんが、各ブロックが独立して実行されるため、1つのブロックに属するスレッドを同期します。

グリッドを同期する場合は、カーネル（K）を2つのカーネル（K1とK2）に分割し、両方を呼び出します。それらは同期されます（K1が終了した後にK2が実行されます）。

于 2016-12-31T15:20:45.647 に答える
17

__syncthreads()同じブロック内のすべてのスレッドがコマンドに到達し、ワープ内のすべてのスレッドが到達するまで待機します。つまり、スレッドブロックに属するすべてのワープがステートメントに到達する必要があります。

カーネルで共有メモリを宣言すると、配列は1つのスレッドブロックにのみ表示されます。したがって、各ブロックには独自の共有メモリブロックがあります。

于 2013-03-06T06:28:50.787 に答える
7

既存の回答は、どのように機能するか（ブロック内__syncthreads()同期を可能にする）に答える素晴らしい仕事をしました。ブロック間同期の新しい方法があるという更新を追加したかっただけです。CUDA 9.0以降、ブロックのグリッド全体を同期できる「協調グループ」が導入されました（Cudaプログラミングガイドで説明されています）。これにより、（上記のように）新しいカーネルを起動するのと同じ機能が実現されますが、通常はオーバーヘッドが低くなり、コードが読みやすくなります。

于 2020-04-05T17:35:44.247 に答える
2

詳細を提供するために、回答は別として、seibertを引用します。

より一般的には、__ syncthreads（）は、ブロック内の書き込み後の読み取りメモリの競合状態からユーザーを保護するように設計されたバリアプリミティブです。

使用規則は非常に単純です。

  1. 別のスレッドが書き込んだメモリ位置をスレッドが読み取る可能性がある場合は、書き込み後、読み取り前に__syncthreads（）を配置します。

  2. __syncthreads（）はブロック内のバリアにすぎないため、同じブロック内のスレッド間でのみ発生する可能性のある競合がない限り、グローバルメモリでの書き込み後の競合状態からユーザーを保護することはできません。__syncthreads（）は、ほとんどの場合、書き込み後の共有メモリを保護するために使用されます。

  3. すべてのスレッドが同じ__syncthreads（）呼び出しに到達することが確実になるまで、ブランチまたはループで__syncthreads（）呼び出しを使用しないでください。これにより、ifブロックをいくつかの部分に分割して、すべてのスレッド（if述語に失敗したスレッドを含む）が実行するトップレベルに__syncthread（）呼び出しを配置する必要がある場合があります。

  4. ループ内の書き込み後の読み取り状況を探す場合、__ syncthread（）呼び出しを配置する場所を見つけるときに、頭の中でループを展開するのに役立ちます。たとえば、ループ内の同じ共有メモリ位置への異なるスレッドからの読み取りと書き込みがある場合、ループの最後に追加の__syncthreads（）呼び出しが必要になることがよくあります。

  5. __syncthreads（）はクリティカルセクションをマークしないので、そのように使用しないでください。

  6. カーネル呼び出しの最後に__syncthreads（）を配置しないでください。その必要はありません。

  7. 2つの異なるスレッドが同じメモリ位置にアクセスすることはないため、多くのカーネルは__syncthreads（）をまったく必要としません。

于 2021-01-08T21:41:42.647 に答える