CUDA原子操作修改标志位

Question

CUDA原子操作修改标志位

4

我有一段串行代码，它执行的操作类似于这样：

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

我理解在串行中这个操作可以完美地执行一次。在CUDA中，正确的原子操作是什么？

- ThatQuantDude

为了避免多个线程同时修改同一变量，从而导致未定义的行为。 - sgarizvi

sgar91，是的，正确的，这是旧的遗留代码，我无法改变它的结构。因此，基本上第一个被执行的线程应该执行它，将其阻塞在warp中的其他线程，并将变量更改为true，以便没有其他线程会再进入该部分。 - ThatQuantDude

2

你所描述的并不像是原子函数，而更像是一个临界区。你可以在右上角搜索“cuda临界区”以获取一些想法。不幸的是，我的一篇关于cuda临界区的帖子被删除了。如果你愿意，我可以在这里发布它作为答案。“原子”函数只允许在“做一些初始化”区域中对通常单个变量进行有限的操作。因此，如果该区域涉及任何操作，它可能无法使用原子服务，尽管原子有助于构建临界区。 - Robert Crovella

嗨，罗伯特，你的帖子将会受到很高的赞赏。 - ThatQuantDude

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

在我看来，您需要的是代码中的“关键区域(critical section)”。关键区域允许一个线程执行一系列指令，同时防止任何其他线程或线程块执行这些指令。

例如，可以使用关键区域控制对内存区域的访问，以便单个线程无冲突地访问该区域。

原子操作本身只能用于非常有限的、基本上是单个变量的操作。但是可以使用原子操作来构建关键区域。

您应该在内核中使用以下代码来控制线程对关键区域的访问：

__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()、共享内存等）控制在获胜的线程块中的行为。

请注意，这种方法将对性能产生显著影响。仅当无法找到其他并行算法时才应使用关键部分。

最后，警告一句。与任何线程并行架构一样，不当使用关键部分可能导致死锁。特别地，假设线程块和/或warp内的执行顺序是一种有缺陷的方法。

这里是使用binary_semaphore实现单个设备全局“锁定”的示例，可用于访问控制到关键部分。