Cuda 原子更改标志

Question

我有一段串行代码，它做这样的事情

if( ! variable )
{
  do some initialization here 
  variable = true;
}

我知道这在串行中工作得很好，并且只会执行一次。在 CUDA 中，哪种原子操作是正确的？

Answer 1

在我看来，您想要的是代码中的“关键部分”。临界区允许一个线程执行一系列指令，同时阻止任何其他线程或线程块执行这些指令。

例如，临界区可用于控制对内存区域的访问，从而允许单个线程对该区域的无冲突访问。

原子本身只能用于非常有限的，基本上是单一的操作，对单个变量。但是原子可以用来构建临界区。

您应该在内核中使用以下代码来控制线程对临界区的访问:

__syncthreads();
if (threadIdx.x == 0)
  acquire_semaphore(&sem);
__syncthreads();
  //begin critical section
  // ... your critical section code goes here
  //end critical section
__threadfence(); // not strictly necessary for the lock, but to make any global updates in the critical section visible to other threads in the grid
__syncthreads();
if (threadIdx.x == 0)
  release_semaphore(&sem);
__syncthreads();

在内核定义这些辅助函数和设备变量之前:

__device__ volatile int sem = 0;

__device__ void acquire_semaphore(volatile int *lock){
  while (atomicCAS((int *)lock, 0, 1) != 0);
  }

__device__ void release_semaphore(volatile int *lock){
  *lock = 0;
  __threadfence();
  }

我已经测试并成功使用了上面的代码。请注意，它本质上使用每个线程块中的线程 0 作为请求者在线程块之间进行仲裁。如果您只希望获胜线程块中的一个线程执行临界区代码，则应进一步调整(例如 if (threadIdx.x < ...) )临界区代码。

在一个信号量的 warp 仲裁中拥有多个线程会带来额外的复杂性，所以我不推荐这种方法。相反，让每个线程块仲裁，就像我在这里展示的那样，然后使用普通的线程块通信/同步方法(例如 __syncthreads()、共享内存等)控制您在获胜线程块内的行为

请注意，这种方法的性能成本会很高。当您无法弄清楚如何以其他方式并行化您的算法时，您应该只使用临界区。

最后，警告一下。与任何线程并行架构一样，临界区的不当使用会导致死锁。特别是，对线程块和/或线程块内的扭曲的执行顺序进行假设是一种有缺陷的方法。