How to make the threads idle? Histogramming application

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi guys,

I will try to be as specific as possible. I need to obtain a color histogram of an image. To do so, I have a window of 15x15 (constant, cannot change it). In other words, every thread in my application is going to move that 15x15 = 225 pixels in the original image to the shared memory, and do the color histogramming in the shared memory. I also do some r,g,b to H,S,V color conversion in the shared memory. Due to the size of the image, it seems I need 71x47 threads. 71 in the horizontal and 47 in vertical. So I set the block size to 16x16. Using this block size, I have to have a grid dimension of 5x5 blocks, which resuls in 80x80 threads and that is larger than 71x47. Finally, here is my question:

How can I make the excessive threads idle? For example, would this be a legitimate code if I put this at the beginning of my kernel:

if (blockID.x + threadID.x > 71) && (blockID.y + threadID.y > 47)

{

      return;

}

This basically says, if the thread number in the horizontal or vertical direction is bigger than some number, just return and do not do anything.

The reason that I am asking is somewhere in code makes my GPU hang. I thought that I might have a programming model flaw, which I am suspicious about this part in my code. Also, if you think there is a problem in my coding logic (block or grid sizes,inefficient computing) would you let me know?

Thanks in advance for the reply.

2

Yes you can exit any thread any time.

There is one exception: The early-out may not skip any __syncthreads() call as the other threads will then wait forever for the one that has already exited. This also applies to any other flow control structure. So if you have a sync inside an if() statement, you need to be sure that either all threads take the if-branch or all take the else-branch. Same for for-loops, all need to do same number of iterations if there is a sync inside.

Peter