__syncthreads(); doesn't work more than 32 threads in a block. Shared memory: make sure the shared

Accelerated Computing CUDA CUDA Programming and Performance
1

I have a shared memory variable in the Kernel named “backgNoise”, and the value of backgNoise can only be gotten in 1st thread. How can I promise other threads use it after it update the value?

dim3 block_dimension(10, 10);
thread_dimension = 60;

kernel<<<block_dimension, thread_dimension >>>(dataArray_device, valueTauCMM_device);

kernel{
shared float backgNoise;

      if(threadIdx.x == 0) {
               backgNoise = a;
      }
      b[threadIdx.x] = b[threadIdx.x] - backgNoise;

}

I’ve tried syncthreads. In these two ways:
1.
if(threadIdx.x == 0) {
backgNoise = a;
__syncthreads();
}

if(threadIdx.x == 0) {
backgNoise = a;

}
__syncthreads();

But they just work when block has less than 32 threads. When I use more than 32 threads,it just fail.

2

Use syncthreads after the if statement

3

Hi, thanks.
Yes, I’ve tried syncthreads. In these two ways:
1.
if(threadIdx.x == 0) {
backgNoise = a;
__syncthreads();
}

if(threadIdx.x == 0) {
backgNoise = a;

}
__syncthreads();

But they just works when block has less than 32 threads. When I use more than 32 threads,it just fail.

4

Rule of thumb is to never put __syncthreads into conditional branch - all threads should be making this call.

5

Hello,

The function __synchthreads must be executed by all threads in order to obtaine the synchronization. Other wise you can use the threadfence function which makes sure that result of a specific line is seen by all threads before. The threadfence function can be used in if without problems.

6

Thanks everybody.