GPU memory bandwidth optimization

Overview

GPU memory bandwidth optimization is the practice of structuring memory access patterns to maximize data transfer rates between GPU cores and memory, critical because memory bandwidth is often the primary performance bottleneck in GPU computing.

The fundamental problem: Modern GPUs have thousands of cores but limited memory buses. A single NVIDIA A100 has 6,912 CUDA cores sharing a ~2 TB/s memory bandwidth. If each core requests data independently, the memory system becomes congested. Smart access patterns can achieve 10-100× speedups.

Core Concepts

Why it exists: GPU memory controllers read data in large chunks (32-128 bytes). When thread0 reads address A, thread 1 reads A+4, thread 2 reads A+8, etc., the hardware fetches one 128-byte block instead of 32 separate 4-byte reads.

Derivation of coalesced bandwidth:

For a warp of $w = 32$ threads:

• Each thread requests $b$ bytes
• Memory transaction size: $T$ bytes (typically 128B)
• Coalesced case: addresses span $\leq T$ bytes

$\text{Transactions needed} = \lceil \frac{w \cdot b}{T} \rceil$

Best case (fully coalesced): All32 threads access consecutive floats (4 bytes each): $\text{Span} = 32 \times 4 = 128 \text{ bytes} = 1 \text{ transaction}$

Worst case (strided by 32): Thread $i$ accesses address $A +32i \times 4$: $\text{Span} = 32 \times 32 \times 4 = 4096 \text{ bytes} = 32 \text{ transactions}$

Effective bandwidth ratio: $\text{Efficiency} = \frac{\text{Data used}}{\text{Data transferred}} = \frac{w \cdot b}{\text{Transactions} \times T}$

Coalesced: $\frac{128}{128} = 100\%$ | Strided: $\frac{128}{4096} = 3.125\%$

$M = B \times \eta$

where: $\eta = \frac{\text{Useful bytes loaded}}{\text{Total bytes transferred from DRAM}}$

Derivation from hardware constraints:

1. Physical bandwidth: $B = \frac{\text{Bus width} \times \text{Clock rate} \times 2 \text{ (DDR)}}{\text{8 bits/byte}}$
2. Transaction overhead: Each memory request has minimum size $T_{min}$
3. Warp divergence: When threads access non-contiguous addresses, each subset needs separate transaction
4. Bank conflicts: In shared memory, parallel access to same bank serializes

Combined efficiency: $\eta_{total} = \eta_{coalesce} \times \eta_{cache} \times \eta_{occupancy}$

• $\eta_{coalesce}$: Coalescing efficiency (derived above)
• $\eta_{cache}$: Cache hit rate (reduces DRAM traffic)
• $\eta_{occupancy}$: Fraction of time memory controllers are busy

Why banks exist: Hardware cost. To provide32-way parallel access to arbitrary addresses would require full crossbar ($32 \times 32$ connections). Banks partition memory so each thread maps to one bank, enabling simpler hardware with $32 \times 1$ connections per bank.

Bank assignment: For address $A$ and $N$ banks: $\text{Bank} = \left(\frac{A}{\text{word size}}\right) \bmod N$

Conflict rule: $k$ threads accessing bank $b$ take $k$ cycles instead of 1.

Broadcasting exception: If all threads read the same address, hardware broadcasts in1 cycle (no conflict).

Optimization Techniques

1. Array-of-Structures vs Structure-of-Arrays

Array of Structures (AoS):

struct Particle { float x, y, z, vx, vy, vz; };
Particle particles[N];
 
// Thread i accesses:
float x = particles[i].x;  // Address: base + i*24 + 0
float y = particles[i].y;  // Address: base + i*24 + 4

Why this is slow: Thread 0 reads address A, thread 1 reads A+24, thread 2 reads A+48. The field the threads want (x) is scattered every 24 bytes. The hardware still fetches 128-byte lines, but a warp reading one float per thread now touches $32 \times 24 = 768$ bytes of address span. That requires $\lceil 768 / 128 \rceil = 6$ transactions per field — but of the 768 bytes fetched, only $32 \times 4 = 128$ bytes are useful (efficiency $\approx 16.7\%$). So AoS partially coalesces; it is not fully serialized into 32 separate reads, but it wastes ~5/6 of the bandwidth.

Structure of Arrays (SoA):

struct Particles {
    float x[N], y[N], z[N];
    float vx[N], vy[N], vz[N];
};
Particles particles;
 
// Thread i accesses:
float x = particles.x[i];  // Address: base_x + i*4
float y = particles.y[i];  // Address: base_y + i*4

Why this is fast: Thread 0 reads A, thread 1 reads A+4, thread 2 reads A+8. Addresses are consecutive, so hardware combines into 1 transaction per warp per field.

Speedup calculation: For 6 fields, per warp:

• AoS: $6$ fields $\times 6$ transactions/field $= 36$ transactions
• SoA: $6$ fields $\times 1$ transaction/field $= 6$ transactions

$\text{Reduction} = \frac{36}{6} = 6\times \text{ fewer memory transactions}$

(The gap grows for wider structs: a struct spanning $S$ bytes per element gives $\lceil 32S/128 \rceil$ transactions per field for AoS vs 1 for SoA.)

2. Tiling with Shared Memory

Problem: For $N \times N$ matrices, each thread loads $2N$ elements from global memory. Total: $N^2 \times 2N = 2N^3$ global element reads (many redundant).

Tiled approach:

1. Divide matrices into $T \times T$ tiles
2. Load one tile of $A$ and one tile of $B$ into shared memory once per block
3. All threads in block reuse the shared data
4. Move to next tile

Why this step? Shared memory is ~100× faster than global memory and shared within a block. By loading tiles cooperatively, we convert $N$ redundant global reads into 1 global read + $N$ fast shared reads.

Derivation of speedup:

Without tiling: $\text{Global reads per output element} = 2N$

With $T \times T$ tiles: $\text{Tiles along the k-dimension} = \frac{N}{T}$ $\text{Global reads per output element} = \frac{2N}{T} \text{ (global)} + 2N \text{ (shared)}$

Global memory traffic reduction: $\text{Reduction factor} = \frac{2N}{2N/T} = T$

For $T = 32$: 32× less global memory traffic.

Full derivation of arithmetic intensity:

Operation count: $2N^3$ FLOPs (each multiply-add = 2 FLOPs, done $N^3$ times).

Global data movement (tiled), counting bytes: each output element triggers $\frac{2N}{T}$ global element reads, and each element is a 4-byte float: $\text{Bytes moved} = N^2 \times \frac{2N}{T} \times 4 = \frac{8N^3}{T} \text{ bytes}$

Arithmetic intensity: $I = \frac{\text{FLOPs}}{\text{Bytes}} = \frac{2N^3}{8N^3/T} = \frac{T}{4} \text{ FLOPs/byte}$

For $T = 32$: $I = 8$ FLOPs/byte.

Compute-bound vs bandwidth-bound (Roofline): A kernel is compute-bound when its arithmetic intensity exceeds the machine balance $I^* = \frac{\text{peak FLOP/s}}{\text{peak bytes/s}}$. On a GPU with 20 TFLOP/s peak and 1 TB/s bandwidth: $I^* = \frac{20 \times 10^{12}}{1 \times 10^{12}} = 20 \text{ FLOPs/byte}$

• Tiled ($T=32$): $I = 8$ FLOPs/byte $< 20$ → still bandwidth-bound, but 32× less traffic than naive.
• Larger tiles / register blocking push $I$ toward and past $20$ to become compute-bound.
• Naive (no tiling, $T=1$): $I = 0.25$ FLOPs/byte $\ll 20$ → heavily bandwidth-bound.

3. Avoiding Bank Conflicts in Shared Memory

Naive approach:

__shared__ float tile[32][32];
 
// Load: thread i reads column i → coalesced
tile[threadIdx.y][threadIdx.x] = input[...];
 
// Store: thread i writes row i → each thread different bank
output[...] = tile[threadIdx.x][threadIdx.y];  // CONFLICT!

Why conflict occurs: When thread 0 reads tile[0][0], thread 1 reads tile[1][0], .., thread 31 reads tile[31][0]. All access column 0, which maps to the same bank (bank 0) → 32-way conflict → 32× slower.

Bank mapping: Address tile[row][col] = base + (row*32 + col)*4 bytes $\text{Bank} = (row \times 32 + col) \bmod 32 = col \bmod 32$

All threads accessing column 0 hit bank 0.

Solution: Padding

__shared__ float tile[32][33];  // Extra column

Why this works: Now address tile[row][col] = base + (row*33 + col)*4 $\text{Bank} = (row \times 33 + col) \bmod 32$

For column 0: thread $i$ accesses bank $(i \times 33) \bmod 32$. Since $\gcd(33, 32) = 1$, consecutive threads map to different banks → no conflict!

Cost: 3% memory overhead (33/32) → 32× speedup in transpose.

Why it feels right: On CPUs, caches are large relative to working sets and provide good automatic optimization.

The reality: GPU caches are small per-thread (128 KB L1 per SM shared by 1024+ threads = 128 bytes per thread). With 32-thread warps executing in lockstep, if threads access scattered addresses, the cache gets thrashed before data can be reused.

Example: Matrix multiply without tiling. Each thread needs $2N$ elements, but cache holds only ~32 floats per thread. For $N = 1024$, 99.7% of accesses miss the cache → bandwidth-bound.

The fix: Explicitly use shared memory for deliberate data reuse. Don't rely on automatic caching.

Why it feels right: The addresses are strided (every 24 bytes), so it seems nothing coalesces.

The reality: The hardware always fetches full 128-byte cache lines. A warp reading one float per thread from a 24-byte-stride struct spans $32 \times 24 = 768$ bytes → $\lceil 768/128 \rceil = 6$ transactions, not 32. It partially coalesces but wastes ~83% of the bandwidth (only 128 of 768 bytes useful).

The fix: Use SoA to collapse to 1 transaction/field. The real penalty of AoS is wasted bandwidth (efficiency ~17%), not "32× more transactions."

The reality: Memory transactions must be aligned to 32/64/128-byte boundaries. If array starts at address 0x102 (misaligned), even consecutive accesses split across boundaries.

Example: Array starts at 0x102. Warp reads 0x102-0x181 (128 bytes). This spans boundaries 0x100 and 0x180 → 2 transactions instead of 1.

The fix: Use cudaMalloc (returns aligned pointers) or __align__ attribute:

__align__(128) float data[N];

Active Recall Practice

GPUs are like that library. They have thousands of "students" (threads) all working at once. The memory is the library shelves. If threads ask for data randomly, the memory system gets overwhelmed. But if nearby threads ask for nearby data, the memory system can deliver everything in one trip.

There are tricks: (1) organize data so nearby threads need nearby data (like arranging students by their shelf), (2) use a small fast "desk" (shared memory) to share data between students in the same room, (3) make sure students don't all grab from the same drawer at once (bank conflicts). These tricks make GPU programs 10-100× faster!

Connections

• GPU-Warp-Scheduling: coalescing works within warps (32 threads)
• Cache-Hierarchy: L1/L2 caches complement but don't replace explicit optimization
• Roofline-Model: bandwidth optimization moves programs toward compute-bound regime
• Parallel-Algorithm-Design: memory access patterns influence algorithm choice
• CUDA-Shared-Memory: hardware details of banking and conflicts
• Memory-Latency-Hiding: high occupancy allows latency tolerance

Summary

GPU memory bandwidth optimization centers on coalescing (making threads access consecutive addresses), tiling (reusing data in fast shared memory), and avoiding bank conflicts (structuring shared memory access patterns). The fundamental trade-off: small changes in data layout yield 10-100× speedups because memory bandwidth is the GPU's scarcest resource. Every optimization derives from the hardware truth that loading 128 bytes costs the same as loading 4 bytes – so pack threads' requests together.

#flashcards/hardware

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)