CODEWITHSUNDEEP
cudagpgpugpu-warp
Gabriel
Gabriel

Reputation: 9432

How are 2D / 3D CUDA blocks divided into warps?

If I start my kernel with a grid whose blocks have dimensions:

dim3 block_dims(16,16);

How are the grid blocks now split into warps? Do the first two rows of such a block form one warp, or the first two columns, or is this arbitrarily-ordered?

Assume a GPU Compute Capability of 2.0.

Upvotes: 21

Views: 7288

Answers (2)

Mohsen
Mohsen

Reputation: 183

To illustrate @talonmies's answer through 'Visual Studio WarpWatch' window for two consecutive warps (dim3 block_dims(16,16); and WarpSize = 32):

First Warp Second Warp

Upvotes: 6

talonmies
talonmies

Reputation: 72348

Threads are numbered in order within blocks so that threadIdx.x varies the fastest, then threadIdx.y the second fastest varying, and threadIdx.z the slowest varying. This is functionally the same as column major ordering in multidimensional arrays. Warps are sequentially constructed from threads in this ordering. So the calculation for a 2d block is

unsigned int tid = threadIdx.x + threadIdx.y * blockDim.x;
unsigned int warpid = tid / warpSize;

This is covered both in the programming guide and the PTX guide.

Upvotes: 38

Related Questions