19

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.

einpoklum
  • 118,144
  • 57
  • 340
  • 684
Gabriel
  • 8,990
  • 6
  • 57
  • 101

2 Answers2

37

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.

talonmies
  • 70,661
  • 34
  • 192
  • 269
  • 12
    Note "column-major order" assumes that the dim3 is an array, rather than a struct. A more precise description is that `.x` is the fastest varying of the dimensions, `.y` is second-fastest varying, and `.z` varies slowest. How you associate `.x`, `.y`, and `.z` with rows, columns, slices, offsets, tree levels, or any other addressing within memory is up to you. – harrism Sep 24 '12 at 23:18
4

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

Mohsen
  • 153
  • 11