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

Solution 1:

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.

Solution 2:

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