海报已经找到了自己问题的答案。然而，在下面的代码中，我提供了一个在CUDA中实现“临界区”的通用框架。更详细地说，该代码执行块计数，但很容易修改以托管在“临界区”中执行的其他操作。下面，我还报告了一些对代码的解释，并列出了在CUDA中实现“临界区”时的一些“典型”错误。
代码如下：
#include <stdio.h>
#include "Utilities.cuh"
#define NUMBLOCKS 512
#define NUMTHREADS 512 * 2
struct Lock {
int *d_state;
Lock(void) {
int h_state = 0;
gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int)));
gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice));
}
__host__ __device__ ~Lock(void) {
#if !defined(__CUDACC__)
gpuErrchk(cudaFree(d_state));
#else
#endif
}
__device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); }
__device__ void unlock(void) { atomicExch(d_state, 0); }
};
__global__ void blockCountingKernelNoLock(int *numBlocks) {
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
}
__global__ void blockCountingKernelLock(Lock lock, int *numBlocks) {
if (threadIdx.x == 0) {
lock.lock();
numBlocks[0] = numBlocks[0] + 1;
lock.unlock();
}
}
__global__ void blockCountingKernelDeadlock(Lock lock, int *numBlocks) {
lock.lock();
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
lock.unlock();
}
int main(){
int h_counting, *d_counting;
Lock lock;
gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));
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);
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，可以更改此状态。只有每个块内的单个线程，特别是具有本地线程索引threadIdx.x == 0的线程可以调用lock和unlock方法。
在执行期间的随机时间，具有本地线程索引threadIdx.x == 0和全局线程索引为t的线程之一将首先调用lock方法。特别是，它将启动atomicCAS(d_state, 0, 1)。由于最初d_state == 0，然后d_state将被更新为1，atomicCAS将返回0，线程将退出lock函数，传递到更新指令。同时，所有其他块的所有其他线程，其中threadIdx.x == 0，将执行lock方法。但是，它们将发现d_state的值等于1，因此atomicCAS(d_state, 0, 1)不会执行更新，并将返回1，从而使这些线程运行while循环。在该线程执行上述操作时，线程最终完成更新，然后执行unlock函数，即atomicExch(d_state, 0)，从而将d_state恢复为0。此时，随机地，另一个具有threadIdx.x == 0的线程会再次锁定该状态。
上述代码还包含了第三个内核函数，即
blockCountingKernelDeadlock。然而，这是另一种错误的临界区实现方式，会导致死锁。事实上，我们需要注意的是warp以锁步操作，并在每条指令后进行同步。因此，当我们执行
blockCountingKernelDeadlock时，有可能一个warp中的某个线程（例如一个本地线程索引为
t≠0的线程）会锁定状态。在这种情况下，同一个warp中的其他线程，包括那些
threadIdx.x == 0的线程，都会执行与线程
t相同的while循环语句，因为同一个warp中的线程是以锁步方式执行的。因此，所有线程都将等待有人解锁状态，但没有其他线程能够这样做，代码就会陷入死锁。
locks[id] = 0u;替换atomicExch(&(locks[id]),0u);？（尝试过，但不起作用） - whenov