Bug in emulation mode and __syncthreads()? Kernel stops abruptly

Hi all

I’m having a problem with running my kernel in emulation mode. My kernel is as follows:

[codebox]while (condition)

perform some operations

__syncthreads

if (condition)

while (condition)

some other operations

__syncthreads

rest of the kernel[/codebox]

The problem is that if the second __syncthreads is ever reached, the kernel terminates instantly without the rest of the code being run. Is this a bug in emulation mode or is there something I have misunderstood? The same thing happens if I insert a __syncthreads immediately after the first while.

If I run it outside emulation mode, the kernel executes properly.

Any ideas?

Can the second ‘condition’ be different for different threads? And are you absolutely sure that your non-emulation runs really are correct?

You should insert

cudaThreadSynchronize();

printf("%s\n", cudaGetErrorString(cudaGetLastError()));

after your kernel call and examine the output.