初版

流传甚广的版本，来自《GPU高性能编程CUDA实战》中的sample

class CudaKernelLock {
private:
    int* mutex;
public:
    CudaKernelLock(void) {
        int state = 0;
        cudaError_t ret = cudaMalloc((void**)&mutex, sizeof(int));
        ret = cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
    }
    ~CudaKernelLock() {
        cudaFree(mutex);
    }

    __device__ void lock(void) {
        while (atomicCAS(mutex, 0, 1) != 0);
    }
    __device__ void unlock(void) {
        atomicExch(mutex, 0);
    }
};

这个方法思路是可行的，即通过原子比较交换操作atomicCAS竞争mutex，atomicCAS读取mutex的值，计算(*mutex==0?1:*mutex),并将结果存储在原内存地址。这三个操作为一个原子事务中执行。函数返回交换前的*mutex值。这样就在一个线程获取mutex并置值后，其他线程一直在等待。直到atomicExch置mutex为0后可再次获取。故在核函数中lock和unlock函数之间的操作是串行的。

存在问题【踩坑&填坑】

但上述实现有以下几个问题：

该锁的设计只能在CUDA的block之间加锁。

一个block只能有一个thread，即核函数只能设计成kernel<<<128,1>>>(...)而不能是kernel<<<1,128>>>(...)，否则则会死锁。原因是因为CUDA是以warp为单位运行的，而warp的运行遵循同步执行规则（locked-step execution），即一个warp中的线程同时执行一个函数，并同时退出一个函数（SIMT）。一个warp（通常是32个thread）其中一个线程获得锁之后，等待其他线程一起退出lock函数，但其他线程等待该线程执行unlock来释放锁，从而出现了死锁。但当一个block中只有一个thread时，一个warp中只有一个thread，进而不会死锁。

lock和unlock函数之间的非原子操作，并不一定是严格线程安全的。

如下所示为一个计算数组中大于阈值的元素个数的操作。然而并不能得到正确的cnt结果。

__global__ void GtThrCntKernel(float* src, int len, float thr, int* cnt, CudaKernelLock Lock) {
    int pos_start = blockDim.x * blockIdx.x + threadIdx.x;
    int pos_step = blockDim.x * gridDim.x;

    for (int i = pos_start;i < len;i += pos_step) {
        if (src[i] <= thr) {
            continue;
        }
        Lock.lock();
        *cnt = *cnt + 1;
        Lock.unlock();
    }
}
    1. 原因是一个线程对一个全局存储器或共享存储器的修改并不一定对其他block中线程是立即可见的，也就是说可能有些线程仍然是读取的若干次自增前的cnt。
    1. 解决这个问题有三种方法：
      1. 在自增后加个延时函数，但加多长合适，加少了不安全，加多了影响效率，太蠢了；
      1. 将*cnt = *cnt + 1;替换为atomicAdd(cnt, 1);,以确保读写操作都是原子的；
      1. 更加通用的方法是使用__threadfence()，对该命令的官方解释是能确保执行该命令的线程在在该语句前对全局存储器或共享存储器的访问已经全部完成（不是保证所有线程运行到同一位置），即该命令之前的所生产的数据能够安全地被其他线程消费，执行结果对gird中所有线程可见（__threadfence_block()是执行结果对block中所有线程可见）。可以在lock和unlock中加上该命令。

如果采用如下所示的调用方式，该锁设计是不可重入的。

void callKernel(){
    CudaKernelLock lock_;
    // prepare the data
    ...
    // launch kernel
    GtThrCntKernel << <128, 1 >> > (dev_a, len, thr, dev_c, lock_);
    CudaError_t ret = cudaGetLastError();
    if (ret != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(ret));
    }
    // launch kernel again
    GtThrCntKernel << <128, 1 >> > (dev_a, len, thr, dev_c, lock_);
    ret = cudaGetLastError();
    if (ret != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(ret));
    }
}
    1. 首先这里核函数调用时对lock_只能传值，不能传引用或指针，因为lock_对象是host端的量，传引用或指针给核函数，核函数中是没有办法用的。
    1. 传值的话就涉及到拷贝构造，即核函数中生成一个自己的device端的lock对象，并执行默认的浅拷贝mutex指针，浅拷贝本身没有问题，本身也希望mutex只有一份，但问题在于核函数退出时，析构自己的lock对象会把唯一的一个mutex的内存给free掉，这就导致lock_出来后就不可用了，再次将其传入kernel函数时，mutex指针就是个野指针了，此时cudaGetLastError会报错。
    1. 解决办法1：提前将lock对象拷贝到device端，核函数传device端的指针，就避免函数内拷贝构造和析构问题。
__global__ void GtThrCntKernel(float* src, int len, float thr, int* cnt, CudaKernelLock* Lock) {
    int pos_start = blockDim.x * blockIdx.x + threadIdx.x;
    int pos_step = blockDim.x * gridDim.x;

    for (int i = pos_start;i < len;i += pos_step) {
        if (src[i] <= thr) {
            continue;
        }
        Lock->lock();
        atomicAdd(cnt, 1);
        Lock->unlock();
    }
}

void callKernel(){
    CudaKernelLock lock_;
    CudaKernelLock *dev_lock_;
    cudaMalloc((void**)&dev_lock_, sizeof(CudaKernelLock));
    cudaMemcpy(dev_lock_, &lock_, sizeof(CudaKernelLock), cudaMemcpyHostToDevice);
    // prepare the data
    ...
    // launch kernel
    GtThrCntKernel << <128, 1 >> > (dev_a, len, thr, dev_c, dev_lock_);
    CudaError_t ret = cudaGetLastError();
    if (ret != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(ret));
    }
    // launch kernel again
    GtThrCntKernel << <128, 1 >> > (dev_a, len, thr, dev_c, dev_lock_);
    ret = cudaGetLastError();
    if (ret != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(ret));
    }
}
    1. 解决办法2：单独封装Mutex类，Lock类中用shared_ptr管理Mutex，拷贝构造时只会让shared_ptr的use_count加1，不会生成新的Mutex副本，且只有最后一个shared_ptr析构的时候（use_count为0时），才会析构Mutex，可以解决我们的问题。但是遗憾的是，__deivice__函数中不支持shared_ptr重载的->操作符，我们只能在__device__函数外额外定义变量mutex_ptr取得Mutex类对象中的mutex指针，还得自己实现拷贝构造函数(默认的拷贝构造也行)，就有了如下略显丑陋的实现。
class Mutex {
private:
    int* mutex;
public:
    Mutex() {
        int state = 0;
        cudaError_t ret = cudaMalloc((void**)&mutex, sizeof(int));
        ret = cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
    }
    ~Mutex() {
        cudaFree(mutex);
    }
};
class CudaKernelLock {·
public:
    std::shared_ptr<Mutex> mutex = nullptr;
    int* mutex_ptr = nullptr;
public:
    CudaKernelLock(void) {
        mutex = std::make_shared<Mutex>();
        mutex_ptr = mutex->mutex;
    }
    CudaKernelLock(const CudaKernelLock& lock) {
        mutex = lock.mutex;
        mutex_ptr = mutex->mutex;
    }
    ~CudaKernelLock(void) {
    }

    __device__ void lock(void) {
        while (atomicCAS(mutex_ptr, 0, 1) != 0);
        __threadfence();
    }
    __device__ void unlock(void) {
        __threadfence();
        atomicExch(mutex_ptr, 0);
    }
};
    1. 解决办法3：将cudaFree封装成lamda函数，shared_ptr创建时传入销毁函数，优点是不用单独封装Mutex类，缺点还是绕不开需要一个裸指针mutex_ptr去给__device__函数使用。
class CudaKernelLock {
private:
    std::shared_ptr<int> mutex;
    int* mutex_ptr = nullptr;
public:
    CudaKernelLock() {
        int state = 0;
        cudaError_t ret = cudaMalloc((void**)&mutex_ptr, sizeof(int));
        ret = cudaMemcpy(mutex_ptr, &state, sizeof(int), cudaMemcpyHostToDevice);
        // auto lamdaFree = [](int* p) {
        //    cudaFree(p);
        // };
        // std::shared_ptr<int> tmp(mutex_ptr, lamdaFree);
        std::shared_ptr<int> tmp(mutex_ptr, [](int* p) {cudaFree(p);});
        mutex = std::move(tmp);//std::make_shared<int>(dev_mutex, lamdaFree);
    }
    ~CudaKernelLock() {}

    __device__ void lock(void) {
        while (atomicCAS(mutex_ptr, 0, 1) != 0);
        __threadfence();
    }
    __device__ void unlock(void) {
        __threadfence();
        atomicExch(mutex_ptr, 0);
    }
};

真·线程锁设计

好看不一定好用

上述的通用互斥锁类的设计虽然封装的好，但还是解决不了block内thread死锁的问题，因此只能kernel<<<n,1>>>(...)这样去使用，测试下来，虽然比CPU串行遍历计算要快，但还是没有完全发挥CUDA并行计算的能力，毕竟一个warp执行单元中只有一个thread。因此实现了如下直接嵌在核函数里的丑陋但好用的真·线程互斥锁：

__global__ void GtThrCntKernel(float* src, int len, float thr, int* cnt, int* mutex) {
int pos_start = blockDim.x * blockIdx.x + threadIdx.x;
int pos_step = blockDim.x * gridDim.x;

for (int i = pos_start;i < len;i += pos_step) {
    if (src[i] <= thr) {
        continue;
    }
    bool blocked = true;
    while(blocked){
        if (0 == atomicCAS(mutex, 0, 1)){
            // **** critical section ****//
            atomicAdd(cnt, 1);
            __threadfence();
            // **** critical section ****//
            atomicExch(mutex, 0);
            blocked = false;
        }
    }
}
}
void callKernel(){
int mutex_state = 0;
int *mutex_;
CudaError_t ret = cudaMalloc((void **)&mutex_, sizeof(int));
ret = cudaMemcpy(mutex_, &mutex_state, sizeof(int), cudaMemcpyHostToDevice);

// prepare the data
...
// launch kernel
GtThrCntKernel <<<4, 128>>> (dev_a, len, thr, dev_c, mutex_);
ret = cudaGetLastError();
if (ret != cudaSuccess) {
    fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(ret));
}

cudaFree(mutex_);
}
  • 如上所示，思路还是一样的，只不过加/解锁的操作没了函数的封装，直接嵌在核函数里，也就没有了warp必须同进同出函数的问题，信息量mutex仍然需要预先在外部分配好，但是没有了类的封装，自然也不会存在构造/析构的问题。同样临界区中block之间的线程安全问题仍需要原子操作或__threadfence()去解决。
  • 实测下来，不再会出现死锁问题，不会有重入问题，且速度会快不少。