GPU Architecture and CUDA Programming
Follow an output owner. Trace its addresses. Reuse values before moving them again.
On this page
A tensor operation becomes threads and data movement
A tensor expression states the desired result. A kernel must also decide which thread computes each value, where operands live, and when another thread can use them. Fast arithmetic alone does not answer those questions.
Start with one output element. Give it an owner. Follow that owner’s loads, arithmetic and store. Then inspect neighboring threads: can their accesses share memory transactions, operands or partial results?
This note uses CUDA C++ and NVIDIA terminology. The interactive models are original, synthetic explanations. Their integer values, address sets and schedules are exact under their stated assumptions. Their playback speed is not GPU execution time.
Separate the machine from the programming model
The host is the CPU side of the application. It allocates or manages storage, prepares inputs and submits work. The device executes kernels. CUDA submission is often asynchronous: the host can continue while previously submitted device work remains unfinished.
| Programming concept | Hardware relationship | Important boundary |
|---|---|---|
| Grid | All blocks launched by one kernel invocation | No implicit whole-grid barrier inside an ordinary kernel |
| Thread block / CTA | A cooperating group scheduled onto one streaming multiprocessor (SM) | Block threads can share storage and use block barriers |
| Warp | 32 threads from one block, with consecutive linear thread IDs | Lane identity describes position in a warp |
| SM | Execution resources, warp scheduling, registers and shared-memory capacity | An SM can host several resident blocks when resources permit |
An SM contains several execution resources; a “CUDA core” is not an independent CPU running one complete CUDA thread. Warp scheduling issues instructions for eligible work. Some instructions wait on data or dependencies; other ready work can use the available resources.
Blocks are designed to execute independently. Do not rely on block 3 starting before block 4, or on the same launch receiving the same SM assignment. Thread-block clusters add specific cooperation features; they do not make all blocks globally synchronized. CUDA programming model ↗.
Follow one thread from launch to output
For a one-dimensional launch, CUDA exposes a block index, the number of threads per block, and a thread index within that block. Flattening those coordinates gives one global output index.
Warp 1 · 32 lanes
Warp 1 · 32 lanes
Warp 1 · 32 lanes
Block 2 → warp 0 → lane 0 → index 128. A[128]=128 + B[128]=256 gives C[128]=384.
The diagram uses two synthetic SMs with two block slots each. Its numbered waves make allocation visible; they are not a CUDA scheduling guarantee. Selecting a block changes the dispatch wave. Selecting a warp exposes its 32 lanes.
With N=130 and 64 threads per block, three blocks launch 192 threads. Only indices 0–129 own valid outputs. Block 2, warp 0, lanes 0 and 1 are valid; the remaining lanes are masked. A 32-thread warp still exists even when many lanes perform no useful arithmetic.
In multidimensional blocks, the x thread coordinate varies fastest when forming the linear thread ID. A convenient two-dimensional drawing does not change which consecutive threads form a warp. Address layout still determines memory behavior. SIMT kernels and indexing ↗.
Write a small kernel with a complete ownership rule
The simplest ownership rule assigns one output to one thread. Each valid thread reads the corresponding input pair and writes its own result. No thread reads another thread’s intermediate value, so this kernel needs no block barrier.
__global__ void add(const float* a, const float* b,
float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) c[i] = a[i] + b[i];
}Here a, b and c refer to valid device storage containing at least n floats. Assume n is a positive int whose indexing fits the expression. The example uses separate output storage. Memory allocation, input transfer and cleanup belong to the calling program.
int threads = 256;
int blocks = n / threads + (n % threads != 0);
add<<<blocks, threads>>>(a, b, c, n);
check(cudaGetLastError()); // check launch submission
check(cudaDeviceSynchronize()); // check execution completionThe illustrative check helper must report errors and stop the failing operation; it is not a CUDA API. A successful launch submission does not prove successful completion. The final synchronization is useful for a correctness check, but adding it after every production kernel can prevent overlap.
Divergence means lanes take different control-flow paths. It can reduce useful parallel work, but a bounds guard is necessary here. Removing the guard to keep lanes uniform would introduce invalid memory accesses. Safe tail handling comes before a performance preference. CUDA C++ kernel launch basics ↗.
Memory spaces define visibility and reuse
| Storage | Who can use it? | Why it matters |
|---|---|---|
| Registers | A thread’s working values | Fast local arithmetic; large live state can limit residency or spill |
| Shared memory | Threads in one block; explicitly supported cluster access is a separate contract | Program-controlled reuse and communication |
| L1 / L2 caches | Hardware-managed according to architecture and access policy | Reduce some accesses to more distant storage; not an explicit sharing protocol |
| Global memory | Addressable device memory available across kernels and blocks | Large capacity, but traffic and dependencies can dominate |
| Local memory | Thread-private address space backed by device memory, possibly cached | The name does not imply on-chip register storage |
A variable written in source code is not guaranteed to remain in a physical register. Dynamic indexing, register pressure or compiler decisions can place thread-private state in local memory. A large per-thread array therefore needs inspection, not an assumption of cheap storage.
On H100, the SM register file contains 64K 32-bit registers: 256 KiB shared among resident execution contexts. The architectural per-thread register limit is 255 registers. This is not 256 KiB per thread. Shared-memory capacity and allocation limits are separate resources. Hopper resource limits ↗.
Distributed shared memory exposes storage within a supported thread-block cluster. Its synchronization and lifetime rules remain necessary. It is not permission for any SM on the GPU to read arbitrary live shared-memory allocations.
Follow addresses into memory sectors
A warp’s global-memory instruction produces a set of lane addresses. The memory system combines accesses into transactions. Count the aligned regions those addresses touch; do not infer traffic from the number of source-level loads alone.
4 sectors × 32 B = 128 B requested. 128 B of unique values used · 100% sector utilization.
Lane 0 requests float 0, byte offset 0, in sector 0. Each small sector cell is one float. Reverse lane order preserves the aligned sector set.
The model fixes 32 active lanes, one aligned 4-byte float per lane, and a 32-byte-aligned base address. Aligned consecutive floats cover four sectors. Shifting the base by one float touches five. A stride of eight floats touches 32 sectors. Reversing the aligned lane order still touches the same four sectors.
“Consecutive lane numbers” and “nearby addresses” are different properties. The selected connection lets you follow one lane while the sector map shows the whole warp’s footprint. Each blue word cell is used; empty cells belong to a requested sector but are unused by this instruction.
The receipt counts sector requests for this simplified access pattern. It is not a measurement of HBM bytes or latency. Cache hits, repeated accesses, instruction width, alignment and the architecture can change lower-level traffic. Two equal sector counts need not produce identical elapsed time. Global-memory access patterns ↗.
Shared memory is a banked service
Shared memory is explicitly managed, but not infinitely multiported. For the scalar 32-bit read model below, consecutive words map to 32 successive banks. Different word addresses in one bank require separate service; the same word can be broadcast to several readers.
32 logical service rounds. Showing 0/32. 1 banks receive requests.
No requests have been served yet. Different words in one bank queue separately. Readers of one identical word share a request.
A row with word stride 1 distributes requests across all banks. A column of a 32-wide row-major array has word stride 32, so every lane addresses bank 0 at a different word. Padding the row width to 33 changes that column’s bank mapping to one bank per lane.
Padding is a layout change, not a universal optimization. It consumes storage and changes address calculations. Swizzling is another layout transformation that can distribute accesses while meeting a compute instruction’s operand layout.
The broadcast preset is the important exception: 32 readers of one identical word do not create a 32-way read conflict. The queue rack groups those readers together. Its service rounds visualize distinct word requests; they are not measured cycles or a model of all shared-memory instruction widths. Shared-memory banks and broadcast ↗.
A barrier protects a producer–consumer boundary
Shared storage creates a producer–consumer dependency. A thread can finish its own load while a neighboring thread has not yet produced a value it needs. A block barrier makes cooperating threads wait and provides the required shared-memory ordering for this pattern.
The first barrier protects consumers from incomplete loads. The second protects prior readers from an early overwrite by threads loading the next tile. A buffer can be ready for consumption without being free for reuse.
Keep block barriers on a control-flow path reached consistently by the block. In a tiled kernel, out-of-range threads can write zeros into their assigned shared slots and still participate. Returning early from those threads can break the block’s cooperation contract.
Warp communication has a different scope. Shuffle instructions exchange register values within a warp. The participation mask and executing lanes must satisfy the instruction’s contract. Independent thread scheduling means old assumptions of implicit lockstep are not a replacement for the required synchronization. Advanced kernel synchronization ↗.
Matrix multiplication is a reuse problem
Matrix multiplication exposes the cost of missing reuse. Let A have shape M×K and B have shape K×N. An output element takes a dot product of one A row and one B column.
A naive output owner loads both operands for each k. Neighboring outputs need many of the same operands. A cooperative block can load an A tile and a B tile once into shared memory, then reuse them for several output accumulators.
Stage 0/4. Select the top-left output tile. No operands have moved.
4 threads · 1 output accumulator per thread.
For the full 4×4 product: 64 global float loads and 128 shared operand reads.
The full-product counts exclude 16 final C stores and ignore caches. Larger register tiles also enlarge this example’s block tile. T labels identify output owners; the scene does not claim a measured speedup.
The walkthrough uses A entries 1–16 and a fixed small B matrix. It reveals two k positions at a time, then the remaining two. C[0,0] first contains 1, then reaches 11. Other output tiles are dimmed; they require their own owners.
For a full FP32 product, a naive requested-byte count includes two float loads per multiply-accumulate and one final float store per output. For a square T×T output tile, assume cooperative input loading with no redundant loads within that block:
A 16×16 tile therefore approaches 4 FLOPs per requested byte for large K, rather than a quantity growing as 4K. Requested load reduction explains one opportunity; measured bandwidth and execution dependencies still determine time. Original iterative FP32 GEMM worklog ↗.
Give a thread more outputs to reuse its operands
Shared tiling reuses values across threads. Register tiling reuses values inside a thread. If one thread owns a 2×2 output region, two A values and two B values can update four accumulators for one k position.
In the walkthrough, the shared strategy uses four threads for a 2×2 output block. The register strategy keeps four threads but gives each a 2×2 region, forming a 4×4 block. The larger block increases global reuse, while per-thread operand reuse reduces shared reads. These are two linked changes, not one isolated optimization.
| Strategy | Global float loads | Shared operand reads | Accumulators per thread |
|---|---|---|---|
| Naive | 128 | 0 | 1 |
| 2×2 shared output tiles | 64 | 128 | 1 |
| 4×4 block; 2×2 register regions | 32 | 64 | 4 |
All strategies perform the same 64 multiply-accumulates and produce the same 16 final values. The counts exclude 16 final output stores and ignore caching. More register state can reduce resident work or spill. A lower load count is a hypothesis for a faster kernel, not proof.
Vectorized accesses can reduce instruction count when alignment and layout permit. A float4 access requires a valid aligned 16-byte region. Zeroing the entire vector at a partial tail loses valid elements unless a separate scalar or masked tail handles them. Inspect actual generated instructions and test odd dimensions. CUDA memory optimization guidance ↗.
Reuse also depends on the route through output tiles
A and B operands can also be reused between output tiles. Two C tiles in the same output row need the same A row tile. Two C tiles in the same output column need the same B column tile. Their order of execution can therefore affect cache reuse.
Visit 1/16: C tile (0,0). So far: 2 operand-tile loads, 0 cache hits. Full traversal: 20 loads.
Each C tile requests A[row] then B[column]. The unified, fully associative LRU cache starts empty; every operand tile has equal size. This is a controlled-order reuse model, not a prediction of GPU cache misses or block scheduling.
A row-major traversal keeps returning to one A row while moving through B columns. A Hilbert traversal stays near recently visited output tiles in both dimensions. The result matrix is unchanged. Whether either order wins depends on cache capacity, tile size, competing work and the execution schedule.
An ordinary CUDA grid does not promise a chosen block execution order. Persistent kernels make a bounded set of blocks process multiple tiles through an explicit work-assignment loop. Such a design can control tile traversal more directly, but also needs load balancing, tail handling and an appropriate resource budget.
Gordić connects persistent scheduling, cache-aware tile traversal and cluster reuse in his Hopper discussion. The cache experiment isolates the reuse mechanism and does not reproduce his referenced benchmark. Gordić: asynchronous Hopper kernels ↗.
Resources limit residency; residency is not the objective
A block size of 128 threads can be a useful tuning choice, particularly for a four-warp WGMMA group. It is not a general CUDA minimum: several smaller resident blocks can collectively supply eligible warps to an SM.
Occupancy is the fraction of the maximum supported warps that are resident on an SM. Resident does not mean issuing useful arithmetic every cycle. A resident warp can wait for dependencies, and active lanes can be masked.
The block’s threads, register allocation, shared-memory allocation and architectural block limits constrain how many blocks can reside together. Real allocations have granularity and architecture-specific rules. Use the occupancy APIs or profiler rather than dividing nominal capacities and treating the result as exact.
| Change | Potential benefit | Potential cost |
|---|---|---|
| More output accumulators | More operand reuse and instruction-level parallelism | More register pressure |
| Larger shared tile | Fewer repeated global loads | More shared storage and fewer resident blocks |
| More resident warps | More independent work to hide latency | Does not remove bandwidth or dependency limits |
| Fusion | Avoid a global-memory intermediate | More live values, scheduling work or spills |
Tile quantization leaves masked work at matrix edges. A 33×33 output covered by 32×32 tiles needs four tiles: 4096 output slots for 1089 valid elements. Wave quantization describes a different tail: a final group of blocks may leave many SMs without work. Neither effect predicts elapsed time without assumptions about block duration and scheduling.
A kernel with lower occupancy can finish sooner because it moves fewer bytes. A kernel with high occupancy can still waste transactions. Compare useful completion time, correctness and the expected mechanism. Profiling: interpret occupancy and performance budgets own the measurement framework.
A reduction changes the ownership of partial results
A dot product and a softmax both need reductions. The ownership rule changes as partial results combine: several input owners produce fewer partial-result owners, until one complete result remains.
Stage 0/5 · 32 active partial sums. Each tree level combines two previously disjoint sets of contributions.
This exact 32-value tree sums integers 1 through 32. Each level combines disjoint contributor sets. The faded suffix keeps historical values for context; only the active prefix holds the current partial sums.
// All 32 lanes of one warp execute every shuffle.
float sum = value;
int lane = threadIdx.x % 32;
for (int offset = 16; offset > 0; offset /= 2) {
float other = __shfl_down_sync(0xffffffffu, sum, offset);
if (lane < offset) sum += other;
}
// Lane 0 now owns this warp's sum.This snippet assumes one-dimensional indexing and full-warp participation. A partial warp or divergent calling pattern needs a correct mask and reduction strategy; replacing the mask blindly is unsafe. Floating-point sums can change with reduction order even when the mathematical sum is identical.
For a block reduction, warp leaders can store partials in shared memory. A block barrier precedes reading those partials. The final reduction must include only initialized entries or neutral values. A fixed 32-entry reduction is invalid when fewer than 32 warp partials have been initialized.
A grid-wide reduction needs another supported coordination method, such as another kernel, atomics under a valid protocol, or a cooperative launch with its constraints. For communication across GPUs, continue with Collective Communication. For production building blocks, CUB ↗ provides established device and block primitives.
Overlap only work with independent storage
Shared tiling still waits for each tile’s operands. A pipeline can load the next tile while computing the current tile. The two operations need independent resources, and the next load needs storage that no current consumer still owns.
Serial finishes at 20; two buffers finish at 14. Cursor t=10. A: tile 2 being read · B: tile 3 loading.
A buffer stays owned from the start of its load until its final compute read finishes. Copy and compute use separate ideal resources. No setup cost, contention or GPU performance measurement is modeled.
Each pale buffer-ownership band extends from the load’s start to the end of the corresponding compute read. When compute is slower than loading, a loaded tile can wait. Loading the tile two positions ahead must wait for that same slot’s previous consumer to finish.
Ampere introduced asynchronous global-to-shared copies that can avoid intermediate register staging. Hopper adds the Tensor Memory Accelerator (TMA) for structured bulk transfers. A producer can initiate a transfer, but consumers still need the specified completion and visibility protocol.
A copy instruction returning is not permission to read its destination. A ready notification is not permission to overwrite a buffer still being consumed. The exact barrier, transaction-count and async-proxy rules depend on the instruction family. CUDA asynchronous data copies ↗ and Hopper TMA ↗.
XOR changes storage addresses without changing the matrix
Padding changes a row’s pitch. A swizzle instead permutes positions within a storage layout. Consumers must apply the same logical-to-physical address mapping; the mathematical matrix does not change.
The highlighted logical column is unchanged. Row r moves chunk c to physical chunk c XOR r. The eight readers track the first 32-bit word of each highlighted chunk. The left addresses share a bank; the right addresses use distinct banks. This illustrates one aligned swizzle region, not an entire Tensor Core instruction.
// One 1024-byte-aligned region, with offsets from 0 through 1023.
unsigned swizzled = offset ^ ((offset & 0x380u) >> 3);For eight rows of 128 bytes, divide each row into eight 16-byte chunks. This mapping keeps the row and the low four byte bits fixed. It XORs the three chunk-column bits with three row bits. Applying it twice restores the original offset.
The illustration tracks eight logical readers, one per row, rather than claiming to simulate every lane or transaction of a real Tensor Core instruction. Each 16-byte chunk spans four 32-bit banks. The first word of logical column 3 maps to bank 12 in the linear layout and to eight distinct banks after swizzling.
Real TMA tensor maps require the documented swizzle mode, alignment, shape and descriptor constraints. Arbitrary XOR addressing is not automatically a valid TMA layout. Gordić provides the motivating chunk diagrams; the NVIDIA data-copy guide defines the supported layout contracts. Gordić: swizzle derivation ↗; TMA swizzle rules ↗.
Tensor Cores change the compute contract
Tensor Cores accelerate supported matrix multiply-accumulate operations. They require suitable element types, shapes and operand layouts. They are not arbitrary scalar arithmetic units that can replace every CUDA instruction.
The cooperation scope depends on the instruction. Warp-level matrix instructions involve one warp. Hopper WGMMA instructions involve a warpgroup: four contiguous warps, aligned as specified by the ISA. “Tensor Cores always need four warps” is therefore too broad.
Operand descriptors and layouts must match the instruction’s requirements. Shared-memory swizzles can help supply operands without the conflicts of a naive layout. Asynchronous matrix operations also require their own completion rules before the program consumes results or reuses operand storage.
Treat architecture generations separately. A Hopper TMA/WGMMA design is not a portable recipe for every NVIDIA GPU. Compiler or library support does not remove the need to validate the chosen dtype, accumulation behavior and tolerance. PTX matrix instruction contracts ↗; Precision and Quantization.
Read the instructions that answer your question
CUDA C++ is a source language. PTX is a virtual instruction set. The toolchain translates work into native instructions for the selected GPU. PTX virtual registers are not a direct count of allocated physical registers.
nvcc -O3 -lineinfo -Xptxas=-v -arch=sm_90 kernel.cu -o kernel
cuobjdump --dump-sass kernelThese illustrative commands target Hopper’s sm_90. Choose a target supported by the installed toolkit and the GPU that will run the program. Retain compiler options with benchmark results.
Inspect a concrete hypothesis: did a vectorized access appear; did registers spill; did a loop unroll; did an unexpected conversion remain? A shorter PTX listing alone does not establish lower native instruction cost or faster completion.
The compiler’s resource report, native instructions and profiler answer different questions. Read them together with the measured kernel. NVCC compilation model ↗; CUDA binary utilities ↗.
Validate correctness before interpreting speed
Start with a reference result and awkward shapes. Vector lengths need not divide the block size. Matrix dimensions need not divide the tile size. Padding must contribute neutral values, rather than read unrelated storage or remove valid tail elements.
| Check | Example failure it catches |
|---|---|
| Compare values with a trusted reference | Wrong index, missed term or incorrect rescaling |
| Test boundaries and nonmultiples | Out-of-bounds vector tail or incomplete matrix tile |
| Define floating-point tolerances | Different reduction order mistaken for an indexing defect |
| Check every CUDA API and launch result | Invalid launch or allocation failure hidden by later work |
| Use memory/race/synchronization checking tools | Invalid access or inconsistent cooperation |
| Inspect register/shared resource use | Unexpected spills or an infeasible configuration |
| Measure representative work after warmup | Initialization or compilation mistaken for kernel cost |
compute-sanitizer --tool memcheck ./kernel
compute-sanitizer --tool racecheck ./kernel
compute-sanitizer --tool synccheck ./kernelThe tools expose particular classes of failures; a clean run is not a proof for every input or scheduling interleaving. Compute Sanitizer ↗.
Measure asynchronous GPU work with a completion boundary that matches the question. CUDA events can time a declared stream interval. Host wall time around an unsynchronized launch mainly measures submission. End-to-end timing should include the application work that matters, rather than only the most convenient kernel.
Warm up, repeat, retain hardware/software/dtype/shape details, and check numerics after each change. Nsight Systems helps identify scheduling and dependency gaps; Nsight Compute helps investigate a selected kernel’s resource and instruction behavior. Profiling and Kernel Optimization provides the complete measurement workflow.
Connect the mechanism to AI workloads
Matrix multiplication supplies the compute; reductions and data lifetime supply the rest. In attention, explicitly storing a large score matrix can create costly global-memory writes and reads. Tiling and online normalization allow partial results to be combined without materializing every intermediate in device memory.
A running softmax state must preserve the maximum, normalization sum and weighted output consistently. When a larger maximum arrives, old contributions must be rescaled. Uninitialized warp partials, incorrect masks and all-masked rows can break an otherwise plausible implementation. Attention owns the exact equations and masking rules.
The same questions recur in LayerNorm, softmax, activation fusion and quantized kernels: who owns the reduction, which values can stay on chip, and what must remain live? Specialized instructions are useful only when their constraints match that dataflow.
These kernels operate within a GPU. FSDP, tensor parallelism and collective communication add ownership across GPUs. Distinguish the two levels before interpreting a distributed performance trace. Distributed Training and Parallelism; GPU Communication.
Sources, scope and next investigations
Two complementary starting points were Gordić: Inside NVIDIA GPUs ↗ and Making CUDA less overwhelming ↗. Their visual progressions suggested the questions; the explanations and all interactive figures here are original. No blog benchmark numbers are reproduced as new measurements.
- CUDA programming model ↗ and Writing SIMT kernels ↗ — execution, address mappings and memory cooperation.
- CUDA Best Practices Guide ↗ — layouts, reuse and performance tradeoffs.
- Hopper tuning guide ↗ — H100-specific resource limits, clusters and TMA.
- PTX ISA reference ↗ — instruction participation, layouts and synchronization contracts.
- Boehm’s FP32 GEMM worklog ↗ — original measured optimization sequence; its hardware and precision are separate from these toy models.
- CUB ↗ and Compute Sanitizer ↗ — implementation primitives and debugging.
The browser models test exact address sets, contribution sums, output values and storage lifetimes. They do not execute CUDA. The CUDA snippets are reference examples, not GPU-validated benchmark artifacts. Compile, check and measure on the target GPU before relying on performance or deployment behavior.
Continue with Profiling, Performance, Attention or GPU Infrastructure.