When do we need two dimension threads in CUDA?

Question

I was wondering when we should use x and y coordinates for threads in CUDA? I've seen some codes when they have nested loops, they use x and y coordinates. Is there any general rules for that?Thanks

Answer 1

The answer to the question in the title is simple: Never. You never really need the 2D coordinates.

However, there are several reasons why they are actually present. One of the main reason is that it simplifies the modelling of certain problems. Particularly, of problems that GPUs are "good at", or that they have been used for, for "historical" reasons. I'm thinking about things like image processing or matrix operations here. Writing an image processing or matrix multiplication CUDA kernel is far more intuitive when you can clearly say:

int pixelX = threadIdx.x + blockIdx.x * blockDim.x;
int pixelY = threadIdx.y + blockIdx.y * blockDim.y;

and from that on only deal with the simple pixel coordinates. How much this actually simplifies the index hassle becomes even more obvious when shared memory is involved, for example, during a matrix multiplication, and you want to slice-and-dice a set of rows+columns out of a larger matrix, to copy it to local memory. If you only had 1D indices and had to fiddle around with offsets and strides, this would be error prone.

The fact that CUDA actually does not only support 2D, but also 3D kernels might stem from the fact that 3D textures are frequently used for things like Volume Rendering, which is also something that can be greatly accelerated with GPUs (Websearches including keywords like "volume ray casting" will lead you to some nice demos here).

(Side note: In OpenCL, this idea has even been generalized. While CUDA only supports 1D, 2D and 3D kernels, in OpenCL, you only have ND kernels, where the N is explicitly given as the work_dim parameter)

_{(Another side note: I'm pretty sure that there also are more low-level, technical reasons, that are related to hardware achitectures of GPUs or the caching of video memory, where the localities of 2D kernels may easily be exploited and be beneficial for the overall performance - but I'm not familiar with that, so this is only a guess until now)}