使用atomicCAS实现CUDA互斥锁的行为咨询

最近我在学习CUDA编程，尝试用atomicCAS函数实现一个互斥锁，让所有线程依次访问并递增一块共享数据。下面是我写的代码和对应的运行结果，想和大家交流下这个实现的行为是否符合预期，或者有没有可以优化的地方~

实现代码 #

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "book.h"
#include <time.h>

// 全局变量：0表示未锁定，1表示锁定
__device__ int lock;

__global__ void mykernel(float* data){
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    do {
        __syncthreads(); 
        // __threadfence();
        printf("Thread %d of block %d waiting since lock value is %d...\n", threadIdx.x, blockIdx.x, lock);
    } while(atomicCAS(&lock,0,1));

    printf("Locked!\n");
    data[0] += 1;
    printf("Now data is %f\n", data[0]);

    __threadfence(); // 等待写入操作完成
    lock = 0;
    __threadfence(); // 等待写入操作完成
    printf("Unlocked!\n");
}

int main(void) {
    float* d_data;
    float *h_data;
    int n_block = 4;
    int n_thread_per_block = 4;

    // 初始化共享数据
    h_data = (float*) malloc(sizeof(float));
    h_data[0] = 0;

    // 初始化互斥锁
    cudaMemset((int*) lock, 0, sizeof(int));
    printf("Inital value of lock is: %d\n", lock);

    HANDLE_ERROR(cudaMalloc((void**)&d_data, sizeof(float)));
    HANDLE_ERROR(cudaMemcpy(d_data, h_data, sizeof(float), cudaMemcpyHostToDevice));

    mykernel <<<n_block, n_thread_per_block>>> (d_data);
}

运行结果 #

Inital value of lock is: 0
Thread 0 of block 2 waiting since lock value is 0...
Thread 1 of block 2 waiting since lock value is 0...
Thread 2 of block 2 waiting since lock value is 0...
Thread 3 of block 2 waiting since lock value is 0...
Thread 0 of block 1 waiting since lock value is 0...
Thread 1 of block 1 waiting since lock value is 0...
Thread 2 of block 1 waiting since lock value is 0...
Thread 3 of block 1 waiting since lock value is 0...
Thread 0 of block 3 waiting since lock value is 0...
Thread 1 of block 3 waiting since lock value is 0...
Thread 2 of block 3 waiting since lock value is 0...
Thread 3 of block 3 waiting since lock value is 0...
Thread 0 of block 0 waiting since lock value is 0...
Thread 1 of block 0 waiting since lock value is 0...
Thread 2 of block 0 waiting since lock value is 0...
Thread 3 of block 0 waiting since lock value is 0...
Locked!
Thread 0 of block 3 waiting since lock value is 1...
Thread 1 of block 3 waiting since lock value is 1...
Thread 2 of block 3 waiting since lock value is 1...
Thread 3 of block 3 waiting since lock value is 1...
Thread 0 of block 1 waiting since lock value is 1...
Thread 1 of block 1 waiting since lock value is 1...
Thread 2 of block 1 waiting since lock value is 1...
Thread 3 of block 1 waiting since lock value is 1...
Thread 0 of block 0 waiting since lock value is 1...
Thread 1 of block 0 waiting since lock value is 1...
Thread 2 of block 0 waiting since lock value is 1...
Thread 3 of block 0 waiting since lock value is 1...
Now data is 1.000000
Thread 0 of block 3 waiting since lock value is 1...
Thread 1 of block 3 waiting since lock value is 1...
Thread 2 of block 3 waiting since lock value is 1...
Thread 3 of block 3 waiting since lock value is 1...
Thread 0 of block 1 waiting since lock value is 1...
Thread 1 of block 1 waiting since lock value is 1...
Thread 2 of block 1 waiting since lock value is 1...
Thread 3 of block 1 waiting since lock value is 1...
Thread 0 of block 0 waiting since lock value is 1...
Thread 1 of block 0 waiting since lock value is 1...
Thread 2 of block 0 waiting since lock value is 1...
Thread 3 of block 0 waiting since lock value is 1...
Unlocked!
Locked!
Thread 0 of block 3 waiting since lock value is 1...
Thread 1 of block 3 waiting since lock value is 1...
Thread 2 of block 3 waiting since lock value is 1...
Thread 3 of block 3 waiting since lock value is 1...
Thread 0 of block 0 waiting since lock value is 1...
Thread 1 of block 0 waiting since lock value is 1...
Thread 2 of block 0 waiting since lock value is 1...
Thread 3 of block 0 waiting since lock value is 1...
Now data is 2.000000
Thread 1 of block 2 waiting since lock value is 1...
Thread 2 of block 2 waiting since lock value is 1...
Thread 3 of block 2 waiting since lock value is 1...
Thread 0 of block 3 waiting since lock value is 1...
Thread 1 of block 3 waiting since lock value is 1...
Thread 2 of block 3 waiting since lock value is 1...
Thread 3 of block 3 waiting since lock value is 1...
Thread 0 of block 0 waiting since lock value is 1...
Thread 1 of block 0 waiting since lock value is 1...
Thread 2 of block 0 waiting since lock value is 1...
Thread 3 of block 0 waiting since lock value is 1...
Unlocked!
Locked!
Thread 1 of block 2 waiting since lock value is 1...
Thread 2 of block 2 waiting since lock value is 1...

从运行结果来看，线程确实是依次获取锁、修改共享数据、释放锁，数据也能正确递增，整体行为符合我的预期。不过我还是有点疑问：比如__syncthreads在这里的作用是否必要？还有__threadfence的使用是否足够严谨？有没有什么潜在的线程安全问题我没考虑到？

内容来源于stack exchange