Streaming Multiprocessors, Memory Hierarchies & Coalescing
From High-Bandwidth Memory (HBM3) to L2 Cache and on-chip SRAM. Why memory coalescing and warp occupancy dictate 90% of model throughput.
- •Map thread blocks and 32-thread warps to Streaming Multiprocessors (SMs)
- •Structure memory pointers for 128-byte coalesced global memory transactions
- •Eliminate shared memory 32-way bank conflicts through swizzling techniques
- •Analyze warp scheduler stall reasons with NVIDIA Nsight Compute
The Silicon Reality: Why Memory Access Dictates Performance
An NVIDIA H100 SXM5 GPU features 132 Streaming Multiprocessors (SMs) capable of a staggering 1,979 TFLOPS of FP8 tensor compute. However, its High-Bandwidth Memory (HBM3) interface is capped at 3.35 TB/s. If a 32-thread warp attempts to load memory with stride or scatter patterns, the memory controller issues multiple serialized 32-byte sector transactions, throttling bandwidth down to less than 15% of peak. True kernel acceleration requires coalesced loads: 32 threads accessing contiguous 4-byte or 2-byte elements in a single 128-byte burst.
NVIDIA H100 SM90 Silicon Memory Hierarchy
| Memory Tier | Physical Capacity | Peak Bandwidth | Latency Access | Architectural Role |
|---|---|---|---|---|
| Registers (RF) | 64K x 32-bit / SM | 33.0 TB/s | 1 cycle | Hardware buffer stage |
| Shared Memory (SRAM) | 228 KB / SM | 19.3 TB/s | 19 cycles | Hardware buffer stage |
| L2 Cache (SRAM) | 50 MB Shared | 12.0 TB/s | 32 cycles | Hardware buffer stage |
| HBM3 VRAM | 80 GB HBM3 | 3.35 TB/s | 420 cycles | Hardware buffer stage |
| PCIe Gen5 / Host RAM | 2 TB System RAM | 64 GB/s | 12,000+ cycles | Hardware buffer stage |
Notice the latency chasm: fetching a tensor from HBM3 takes ~420 clock cycles, while reading from SRAM takes only ~19 cycles. Writing high-speed kernels is almost entirely the art of keeping data in SRAM.
Production CUDA Memory Architecture Suite
Inspect the production CUDA implementation: comparing uncoalesced strided access against 128-byte coalesced burst transfers, along with a shared memory bank conflict swizzler.
// Coalesced vs Strided Global Memory Load Benchmark
#include <cuda_fp16.h>
#include <cuda_runtime.h>
// Coalesced: Adjacent threads access adjacent memory addresses
__global__ void coalesced_copy_kernel(const float* __restrict__ d_in, float* __restrict__ d_out, int n) {
int idx = blockDim.x * blockIdx.x + threadIdx.x;
if (idx < n) {
// Single 128-byte transaction per warp for 32 floats!
d_out[idx] = d_in[idx];
}
}
// Strided: Non-coalesced access forcing 32 separate serialized transactions
__global__ void strided_copy_kernel(const float* __restrict__ d_in, float* __restrict__ d_out, int stride, int n) {
int idx = (blockDim.x * blockIdx.x + threadIdx.x) * stride;
if (idx < n) {
d_out[idx] = d_in[idx];
}
}Run this lesson's Triton kernels and multi-file code directly in an ephemeral pre-configured PyTorch 2.5 container.