CUDA中实现多线程访问公共临界区的问题

CUDA本来是用于多线程并发计算的，除了device、stream、event、thread级别的synchronize外，没有提供现成的锁的API实现，需要借助atomicCAS()自己写个简单实现。另外加锁功能不能用于同一个block里的线程之间，否则程序会卡死(因为同一个block里的线程属于同一个warp的线程在同一个SM上执行时同一时间周期内执行的指令是相同的！)，也就是说要实现多线程并发访问时对公共临界资源的加锁，那么所有的线程必须分布在不同的block里。

假如需要用CUDA准确实现点云的voxelization功能，那么需要在多线程并发将点按各自坐标分配到各自对应的voxel中去同时对已经分配了点的voxel个数进行记数量，但是记数就涉及到了多线程对同一个计数器的访问，所以需要加锁机制，写了一个模拟程序实验了一下，发现加了加锁功能后，计数是正确的，但是整个性能下降严重，在RTX2080卡上跑，假如有30万个点，需要179.54ms左右，通常模型推理使用的比较多的12-15万个点左右也需要73.93-92.44ms左右，如果拿到Jetson边缘端板子上来说显然耗时太长，Orin上稍快点但是仍然不可接受，所以放弃这种加锁实现的方案，采用先作voxelization时点的分配不受限制，分配完后再采用CUDA 多线程机制进行规约化统计点数非零的voxels，然后根据实际voxels个数和要求的MAX_VOXELS的差值，随机将其中一部分的voxel的点数计数清零以便在下一步复制点云数据时丢弃这部分实际非空的voxels，确保voxels个数不超过MAX_VOXELS指定的值(比如说30000)。

CUDA里实现多线程并发访问临界资源时的加锁机制在对性能要求不高或者GPU本身计算能力强悍的场合下可能还是有用的，所以把模拟实现代码分享一下，有需要时稍微修改一下就可以了。

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

#define BLOCK_NUM 300000

struct Lock{
  int *mutex;
  Lock(void){
    int state = 0;
    cudaMalloc((void**) &mutex, sizeof(int));
    cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
  }
  ~Lock(void){
    cudaFree(mutex);
  }
  __device__ void lock(int blkid,int idx){
    //printf("block %d thread #%d locking\n", blkid, idx);
    while(atomicCAS(mutex, 0, 1) != 0);
    //printf("block %d thread #%d locked\n", blkid, idx);
  }
  __device__ void unlock(int blkid, int idx){
    //printf("block %d thread #%d unlocking\n", blkid, idx);
    atomicExch(mutex, 0);
    //printf("block %d thread #%d unlocked\n", blkid, idx);
  }
};


__global__ void theKernel(Lock lock, unsigned int *counter){
  int idx = blockIdx.x * blockDim.x + threadIdx.x; 

  lock.lock(blockIdx.x, idx);
  //printf("Thread #%i enters the critical section.block %d\n", idx, blockIdx.x);
  *counter += 1;
  //printf("couter %d\n", *counter);
  lock.unlock(blockIdx.x, idx);
}

int main(void)
{
  Lock lock;
  unsigned int c = 0;
  unsigned int *counter = NULL;
  cudaMalloc((void**) &counter, sizeof(int));
  cudaMemcpy(counter, &c, sizeof(int), cudaMemcpyHostToDevice);
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start);
  theKernel<<<BLOCK_NUM, 1>>>(lock, counter);
  cudaDeviceSynchronize(); 
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  float ms;
  cudaEventElapsedTime(&ms, start, stop); 
  cudaFree(counter);
  cudaEventDestroy(start);
  cudaEventDestroy(stop);
  printf("time cost: %f\n", ms);
  return 0;
}