5

I have something like this:

__global__ void globFunction(int *arr, int N) {
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;
    // calculating and Writing results to arr ...
    __syncthreads();
    // reading values of another threads(ex i+1)
    int val = arr[idx+1]; // IT IS GIVING OLD VALUE
}


int main() {
    // declare array, alloc memory, copy memory, etc.
    globFunction<<< 4000, 256>>>(arr, N); 
    // do something ...
    return 0;
}

Why am I getting the old value when I read arr[idx+1]? I called __syncthreads, so I expect to see the updated value. What did I do wrong? Am I reading a cache or what?

3
  • 4
    __syncthreads is a block level synchronization primitive, not grid level. Based on your last few questions, it seems you should spend a bit of time reading the CUDA documentation.
    – talonmies
    Commented Nov 22, 2011 at 22:17
  • Ok so it is impossible to read it?
    – nosbor
    Commented Nov 22, 2011 at 22:20
  • @sidyll: It's C for CUDA. Leave the C tag.
    – Kerrek SB
    Commented Nov 22, 2011 at 22:28

1 Answer 1

7

Using the __syncthreads() function only synchronizes the threads in the current block. In this case this would be the 256 threads per block you created when you launched the kernel. So in your given array, for each index value that crosses over into another block of threads, you'll end up reading a value from global memory that is not synchronized with respect to the threads in the current block.

One thing you can do to circumvent this issue is create shared thread-local storage using the __shared__ CUDA directive that allows the threads in your blocks to share information among themselves, but prevents threads from other blocks accessing the memory allocated for the current block. Once your calculation within the block is complete (and you can use __syncthreads() for this task), you can then copy back into the globally accessible memory the values in the shared block-level storage.

Your kernel could look something like:

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

If you must synchronize threads across blocks, you should be looking for another way to solve your problem, since the CUDA programing model works most effectively when a problem can be broken down into blocks, and threads synchronization only needs to take place within a block.

4
  • 5
    Warning, this is dangerous code. The compiler may choose to reorder the writes to flags and arr, resulting in a race condition. You should probably need a __threadfence() in there to be correct. And in general, interblock communication without atomics must be done with care, if at all -- usually you should find another way...
    – harrism
    Commented Nov 23, 2011 at 7:05
  • Thanks for the info on __threadfence. I've added that to the code, and according to the Nvidia documentation that should work properly. Since only a single thread is writing to a flag (there is a flag per thread), I don't see the need for atomic operations.
    – Jason
    Commented Nov 24, 2011 at 5:26
  • 1
    @Jason: Your second code assumes that the thread idx is scheduled and running at the same time as the thread idx+1. Nothing in the CUDA execution model guarantees that. Sorry, but that code is broken and it will deadlock under easily demonstrable conditions.
    – talonmies
    Commented Nov 24, 2011 at 6:38
  • Okay, I did not realize that about the CUDA runtime. I've changed my answer to reflect this. I'd appreciate any further comments if you see something wrong.
    – Jason
    Commented Nov 24, 2011 at 15:37

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.