我有一段序列代码,它做这样的事情
if( ! variable )
{
do some initialization here
variable = true;
}
我知道这在串行中非常有效,并且只会执行一次。什么原子操作在 CUDA 中是正确的?
在我看来,您想要的是代码中的“关键部分”。临界区允许一个线程执行一系列指令,同时阻止任何其他线程或线程块执行这些指令。
临界区可用于控制对内存区域的访问,例如,以允许单个线程对该区域的非冲突访问。
原子本身只能用于非常有限的,基本上是单一的操作,对单个变量。但是原子可以用来构建一个临界区。
您应该在内核中使用以下代码来控制对关键部分的线程访问:
__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 arbitrate中为一个信号量拥有多个线程会带来额外的复杂性,所以我不推荐这种方法。相反,让每个线程块像我在这里展示的那样进行仲裁,然后使用普通线程块通信/同步方法(例如__syncthreads()
,共享内存等)控制您在获胜线程块中的行为
请注意,这种方法对性能的影响很大。当您无法弄清楚如何并行化您的算法时,您应该只使用关键部分。
最后,一句警告。与任何线程并行架构一样,对关键部分的不当使用可能会导致死锁。特别是,对线程块内的线程块和/或扭曲的执行顺序进行假设是一种有缺陷的方法。