发帖人已经为自己的问题找到了答案。尽管如此，在下面的代码中，我提供了一个通用框架来实现 CUDA 中的关键部分。更详细地说，代码执行块计数，但它很容易修改以托管要在关键部分执行的其他操作。下面，我还报告了代码的一些解释，以及在 CUDA 中关键部分的实现中的一些“典型”错误。
编码
#include <stdio.h>
#include "Utilities.cuh"
#define NUMBLOCKS 512
#define NUMTHREADS 512 * 2
/***************/
/* LOCK STRUCT */
/***************/
struct Lock {
int *d_state;
// --- Constructor
Lock(void) {
int h_state = 0; // --- Host side lock state initializer
gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int))); // --- Allocate device side lock state
gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); // --- Initialize device side lock state
}
// --- Destructor
__host__ __device__ ~Lock(void) {
#if !defined(__CUDACC__)
gpuErrchk(cudaFree(d_state));
#else
#endif
}
// --- Lock function
__device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); }
// --- Unlock function
__device__ void unlock(void) { atomicExch(d_state, 0); }
};
/*************************************/
/* BLOCK COUNTER KERNEL WITHOUT LOCK */
/*************************************/
__global__ void blockCountingKernelNoLock(int *numBlocks) {
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
}
/**********************************/
/* BLOCK COUNTER KERNEL WITH LOCK */
/**********************************/
__global__ void blockCountingKernelLock(Lock lock, int *numBlocks) {
if (threadIdx.x == 0) {
lock.lock();
numBlocks[0] = numBlocks[0] + 1;
lock.unlock();
}
}
/****************************************/
/* BLOCK COUNTER KERNEL WITH WRONG LOCK */
/****************************************/
__global__ void blockCountingKernelDeadlock(Lock lock, int *numBlocks) {
lock.lock();
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
lock.unlock();
}
/********/
/* MAIN */
/********/
int main(){
int h_counting, *d_counting;
Lock lock;
gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));
// --- Unlocked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelNoLock << <NUMBLOCKS, NUMTHREADS >> >(d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the unlocked case: %i\n", h_counting);
// --- Locked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelLock << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the locked case: %i\n", h_counting);
gpuErrchk(cudaFree(d_counting));
}
代码说明
关键部分是必须由 CUDA 线程按顺序执行的操作序列。
假设构造一个内核，其任务是计算线程网格的线程块数。一种可能的想法是让每个块中的每个线程都threadIdx.x == 0增加一个全局计数器。为了防止竞争条件，所有的增加都必须按顺序发生，因此它们必须合并到临界区中。
上面的代码有两个内核函数：blockCountingKernelNoLock和blockCountingKernelLock. 前者不使用临界区来增加计数器，并且可以看到，返回错误的结果。后者将计数器增加封装在关键部分内，因此产生正确的结果。但是临界区是如何工作的呢？
临界区由全局状态管理d_state。最初，状态是0。此外，两种__device__方法lock和unlock可以改变这种状态。和方法只能由每个块内的单个线程调用，特别是由具有本地线程索引的线程lock调用。unlockthreadIdx.x == 0
在执行期间随机地，具有局部线程索引threadIdx.x == 0和全局线程索引的线程之一t将是第一个调用该lock方法的线程。特别是，它将推出atomicCAS(d_state, 0, 1). 由于最初d_state == 0, thend_state将被更新为1,atomicCAS将返回0并且线程将退出lock函数，传递给更新指令。在这样一个线程执行上述操作的同时，所有其他块的所有其他线程都threadIdx.x == 0将执行该lock方法。然而，它们会找到一个d_state等于的值1，因此atomicCAS(d_state, 0, 1)不会执行更新并返回1，因此让这些线程运行 while 循环。在那个线程之后t完成更新，然后执行unlock功能，即atomicExch(d_state, 0)，从而恢复d_state到0。此时，随机地，另一个线程threadIdx.x == 0将再次锁定该状态。
上面的代码还包含第三个内核函数，即blockCountingKernelDeadlock. 然而，这是另一个错误的临界区实现，导致死锁。事实上，我们记得warp 以同步方式运行，并且它们在每条指令后同步。因此，当我们执行时blockCountingKernelDeadlock，warp 中的一个线程，比如具有本地线程索引的线程t≠0，可能会锁定状态。在这种情况下，同一个 warp 中的其他线程t，包括 withthreadIdx.x == 0的线程，将执行与 thread 相同的 while 循环语句t，即同一个 warp 中的线程的执行步调一致。因此，所有线程都将等待某人解锁状态，但没有其他线程能够这样做，代码将陷入死锁。