The CUDA Handbook ▸ Chapter 7. Kernel Execution

7.2 Blocks, Threads, Warps and Lanes

Kernels are launched as grids of blocks of threads. Threads can further be divided into 32-thread warps, and each thread in a warp may be called a lane.⁵

7.2.1 Grids of Blocks

Thread blocks are separately scheduled onto SMs, and threads within a given block are executed by the same SM. Figure 7-1 shows a 2D grid (8W x 6H) of 2D blocks (8W x 8H); Figure 7-2 shows a 3D grid (8W x 6H x 6D) of 3D blocks (8W x 8H x 4D).

Figure 7-1. 2D Grid and Thread Block

Figure 7-2. 3D Grid and Thread Block

Grids have grown over successive hardware generations: SM 1.x allowed up to 65535×65535 blocks, SM 2.x added a third dimension (65535×65535×65535), and beginning with SM 3.x (Kepler) the x dimension was widened dramatically. Current hardware supports grids of up to \(2^{31}-1\) blocks in the x dimension and 65535 in each of the y and z dimensions⁶. Blocks may be up to 512 threads (SM 1.x) or 1024 threads (SM 2.x and later) in size⁷, and threads within a block can communicate via the SM’s shared memory. Blocks within a grid are likely to be assigned to different SMs; to maximize throughput of the hardware, a given SM can run threads and warps from different blocks at the same time. The warp schedulers dispatch instructions as needed resources become available.

Threads

Each thread gets a full complement of registers⁸ and a thread ID that is unique within the threadblock. To obviate the need to pass the size of the grid and threadblock into every kernel, the grid and block size also are available for kernels to read at runtime. The built-in variables used to reference these registers are given in Table 7-1. They are all of type dim3().

Taken together, these variables can be used to compute which part of a problem the thread will operate on. A global index for a thread—unique across the entire grid—can be computed by folding all six thread- and block-coordinate values into a single mixed-radix, Horner-style expression:

int globalThreadId =
    threadIdx.x + blockDim.x*(
    threadIdx.y + blockDim.y*(
    threadIdx.z + blockDim.z*(
    blockIdx.x  + gridDim.x *(
    blockIdx.y  + gridDim.y * blockIdx.z))));

Reading the nesting from the inside out, each coordinate is multiplied by the size of the dimensions enclosing it: the three threadIdx components are weighted by the block dimensions, and the three blockIdx components by the grid dimensions. Equivalently, it is the block’s linear index within the grid times the number of threads per block, plus the thread’s linear index within its block.

In the overwhelmingly common special case of a one-dimensional grid of one-dimensional blocks, every y and z term vanishes and the expression collapses to the one-liner that appears in most CUDA code:
threadIdx.x + blockIdx.x*blockDim.x.

Table 7-1. Built-In Variables

Warps, Lanes, and ILP

The threads themselves are executed together, in SIMD fashion, in units of 32 threads called a warp, after the collection of parallel threads in a loom⁹. (See Figure 7-3.) All 32 threads execute the same instruction, each using its private set of registers to perform the requested operation. In a triumph of mixed metaphor, the ID of a thread within a warp is called its lane.

Figure 7-3. Loom

The warp ID and lane ID can be computed using a global thread ID as follows:

int warpID = globalThreadId >> 5;
int laneID = globalThreadId & 31;

These relations assume that a thread block holds a whole number of warps—that blockDim.x*blockDim.y*blockDim.z is a multiple of 32, as it almost always is. When it is, the block’s contribution to the global thread ID is itself a multiple of 32, so it advances the (now grid-wide) warp ID cleanly while leaving the low five bits—the lane—undisturbed.

One of the reasons why warps are such an important unit of execution is because they are the granularity with which GPUs can cover latency. It has been well-documented how GPUs use thread-level parallelism (TLP) to cover memory latency: it takes hundreds of clock cycles to satisfy a global memory request, so when a memory read is encountered, the GPU issues the request and then schedules other warps until the data arrives. Once the data has arrived, the warp becomes eligible for execution again. The warps in question historically had been running separate instruction streams within threadblocks.

What has been less well-documented is how GPUs also exploit instruction-level parallelism (ILP) – the independent operations within a single thread’s instruction stream – to cover instruction and memory latencies. Where thread-level parallelism hides latency by keeping many warps resident, and so motivates the drive for high occupancy, ILP hides it within a single thread, using the same fine-grained parallelism that CPUs exploit to maximize throughput: when computing (a+b)*(c+d), the additions a+b and c+d can be evaluated in parallel before the multiplication, and the SMs have ample logic to track such dependencies and issue independent instructions within each thread. Loop unrolling is an effective way to unlock ILP: besides slightly reducing the number of instructions per loop iteration, it exposes more independent work for the schedulers to issue before a dependent stall. Because ILP does not require many resident warps, it lets kernels reach comparable or even higher performance at lower occupancy than the same kernel’s highest-occupancy formulation – overturning the occupancy-maximizing wisdom that TLP once made conventional, as Section 7.3 discusses.

The sample code contains examples of ILP in saxpy_unrolled() and the memory benchmarks such as globalRead.cu. The GlobalReads() function template, for example, includes an unroll factor that enables each thread to issue multiple read requests in quick succession; the hardware lets the thread keep running until it actually needs one of the results, by which time the data has begun to arrive.

Object Scopes

The scopes of objects that may be referenced by a kernel grid are summarized in Table 7-2, from the most-local (registers in each thread) to the most-global (global memory and texture objects are per-grid). Before the advent of dynamic parallelism, thread blocks served primarily as a mechanism for inter-thread synchronization within a thread block (via intrinsics such as __syncthreads()) and communication (via shared memory). Dynamic parallelism adds resource management to the mix, since streams and events created within a kernel are only valid for threads within the same thread block.

* In order to execute, a kernel only needs enough local memory to service the maximum number of active warps.

** Streams and events can only be created by CUDA kernels using dynamic parallelism.

Table 7-2. Object Scopes

7.2.2 Execution Guarantees

It is important that programmers never make any assumptions about the order in which blocks or threads will execute. In particular, there is no way to know which block or thread will execute first, so initialization generally should be performed by code outside the kernel invocation.

Execution Guarantees and Inter-Block Synchronization

To coordinate execution at any granularity coarser than a thread block (where threads are guaranteed to be resident within the same SM, so they can communicate via shared memory and synchronize execution using intrinsics such as __syncthreads()), developers may use global memory; but in doing so, they must take care to ensure that the coordinating threads will be active at the same time.

For much of CUDA’s history, it was considered a best practice to launch far more threads than the GPU possibly could execute at the same time. The GPU has dedicated hardware to launch threads, in the form of warps and thread blocks, so quickly that many CUDA kernel formulations simply compute a global ID for the thread, and compute an output element if the ID is within the problem size.

The problem is that unless the GPU is big enough to hold the entire grid, some thread blocks may execute to completion before other thread blocks have started running. The result is deadlock: the threads fruitlessly polling a global memory location are either preventing other threads in the kernel launch from executing, or attempting to coordinate execution with threads that have already exited.

There are a few special cases when inter-block synchronization can work: if simple mutual exclusion is all that’s desired, atomicCAS() certainly can be used to provide that; and thread blocks can use atomics to signal when they’ve completed, so that the last thread block in a grid can perform some operation before it exits, knowing that all other thread blocks have completed execution. This strategy is employed by the threadFenceReduction SDK sample and the reduction4SinglePass.cu sample (Section 12.2).

The systemic answer to guarantee that all threads, warps, and thread blocks of a kernel launch are resident at the same time is with a Cooperative Kernel Launch, as described in Section 7.6.

7.2.3 Block and Thread IDs

A set of special read-only registers give each thread context in the form of a thread ID and block ID. The thread and block IDs are assigned as a CUDA kernel begins execution; for 2D and 3D grids and blocks, they are assigned in row-major order.

Thread block sizes are best specified in multiples of 32, since warps are the smallest possible granularity of execution on the GPU. Figure 7-4 shows how thread IDs are assigned in 32-thread blocks that are 32×1, 16×2, and 8×4, respectively.

Figure 7-4. Blocks of 32 threads

For blocks with a thread count that is not a multiple of 32, some warps are not fully populated with active threads. Figure 7-5 shows thread ID assignments for 28-thread blocks that are 28×1, 14×2, and 7×4; in each case, 4 threads in the 32-thread warp are inactive for the duration of the kernel launch. For any thread block size not divisible by 32, some execution resources are wasted, as some warps will be launched with lanes that are disabled for the duration of the kernel execution.

Figure 7-5. Blocks of 28 threads.

There is no performance benefit to 2D or 3D blocks or grids, but they sometimes make for a better match to the application.

The reportClocks.cu program illustrates how thread IDs are assigned, and how warp-based execution works in general.

__global__ voidWriteClockValues(     unsigned int *completionTimes,     unsigned int *threadIDs ){    size_t globalBlock = blockIdx.x+gridDim.x*        (blockIdx.y+gridDim.y*blockIdx.z);    size_t globalThread = threadIdx.x+blockDim.x*        (threadIdx.y+blockDim.y*threadIdx.z);        size_t totalBlockSize = blockDim.x*blockDim.y*blockDim.z;    size_t globalIndex = globalBlock*totalBlockSize + globalThread;     completionTimes[globalIndex] = clock();    threadIDs[globalIndex] = threadIdx.y<<4|threadIdx.x;}

WriteClockValues() writes to the two output arrays using a global index computed using the block and thread IDs, and the grid and block sizes. One output array receives the return value from the clock() intrinsic, which returns a high-resolution timer value that increments for each warp. In the case of this program, we are using clock() to identify which warp processed a given value. clock() returns the value of a per-multiprocessor clock cycle counter, so we normalize the values by computing the minimum and subtracting it from all clock cycle values. We call the resulting values the thread’s “completion time.”

Let’s take a look at completion times for threads in a pair of 16×8 blocks (Listing 7-3) and compare them to completion times for 14×8 blocks (Listing 7-3). As expected, they are grouped in 32s, corresponding to the warp size.

0.01 ms for 256 threads = 0.03 us/thread

0.01 ms for 256 threads = 0.03 us/threadCompletion times (clocks):Grid (0, 0, 0) - slice 0:   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   aGrid (1, 0, 0) - slice 0:   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6Thread IDs:Grid (0, 0, 0) - slice 0:   0   1   2   3   4   5   6   7   8   9   a   b   c   d   e   f  10  11  12  13  14  15  16  17  18  19  1a  1b  1c  1d  1e  1f  20  21  22  23  24  25  26  27  28  29  2a  2b  2c  2d  2e  2f  30  31  32  33  34  35  36  37  38  39  3a  3b  3c  3d  3e  3f  40  41  42  43  44  45  46  47  48  49  4a  4b  4c  4d  4e  4f  50  51  52  53  54  55  56  57  58  59  5a  5b  5c  5d  5e  5f  60  61  62  63  64  65  66  67  68  69  6a  6b  6c  6d  6e  6f  70  71  72  73  74  75  76  77  78  79  7a  7b  7c  7d  7e  7fGrid (1, 0, 0) - slice 0:   0   1   2   3   4   5   6   7   8   9   a   b   c   d   e   f  10  11  12  13  14  15  16  17  18  19  1a  1b  1c  1d  1e  1f  20  21  22  23  24  25  26  27  28  29  2a  2b  2c  2d  2e  2f  30  31  32  33  34  35  36  37  38  39  3a  3b  3c  3d  3e  3f  40  41  42  43  44  45  46  47  48  49  4a  4b  4c  4d  4e  4f  50  51  52  53  54  55  56  57  58  59  5a  5b  5c  5d  5e  5f  60  61  62  63  64  65  66  67  68  69  6a  6b  6c  6d  6e  6f  70  71  72  73  74  75  76  77  78  79  7a  7b  7c  7d  7e  7f

Completion times (clocks):Grid (0, 0, 0) - slice 0:   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   8   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   a   c   c   c   c   c   c   c   c   c   c   c   c   c   c   c   cGrid (1, 0, 0) - slice 0:   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   0   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   2   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   4   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6   6Thread IDs:Grid (0, 0, 0) - slice 0:   0   1   2   3   4   5   6   7   8   9   a   b   c   d  10  11  12  13  14  15  16  17  18  19  1a  1b  1c  1d  20  21  22  23  24  25  26  27  28  29  2a  2b  2c  2d  30  31  32  33  34  35  36  37  38  39  3a  3b  3c  3d  40  41  42  43  44  45  46  47  48  49  4a  4b  4c  4d  50  51  52  53  54  55  56  57  58  59  5a  5b  5c  5d  60  61  62  63  64  65  66  67  68  69  6a  6b  6c  6d  70  71  72  73  74  75  76  77  78  79  7a  7b  7c  7dGrid (1, 0, 0) - slice 0:   0   1   2   3   4   5   6   7   8   9   a   b   c   d  10  11  12  13  14  15  16  17  18  19  1a  1b  1c  1d  20  21  22  23  24  25  26  27  28  29  2a  2b  2c  2d  30  31  32  33  34  35  36  37  38  39  3a  3b  3c  3d  40  41  42  43  44  45  46  47  48  49  4a  4b  4c  4d  50  51  52  53  54  55  56  57  58  59  5a  5b  5c  5d  60  61  62  63  64  65  66  67  68  69  6a  6b  6c  6d  70  71  72  73  74  75  76  77  78  79  7a  7b  7c  7d

The completion times for the 14×8 blocks, given in Listing 7-4, underscore how the thread IDs map to warps. In the case of the 14×8 blocks, every warp holds only 28 threads; 12.5% of the number of possible thread lanes are idle throughout the kernel’s execution. To avoid this waste, developers always should try to make sure threadblocks contain a multiple of 32 threads.