Advanced

Shared Memory and Bank Conflicts

Shared memory is the programmer-managed SRAM inside each SM, split into 32 four-byte banks that serve one word each per cycle. When several lanes of a warp hit the same bank at different addresses the access serializes, and a 32-way conflict makes a shared-memory-bound loop run over ten times slower. Padding, XOR swizzles, cp.async and TMA are the tools that decide whether a tiled kernel gets the bandwidth it staged data for.

Unlock the full curriculum — ₹2,000 / $25every concept + every answer · 6 months · no auto-renew

RELATED CONCEPTS

GPU Memory Hierarchy Memory Coalescing Tiled Matrix Multiplication Occupancy and Register Pressure CUTLASS and Tensor Core Kernels FlashAttention Internals

PRACTICE THIS IN REAL QUESTIONS

CUDA, Triton & Kernel EngineeringExplain shared memory bank conflicts with the bank arithmetic, show a kernel that has them, and fix it with padding.→CUDA, Triton & Kernel EngineeringTranspose a large matrix at close to copy bandwidth. Why is the naive version slow, and what does each fix buy?→CUDA, Triton & Kernel EngineeringTake a GEMM from naive to 70% of peak: the steps, the speedup at each, and the arithmetic that says why.→GPU & Accelerator ArchitectureA kernel runs at 25% occupancy. Is that a problem? Walk me through what occupancy buys and when you would leave it low on purpose.→CUDA, Triton & Kernel EngineeringWhat changed between FlashAttention 1, 2 and 3, and why did each change buy what it did?→CUDA, Triton & Kernel EngineeringHow do you overlap host-to-device transfers with compute, and what has to be true before the overlap actually happens?→