在CUDA中实现关键部分

Question

我正在尝试使用原子指令在CUDA中实现一个关键部分,但我遇到了一些麻烦.我创建了测试程序来显示问题:

#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <stdio.h>

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    while (atomicExch(&(locks[id]), 1u) != 0u) {} //lock
    //critical section would go here
    atomicExch(&(locks[id]),0u); //unlock
}

int main(int argc, char** argv) {
    //initialize the locks array on the GPU to (0...0)
    unsigned int* locks;
    unsigned int zeros[10]; for (int i = 0; i < 10; i++) {zeros[i] = 0u;}
    cutilSafeCall(cudaMalloc((void**)&locks, sizeof(unsigned int)*10));
    cutilSafeCall(cudaMemcpy(locks, zeros, sizeof(unsigned int)*10, cudaMemcpyHostToDevice));

    //Run the kernel:
    k_testLocking<<<dim3(1), dim3(256)>>>(locks, 10);

    //Check the error messages:
    cudaError_t error = cudaGetLastError();
    cutilSafeCall(cudaFree(locks));
    if (cudaSuccess != error) {
        printf("error 1: CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
        exit(-1);
    }
    return 0;
}

不幸的是,这段代码很难冻结我的机器几秒钟,最后退出,打印出来的消息:

fcudaSafeCall() Runtime API error in file <XXX.cu>, line XXX : the launch timed out and was terminated.

这意味着其中一个while循环没有返回,但似乎这应该工作.

作为提醒,atomicExch(unsigned int* address, unsigned int val)以原子方式设置存储在地址中的内存位置val的old值并返回该值.所以我的锁定机制背后的想法是它最初是0u,因此一个线程应该通过while循环,所有其他线程应该等待while循环,因为它们将读取locks[id]为1u.然后当线程完成临界区时,它会重置锁,0u以便另一个线程可以进入.

我错过了什么？

顺便说一下,我正在编译:

nvcc -arch sm_11 -Ipath/to/cuda/C/common/inc XXX.cu

Answer 1

好吧,我把它想出来了,这是另一种 - 最重要的 - 范例 - 痛苦.

正如任何优秀的cuda程序员都知道的那样(注意我不记得这使我成为一个糟糕的cuda程序员,我认为)warp中的所有线程必须执行相同的代码.如果不是因为这个事实,我写的代码将完美地工作.但是,实际上,同一个warp中可能有两个线程访问同一个锁.如果其中一个获取了锁,它只会忘记执行循环,但它不能继续循环,直到其warp中的所有其他线程都完成循环.不幸的是,其他线程永远不会完成,因为它正在等待第一个线程解锁.

这是一个内核,可以毫无错误地执行操作:

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    bool leaveLoop = false;
    while (!leaveLoop) {
        if (atomicExch(&(locks[id]), 1u) == 0u) {
            //critical section
            leaveLoop = true;
            atomicExch(&(locks[id]),0u);
        }
    } 
}

Answer 2

发布者已经找到了解决自己问题的答案。不过，在下面的代码中，我提供了一个通用框架来实现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和unlock方法只能由每个块内的单个线程和，特别是通过具有局部线程索引的线程的调用，threadIdx.x == 0。

在执行期间，具有本地线程索引threadIdx.x == 0和全局线程索引的线程之一将随机t地首先调用该lock方法。特别是它将启动atomicCAS(d_state, 0, 1)。从最初开始d_state == 0，然后d_state将被更新为1，atomicCAS将返回0，并且线程将退出lock函数，并传递给update指令。同时，这样的线程执行上述操作，所有其他块中的所有其他线程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。但是，这是关键部分的另一种错误实现，导致死锁。的确，我们记得扭曲在锁步中运行，并且它们在每条指令后都会同步。因此，当我们执行时blockCountingKernelDeadlock，线程束中的一个线程（例如具有本地线程索引的线程t?0）可能会锁定状态。在这种情况下，处于相同warp中的其他线程t（包括带有with的threadIdx.x == 0线程）将执行与thread相同的while循环语句t，即以锁步执行的处于相同warp中的线程的执行。因此，所有线程将等待某人解锁状态，但是其他线程将无法这样做，并且代码将陷入死锁。