what 'incorrect use of __syncthreads()' means ?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi
What the error message ‘incorrect use of __syncthreads()’ means in reality ?

I have kernel that runs fine in hardware, but when run in device emulation mode
some blocks terminate early with ‘incorrect use of __syncthreads()’ message,

In what situations this message is triggered ?

I’m suspicious about this construct:

int device foobar()
{
shared someval;

if (threadIdx.x == 0)
{
// calculate someval
}

__syncthreads();

// early out from function
if (someval == 0)
return 0;

int reg_someval = someval;

for (i = 0; i < reg_someval; i++)
{
for (j = 0; …)
{
}

   __syncthreads();

  if (threadIdx.x == 0)
  {
      // merge values on thread 0 and write result to shared memory
  }

  __syncthreads();

}

return result_from_shared_memory;
}

void global fooKernel()
{
int val = foobar();
if (val > 0)
{
// do something on all threads
if (threadIdx.x == 0)
{
// merge values from shared memory and write result to global memory
}
__syncthreads();
}

// do some other calculations

return;
}

how the ‘__syncthreads’ are realised in emulation mode ?
is the code above abuses ‘__syncthreads’ ? - all should be fine since all threads in
the block always follow that same execution path (they all take or not the branch).

2

__syncthreads() works only if ALL threads of the same block reach it,
ex:

if (val > 0)
{
// do something on all threads
if (threadIdx.x == 0)
{
// merge values from shared memory and write result to global memory
}
__syncthreads();
}

it’s good only if val > 0 for the whole block.

3

and since val is from shared memory it is that same for whole block :), not a case here

interestingly if i’v put OutputDebugString around every __syncthreads in my kernel to see witch one triggers the message … suprise … the message is gone and all blocks finish now corectly in device emulation mode … what is going on ?

4

uuuh… well, I don’t know,

try to stop this annoying alert putting __syncthreads() out of the if(),

and if it’s needed you can replicate the if() statement:

if (val > 0)

{

…

}

__syncthreads();

if (val > 0)

{

…

}

…

5

that’s will only mask the problem, not kill its root ;)
so again question to nvidia ppl ‘when’ this message about __syncthreads is triggered ?