2

シングルスレッドで実行されるステートメントをCUDAカーネルに書き込むにはどうすればよいですか。たとえば、次のカーネルがある場合：

__global__ void Kernel(bool *d_over, bool *d_update_flag_threads, int no_nodes)
{
   int tid = blockIdx.x*blockDim.x + threadIdx.x;
   if( tid<no_nodes && d_update_flag_threads[tid])
   {
     ...
     *d_over=true; // writing a single memory location, only 1 thread should do?
     ...
   }
}

上記のカーネルでは、「d_over」は単一のブールフラグであり、「d_update_flag_threads」はブール配列です。

私が以前に通常行ったことは、スレッドブロックの最初のスレッドを使用することです。

if(threadIdx.x==0)

ただし、ここにフラグ配列があり、関連付けられたフラグが「true」のスレッドのみがifステートメントを実行するため、この場合は機能しませんでした。そのフラグ配列は、以前に呼び出された別のCUDAカーネルによって設定されており、事前にそれについての知識はありません。

つまり、OpenMPの「Single」コンストラクトに似たものが必要です。

4

2 に答える 2

3

考えられるアプローチは、不可分操作を使用することです。更新を実行するためにブロックごとに1つのスレッドのみが必要な場合は、共有メモリでアトミック操作を実行できます（計算機能> = 1.2の場合）。これは、グローバルメモリで実行するよりも一般的にはるかに高速です。

とはいえ、アイデアは次のとおりです。

int tid = blockIdx.x*blockDim.x + threadIdx.x;

__shared__ int sFlag;
// initialize flag
if (threadIdx.x == 0) sFlag = 0;
__syncthreads();

if( tid<no_nodes && d_update_flag_threads[tid])
{
  // safely update the flag
  int singleFlag = atomicAdd(&sFlag, 1);
  // custom single operation
  if ( singleFlag == 0) 
      *d_over=true; // writing a single memory location, only 1 thread will do it
       ...
}

それは単なるアイデアです。私はそれをテストしていませんが、ブロックの最初のスレッドではなく、単一のスレッドによって実行される操作に近いです。

于 2012-06-05T14:10:51.317 に答える
0

t_overが宣言されている場合はatomicCAS（d_over、0、1）を使用するか、int*として型キャストすることができます。これにより、d_over値が0（false）であると見なされる最初のスレッドのみが更新され、他のスレッドは更新されないことが保証されます。

于 2012-06-05T23:47:13.023 に答える