グローバルメモリ変数を初期化した後のグローバル__syncthreads（）

Question

グローバルカウンタをゼロに設定し、いくつかの計算を行い、共有カウンタをインクリメントしてから、グローバルカウンタに異なるブロックの共有カウンタを追加するCUDAカーネルがあります。今私は、グローバルカウンタのインクリメントの前に初期化が行われることを確認したい。ここに私のカーネルの簡単な例です：

__global__ void counter_function(int32_t* global_count, int32_t* other_data) 
{ 
    __shared__ int32_t block_count; 
    bool condition = false; 

    // first thread of first block initializes global counter 
    if (thread.Idx == 0 && block.Idx == 0) { 
     *global_count = 0; 
    } 

    // do some calculation on other_data and possibliy set 
    // condition = true; 

    // increment counter if condition is true 
    if (condition) { 
     atomicAdd(&block_count, 1); 
    } 

    __synchthreads(); 

    // first thread of each block adds its count to the global count 
    if (thread.Idx == 0) { 
     atomicAdd(&global_count, block_count); 
    } 
} 

私はCUDA Cプログラミング・ガイドのMemory Fence FunctionsとSynchronization Functionsを読んでいるが、__syncthreads()が唯一のブロックごとにスレッドを同期して__threadfence()だけの世話をするので、解決策を見つけることができませんでした私が理解している限り、スレッド内のグローバル読み書き操作の順序。また、グローバルカウントを初期化するスレッドを除くすべてのスレッドは、グローバル書き込み操作を1つしか持たないため、私の場合でも使用できません。私は__threadfence()の機能を正しく理解していますか？

1つの方法は、カーネルを2つのカーネルに分割することです。そのため、基本的にカウンタをゼロに設定するだけの別個のカーネルが必要です。しかし、より良い方法がありますか？

これを解決する方法を教えてください。または私は何かを逃していますか？

Answer 1

CUDAには、ブロック間で「同期スレッド」できる機能はありません。また、@ Robert Crovellaがすでに述べたように、スレッドの実行順序は不明です。 2つの別個のカーネルを作成し、前者の実行後にcudaDeviceSynchronizeを呼び出して、要求された操作が実行されたことを確認する必要があります。

さらに、あなたのコードのatomicAddは不要であるようです。各スレッドが1つの別々の要素を更新してから、ベクトル全体に対してReduce操作を実行できるintのベクトルを持つ方がよいでしょう。