Does early exiting a thread disrupt synchronizatio

Question

This question already has an answer here:

• Can I use __syncthreads() after having dropped threads? 2 answers

I am implementing a certain image processing algorithm with CUDA and I have some questions about the thread synchronization issue overall.

The problem at hand can be explained like that:

We have an image with the size of W*H. For each pixel of the image I need to run 9 identical-data parallel processes and each process gives an array of values as the result (the arrays are of the same length for the whole algorithm, lets say N, which is around 20 or 30). For each pixel, these 9 processes will accumulate their results in a final array (a single array for each pixel) after they finish their calculations.

In order to parallelise this, I designed the following structure: I generate blocks with the dimensions of (10,10,9), this means each thread block will process a 10*10 sized sub image and each thread will process 1 of the 9 identical processes for a single pixel. The grid dimension will be (W/10,H/10,1) in this case. For a thread block, I will allocate a shared memory array with the length of 100*N and each thread will write to the appropriate shared memory location according to the coordinates of its current pixel. So, I need a synchronization with an atomicAdd and __synchthreads() here.

The problem here is, if a pixel has a value of zero, then we don't need to process it at all, so I want to exit for such pixels, otherwise I will do unnecessary work since a large portion of the image consists of zeroes (background). So, I thought of writing something like the following:

//X and Y are the coordinates of the current pixel in the input image.
//threadIdx.z gives the index of the process among the 9 for the current pixel. 

int X=blockIdx.x * blockDim.x + threadIdx.x;
int Y=blockIdx.y * blockDim.y + threadIdx.y;
int numOfProcessForTheCurrPixel=threadIdx.z;
int linearIndexOfPixelInBlock=threadIdx.y * blockDim.x + threadIdx.x;

unsigned short pixelValue=tex2D(image,X,Y);
//Here, threads processing zero-pixels will exit immediately.
if(pixelValue==0)
 return;

float resultArray[22];
//Fill the result array according to our algorithm, mostly irrelevant stuff.
ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

for(int i=0;i<22;i++)
    atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

 __syncthreads(); 
 //Then copy from the shared to the global memory and etc. 

What worries me in this situation is what the Programming Guide is saying:

__syncthreads() is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects.

So in my case, if some of the pixels in a 10*10 thread block are zero and some or not, then the threads belonging to the zero pixels will exit immediately at the beginning and the other threads will continue to their processing. What about the synchronization in this case, will it work still properly or will it generate undefined behavior like the Programming Guide says? I thought of making zero pixel threads processing garbage data to keep them busy but this will unnecesarilly increase the processing time if we have blocks which consist of zeroes entirely (and we have them very often). What can be done in this case?

Answer 1

To avoid creating a deadlock, all threads need to hit the _synchthreads() unconditionally. You can do that in your example by replacing the return with an if statement that jumps over the bulk of the function and heads straight for the _syncthreads() for the zero pixel case.

unsigned short pixelValue=tex2D(image,X,Y);
//If there's nothing to compute, jump over all the computation stuff
if(pixelValue!=0)
{

    float resultArray[22];
    //Fill the result array according to our algorithm, mostly irrelevant stuff.
    ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

    for(int i=0;i<22;i++)
        atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

}

__syncthreads(); 

if (pixelValue != 0)
{
    //Then copy from the shared to the global memory and etc. 
}