barrier deadlock on __syncthreads?

[GoogleCodeExporter]

What steps will reproduce the problem?

The error below occurs when encountering the first __syncthreads(); in my 
CUDA kernel code.

What is the expected output? What do you see instead?

==Ocelot== Emulator failed to run kernel "_Z18chiSquaredDistancePfS_S_i" 
with exception:
==Ocelot== [PC 91] [thread 0] [cta 0] bar.sync 0 - barrier deadlock at: 
precomputeMatrix_chikernel.cu:53:0
terminate called after throwing an instance of 
'executive::RuntimeException'
Aborted

What version of the product are you using? On what operating system?

OpenSuse 11.1/GCC4.3.2/CUDA Toolkit 2.3/Ocelot SVN r271 (2009-12-22)

Please provide any additional information below.

__global__ void
chiSquaredDistance(float* C, float* A, float* B, int slabSizeA)
{
    // Thread index
    int tx = threadIdx.x;

    // Block index
    int bx = blockIdx.x;
    int by = blockIdx.y;

    float temp = 0.0f;

    for(int vectorChunk = 0; vectorChunk < VECTORCHUNK_COUNT; vectorChunk+
+)
    {
        temp += 1.0f;  // compute some value for temp, not removed
    }

    __shared__ float sum[VECTORCHUNK_SIZE];
    sum[tx] = temp;
    __syncthreads();
    for(int bit = VECTORCHUNK_SIZE / 2; bit > 0; bit /= 2)
    {
        float t = sum[tx] + sum[tx ^ bit];
        __syncthreads();
        sum[tx] = t;
        __syncthreads();
    }

    // write to global memory
    if(tx == 0)
        C[by * slabSizeA + bx] = sum[tx] / 2;
}

Original issue reported on code.google.com by koenvand...@gmail.com on 15 Jan 2010 at 2:04

[GoogleCodeExporter]

This error is generated when a syncthreads operation is encountered by some, 
but not
all of the threads in a CTA.  I took a look at the assembly generated by your 
code
and it was not obvious where this might occur. 

If you could send me a complete test program, I could try debugging it further 
to see
if this is a problem with the ocelot emulator.

Original comment by gregory....@gatech.edu on 17 Jan 2010 at 8:51

[GoogleCodeExporter]

Fixed when updating to the lastest SVN revision (r309) instead of using r272.

Original comment by koenvand...@gmail.com on 19 Jan 2010 at 12:05

[GoogleCodeExporter]

Original comment by gregory....@gatech.edu on 3 Feb 2010 at 10:34

• Changed state: Fixed