# \_\_syncthreads() + shared memory issue **URL:** \ **Category:** CUDA Programming and Performance\ **Created:** [August 20, 2008, 10:30am UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095 "2008-08-20T10:30:25Z")\ **Posts on this page:** 8\ **Page:** 1
**Author:** ![NitrousUK](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@NitrousUK](https://forums.developer.nvidia.com/u/NitrousUK)\ **Post date:** [August 20, 2008, 10:30am UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/1 "2008-08-20T10:30:25Z")
I’ve got an issue I’ve been banging my head against for a couple of days, and I’m reaching the point of looking to CUDA to be at fault. Reason being I have done a printf of each threads ID, plus the value of a shared variable, directly after a \_\_syncthreads() call, and one thread has a different value from the others. How can this be possible if \_\_syncthreads is supposed to update all threads with the latest value in shared memory?? I’ve also tried making the shared boolean value ‘volatile’ too, but without success. Also note the incorrect usage of \_\_syncthreads error is because thread 9 is entering a branch with a \_\_syncthreads() because its dataToLoad value is 1. Heres the code, and its output: CODE: **shared** bool dataToLoad; \_\_syncthreads(); printf(“Thread:%2d dataLoad:%d\n”,threadIdx.x,dataToLoad); OUTPUT: Thread: 0 dataLoad:0 Thread: 1 dataLoad:0 Thread: 2 dataLoad:0 Thread: 3 dataLoad:0 Thread: 4 dataLoad:0 Thread: 5 dataLoad:0 Thread: 6 dataLoad:0 Thread: 7 dataLoad:0 Thread: 8 dataLoad:0 Thread: 9 dataLoad:1 Error:incorrect use of \_\_syncthreads() From nvidia programming guide, p.21 : “Only after the execution of a \_\_syncthreads() (Section 4.4.2) are writes to shared variables guaranteed to be visible by other threads.” As far as i can see, this should be impossible, I don’t see how any code or errors on my behalf could cause the above output. A shared variable should have the same value amongst all threads after a \_\_syncthreads() call… ps. I can’t paste the whole code since this is work for a company, top secret hush hush stuff ;) ---
**Author:** ![Sarnath](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/sarnath/32/9742_2.png) [@Sarnath](https://forums.developer.nvidia.com/u/Sarnath)\ **Post date:** [August 20, 2008, 11:25am UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/2 "2008-08-20T11:25:14Z")
Device emulation does not model the underlying hardware correctly! Shared memory variables need to be initialized! Operations on Un-initialized variables are just undefined. ---
**Author:** ![NitrousUK](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@NitrousUK](https://forums.developer.nvidia.com/u/NitrousUK)\ **Post date:** [August 20, 2008, 12:51pm UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/3 "2008-08-20T12:51:37Z")
Cheers for the input Sarnath, but I’ve initialised the variable. It’s not until well into the running of the kernel the issue arises. I Don’t suppose if anyone knows of documentation beyond the few sentences in the programming guide for the \_\_syncthreads() function? I’m wondering if threads sync with specific calls to the \_\_syncthreads() function, or perhaps they simply block until all threads hit a call, rather than a specific one? It seems to be whenever I set this shared boolean to true (which causes the ‘if’ following the code I pasted to then be true) i get the cuda error on the first following occurence of \_\_syncthreads(). I wonder if maybe the cuda compiler is erroneously thinking that since that variable is liked to a potentially divergent if, that threads will diverge and miss \_\_syncthread calls? Thought as far as I can see, if \_\_syncthreads does what its supposed to, no threads should diverge. ---
**Author:** ![redpill](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@redpill](https://forums.developer.nvidia.com/u/redpill)\ **Post date:** [August 21, 2008, 8:25pm UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/4 "2008-08-21T20:25:31Z")
Sorry if this is too basic, but: Is one thread executing \_\_syncthreads() more than the others? That would be bad. Also, can you reproduce this behavior OUTSIDE the emulator? ---
**Author:** ![tmurray](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tmurray](https://forums.developer.nvidia.com/u/tmurray)\ **Post date:** [August 21, 2008, 8:41pm UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/5 "2008-08-21T20:41:56Z")
“Also note the incorrect usage of \_\_syncthreads error is because thread 9 is entering a branch with a \_\_syncthreads() because its dataToLoad value is 1.” Yeah, you can’t have \_\_syncthreads() in divergent branches. ---
**Author:** ![smokyboy](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@smokyboy](https://forums.developer.nvidia.com/u/smokyboy)\ **Post date:** [August 23, 2008, 2:47pm UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/6 "2008-08-23T14:47:55Z")
I believe I ran into a similar problem of incorrect thread synchronization in emulation mode. I am using a shared array as well. My code goes like this: ```auto   Â extern __shared__ char array[];   Â int tx = threadIdx.x, ty = threadIdx.y;   Â int offset = SIZE * blockIdx.x + tx + DIM * ty;   Â array[offset] = 3;   Â __syncthreads(); ``` When I run this program in debug+emulation mode, all threads except one seem to block on syncthreads. The rebel thread just comes through somehow, so the program deadlocks. The program runs fine on the GPU. I haven’t really tried if the shared memory is to blame, but what worries me that the debugging is not possible (or can I debug kernels without emulation?). I am using ubuntu, cuda 2.0beta, and netbeans 6.1. ---
**Author:** ![smokyboy](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@smokyboy](https://forums.developer.nvidia.com/u/smokyboy)\ **Post date:** [August 23, 2008, 2:58pm UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/7 "2008-08-23T14:58:40Z")
I tried removing the shared memory write in my code above and the program runs fine. As soon as I uncomment it, it deadlocks again. So this must be some emulation bug, right? ---
**Author:** ![NitrousUK](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@NitrousUK](https://forums.developer.nvidia.com/u/NitrousUK)\ **Post date:** [August 26, 2008, 3:09pm UTC](https://forums.developer.nvidia.com/t/syncthreads-shared-memory-issue/5095/8 "2008-08-26T15:09:23Z")
“Yeah, you can’t have \_\_syncthreads() in divergent branches.” - tmurray Yeah, this is true, however my code should be synchonous as all threads’ condition branchs operate on the same shared memory, so they should all evalute to true/false the same, if \_\_syncthreads is updating the values obtained from shared memory as it should be. smokyboy, I think your describing the same problem I’m having. I’ve managed to get around it by moving the shared boolean to outside the branch. I still havent figured out why it deadlocks. I’m starting to wonder if theres some data structure describing the dependencies between variables and conditional logic between threads running in the background and its erroneously flagging a deadlock situation. I’ve tried putting a \_\_syncthreads before/after every line that operates on shared memory, so I dont see how user error could be causing a deadlock if every operation should be synchonrised. The only way there could be a deadlock then is if I were using the threadIdx variable in an if, which I’m not. I’m hoping nvidia had the foresight to sync up individual \_\_syncthreads calls, and not simply have a counter, otherwise if one were to call one more \_\_syncthreads than the others that would explain why it falls through a call when others hold. But if they were prudent and ID’d each \_\_syncthreads call, it should flag the error upon the next call to \_\_syncthreads. Would be nice if Nvidia gave a bit more documentation about the guts of CUDA…