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.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

您可以使用 atomicCAS(d_over, 0, 1) ，其中 d_over 被声明或类型转换为 int*。这将确保只有第一个看到 d_over 值为 0（假）的线程会更新它，而其他人不会。

于 2012-06-05T23:47:13.023 回答