Threads, blocks, and warps: map work without assuming an order
Trace CUDA indices, partial warps, block independence, and the difference between launched threads and useful elements.
On this page
The question for today
What you’ll understand
- Compute a one-dimensional CUDA element index and launch size.
- Explain block independence and NVIDIA’s 32-thread warp grouping.
- Place bounds checks without violating required block-wide participation.
A GPU kernel describes work executed by many logical threads, but those threads are organized into groups with different cooperation rules. This module follows NVIDIA CUDA Toolkit 12.6.0 documentation and explicitly labels architecture-specific details and hypothetical capacities. NVIDIA CUDA warps contain 32 threads in this model; that number is not a promise about every vendor’s GPU. Start with the logical work mapping, then reason about which threads can communicate and which execution orders the program must tolerate.
The launch defines logical identities
A one-dimensional grid contains blocks, and each block contains threads. The common element index multiplies blockIdx.x by blockDim.x and adds threadIdx.x. Launch enough blocks to cover the input and guard accesses whose index lies beyond the valid range. Integer overflow in host-side size arithmetic can corrupt the launch calculation, so use an appropriate size type and validate device and kernel limits. A valid launch configuration does not by itself prove that its memory addresses are in bounds.
For multidimensional blocks, CUDA linearizes thread coordinates with the x coordinate changing fastest. Warps are formed from consecutive linear thread IDs inside a block; they do not combine leftover threads from different blocks. A block size that is not a multiple of 32 therefore includes a partially populated final warp. Logical threads are not permanently dedicated physical processors. The implementation schedules warp work on available streaming-multiprocessor resources according to architecture and kernel constraints.
A block is a cooperation boundary
Threads in a block can use shared memory and supported block-level synchronization. Ordinary independent blocks must be able to run in any order, concurrently or sequentially; assuming block zero finishes before block one is incorrect. A block is placed on one SM for the ordinary execution model, and several blocks may reside there if resources permit. Specialized cooperative launches and newer cluster features add explicitly documented mechanisms, but they do not make an ordinary launch an implicit grid-wide barrier.
A bounds check must be designed around cooperation. In an elementwise kernel with no collective operations, a thread can skip its out-of-range work. In a tiled algorithm with a block barrier, skipping the barrier in only some non-exited participating threads can break synchronization. Often the right pattern is for every required thread to load a valid value or neutral padding, participate in the collective, and guard only final global stores. The algorithm’s synchronization contract determines whether an early exit is actually safe.
Separate coverage, residency, and readiness
Grid coverage asks whether every intended element has a logical thread or loop iteration. Residency asks how many blocks and warps can be live on an SM under register, shared-memory, and architecture limits. Readiness asks which resident warps have operands and dependencies ready to execute. Those are different quantities. Launching many threads does not guarantee they are simultaneously resident, and a fully populated warp can still spend time waiting for memory or executing work irrelevant to the application.
The modular lab illustrates a limited index-cycle idea. Reset, set modulus m=8 and additive step a=2, predict residues 0,2,4,6 and then zero, and Step four times. This shows that a stride can visit only part of a finite index space. It does not simulate CUDA threads, block placement, warps, or memory accesses. Use the explicit launch arithmetic below for coverage; the lab is a reminder to prove the mapping rather than assume that repeated arithmetic reaches every item.
Work it through
Cover a nonmultiple input
A one-dimensional elementwise kernel processes N=1,000 values with B=256 threads per block. It has no block collectives and guards every data access with i<N.
Compute ceil(1,000/256)=4 blocks. They launch 4×256=1,024 threads, so 24 launched threads must not access input or output elements.
Each block has 256/32=8 warps, giving 32 warp groups across the grid. These groups need not all be resident or execute at the same instant.
In the final block, local thread 231 maps to 3×256+231=999 and is valid. Local thread 232 maps to 1,000 and is the first out-of-range thread; its guarded data operations are skipped.
The launch covers all 1,000 elements while suppressing 24 invalid accesses. Its arithmetic says nothing about simultaneous residency or elapsed time.
The living notebook
Explore the model
Interactive lab · modular
Prove an index cycle reaches what you expect
Reset, set m=8 and a=2, predict 0→2→4→6→0, then Step four times. This additive cycle illustrates finite index coverage; it is not a CUDA thread, warp, memory, or scheduling simulator.
- Residue r0
- 0
- Orbit period · m/gcd(m,a)
- 12
- Distinct residues visited
- 1 of 12
Addition modulo m starts at zero: rₙ = na mod m. Playback advances two additions per second. The orbit has period m/gcd(m,a), including period one for a = 0. This is not a cryptographic security test or a primality algorithm. Parameter changes pause playback and recompute the model at the current elapsed time. Reset restores both parameters and returns time to zero. Playback starts paused. Adjust either parameter, then choose Play or Step.
Look forRepeated index arithmetic needs a coverage argument before it becomes a valid parallel work mapping.
Pause & reason
Check your understanding
Your turn
Try it yourself
An elementwise kernel uses 128 threads per block for 513 values. Compute blocks, launched threads, and out-of-range threads. State why a collective kernel needs more reasoning.
Take a moment to make a prediction or work it out on paper. Open a hint when you need one.
Give me a hint
Round up the block count, then separate valid work from launched lanes.
Compare with the solution
Five blocks launch 640 threads, leaving 640−513=127 out-of-range threads. An elementwise guard can suppress their accesses. If the algorithm has block collectives, required participants must still obey the collective’s contract; the arithmetic alone does not justify placing a barrier inside a nonuniform bounds branch.
Keep the idea
Take it with you
- Logical launch coverage and physical execution are separate questions.
- NVIDIA warp grouping is block-local and uses 32-thread groups here.
- Bounds handling must preserve the algorithm’s required synchronization participation.
Read further
Original sources and deeper explanations to follow at your own pace.
Optional progress, saved in this browser when available. Revisit any lesson at any time.
Lesson not marked complete.