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

Question

Is it possible to prevent a memory address being accessed by other threads for some period? for example:

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

the first thread that finished the other computations and reached __lock_address will lock that memory address untill _unlock_address is called, any other threads that reached __lock_address will have to wait until the first thread unlocks it. The above example is basically equivalent to atomicAdd, but what if I want to do more complicated computation rather than a simple addition?

Edit: mutex in initialized to 0, a is initialized to -1

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

if I do this, a[0] is equal to 1. but it should be 0 since that is the old value of mutex.

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

This is a sanity check, value at a[0] is 0 now. which means mutex is initialized to 0 correctly.

score 3 · Accepted Answer · answered May 19 '21 at 15:03

3

You can use 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);
}

answered May 19 '21 at 15:03

Curious

507
3
16

this is exactly what I was looking for! – omer sahban May 19 '21 at 18:13
Hi, I just tried this, and it's causing the kernel to loop infinitely. seems like for some reason atomicCAS is returning the new value of mutex instead of the old value as it should. – omer sahban May 19 '21 at 19:44
I'll edit the question to give you a minimal code. – omer sahban May 19 '21 at 19:45

score 0 · Answer 2 · answered May 19 '21 at 21:37

Ok, I figured out what's wrong. essentially only the first thread in thread block is getting the old value 0, while simultinuously setting mutex to 1, other threads read mutex after mutex is set to 1 by first thread, then stucking in deadlock. I found this solution that worked for me.

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

2 Answers2