CUDA __syncthreads(); not working; inverse in breakpoint hit order

Question

I have a problem, I think with __syncthreads();, in the following code:

__device__ void prefixSumJoin(const bool *g_idata, int *g_odata, int n)
{
    __shared__ int temp[Config::bfr*Config::bfr];  // allocated on invocation  
    int thid = threadIdx.y*blockDim.x + threadIdx.x;  
    if(thid<(n>>1))
    {
        int offset = 1;
        temp[2*thid] = (g_idata[2*thid]?1:0); // load input into shared memory  
        temp[2*thid+1] = (g_idata[2*thid+1]?1:0); 
        for (int d = n>>1; d > 0; d >>= 1)                    // build sum in place up the tree  
        {   
            __syncthreads();  
            if (thid < d)  
            { 
                int ai = offset*(2*thid+1)-1; // <-- breakpoint B 
                int bi = offset*(2*thid+2)-1;
                temp[bi] += temp[ai];  
            }  
            offset *= 2; 
        } 
        if (thid == 0) { temp[n - 1] = 0; } // clear the last element

        for (int d = 1; d < n; d *= 2) // traverse down tree & build scan  
        {  
            offset >>= 1;  
            __syncthreads();  
            if (thid < d)                       
            {
                int ai = offset*(2*thid+1)-1;  
                int bi = offset*(2*thid+2)-1;
                int t = temp[ai];  
                temp[ai] = temp[bi];  
                temp[bi] += t;   
            }  
        }  
        __syncthreads();
        g_odata[2*thid] = temp[2*thid]; // write results to device memory  
        g_odata[2*thid+1] = temp[2*thid+1]; 
    }
}


__global__ void selectKernel3(...)
{
    int tidx = threadIdx.x;
    int tidy = threadIdx.y;
    int bidx = blockIdx.x;
    int bidy = blockIdx.y;
    int tid = tidy*blockDim.x + tidx;
    int bid = bidy*gridDim.x+bidx;
    int noOfRows1 = ...;
    int noOfRows2 = ...;

    __shared__ bool isRecordSelected[Config::bfr*Config::bfr];
    __shared__ int selectedRecordsOffset[Config::bfr*Config::bfr];

    isRecordSelected[tid] = false;
    selectedRecordsOffset[tid] = 0;


    __syncthreads();
    if(tidx<noOfRows1 && tidy<noOfRows2)
        if(... == ...)
            isRecordSelected[tid] = true;

    __syncthreads();
    prefixSumJoin(isRecordSelected,selectedRecordsOffset,Config::bfr*Config::bfr); // <-- breakpoint A
    __syncthreads();

    if(isRecordSelected[tid]==true){
        {
            some_instruction;// <-- breakpoint C
        ...
        }
    }
}
...
f(){
   dim3 dimGrid(13, 5);
   dim3 dimBlock(Config::bfr, Config::bfr);

   selectKernel3<<<dimGrid, dimBlock>>>(...)

}
//other file

class Config
{
public:
    static const int bfr = 16; // blocking factor = number of rows per block
public:
    Config(void);
    ~Config(void);
};

The prefixSum is from GPU Gems 3: Parallel Prefix Sum (Scan) with CUDA, with little change.

Ok, now I set 3 breakpoints: A, B, C. They should be hit in the order A, B, C. The problem is that they are hit in the order: A, B*x, C, B. So at point C, selectedRecordsOffset is not ready and it causes errors. After A the B is hit a few times, but not all and then C is hit and it goes further in the code and then B again for the rest of the loop. x is different depending on the input (for some inputs there isn't any inversion in the breakpoints so C is the last that was hit).

Moreover if I look on thread numbers that cause a hit it is for A and C threadIdx.y = 0 and for B threadIdx.y = 10. How is this possible while it is the same block so why some threads omit sync? There is no conditional sync. Does anyone have any ideas on where to look for bugs?

If you need some more clarification, just ask.
Thanks in advance for any advice on how to work this out.
Adam

Answer 1

Thou shalt not use __syncthreads() in conditional code if the condition does not evaluate uniformly across all threads of each block.