有没有办法防止内存地址被访问？

Question

是否可以防止某个内存地址在一段时间内被其他线程访问？例如：

__global__ void func(int* a){
  // other computation
  __lock_address(a);
  a[0] += threadIdx.x;
  __unlock_address(a);
}

第一个完成 other computations 并到达 __lock_address 的线程将锁定该内存地址，直到 _unlock_address 被调用，任何其他到达 __lock_address 的线程都必须等到第一个线程解锁它。上面的例子基本上等同于atomicAdd，但是如果我想做更复杂的计算而不是简单的加法呢？

编辑： mutex in初始化为0，a初始化为-1

__global__ void func(int *a, int *mutex){
  a[0] = atomicCAS(mutex, 0, 1); // a[0] = 1
}

如果我这样做，a[0] 等于 1。但它应该是 0，因为那是互斥量的旧值。

__global__ void func(int *a, int *mutex){
  a[0] = mutex[0]; // a[0] = 0
}

这是完整性检查，a[0] 的值现在为 0。这意味着 mutex 被正确初始化为 0。

Answer 1

您可以使用mutex to protect multithreaded access to the memory region. Cuda Programming Guide has a nice example of using atomic operations to implement it (https://docs.nvidia.com/cuda/cuda-c-programming-guide/#scheduling-example)

__device__ void mutex_lock(unsigned int *mutex) {
    unsigned int ns = 8;
    while (atomicCAS(mutex, 0, 1) == 1) {
        __nanosleep(ns);
        if (ns < 256) {
            ns *= 2;
        }
    }
}

__device__ void mutex_unlock(unsigned int *mutex) {
    atomicExch(mutex, 0);
}

Answer 2

好的，我知道出了什么问题。基本上只有线程块中的第一个线程获得旧值 0，同时将互斥锁设置为 1，其他线程读取互斥锁 after 第一个线程将互斥锁设置为 1，然后陷入死锁. 我发现这个 solution 对我有用。

有没有办法防止内存地址被访问？

Is there a way to prevent a memory address being accessed?

cuda