GPU Performance Lab

Explore 70 hands-on lessons in GPU programming and profiling. Read each lesson here, then sign in to run supported exercises.

• Hello CUDA: Profile Your First Kernel

  Start with a small vector-add kernel: the CUDA equivalent of a practical Hello World. Edit the kernel body, choose profiling passes, and use the launch configuration below the kernel to experiment with grid and block dimensions. The host harness supplies A and B, launches your kernel for several input sizes, and verifies that every C[i] equals A[i] + B[i].

• Occupancy Warm-Up: Vector Add

  This vector-add kernel is already correct, but its launch starts with one thread per block. In the launch configuration below the kernel, increase the block's x dimension to give each block more useful warps and use GPUFlight metrics to reach the occupancy target.

• Grid-Stride Loop: Vector Add

  The launch intentionally uses a small grid. The starter kernel only handles one element per thread, so most of the array is left unwritten. Rewrite the body with a grid-stride loop so the same kernel works for every tested N.

• Block Occupancy: Give Each Block More Warps

  This vector-add kernel is correct, but the launch configuration starts with Block X = 32. That gives each block only one warp, so the hardware's resident-block-per-SM limit caps the number of active warps per SM even though warp capacity itself is nowhere near full. Increase the block's x dimension there and use GPUFlight's block occupancy metric to verify the change.

• Shared Memory Occupancy: Size the Reservation to What the Kernel Uses

  This kernel stages each block's inputs in two regions of dynamic shared memory before adding them. The launch reserves 49,152 bytes per block, which caps how many blocks fit on an SM. Trace the shared-memory indices, derive the smallest safe reservation for the configured block, and remove the unused reservation without making either region out of bounds.

• Reduction 3: Replace Global Atomic Contention

  Stage 2 left you with a correct kernel spending almost all of its time waiting, because every one of the 4,194,304 threads queued to update the same 4 bytes. Keep every thread and every value. Change only where the combining happens: fold inside each block first, then issue one atomicAdd per block instead of one per element.

• Reduction 4: Warp-Shuffle Block Sum

  Stage 3 folded a block through shared memory. Inside a single warp you do not need shared memory at all: __shfl_down_sync lets the 32 lanes read each other's registers directly, with no staging and no barrier. Read the driver tab first, because this stage changes what your kernel is asked to produce.

• Tiling at Scale: Read the Traffic, Not the Clock

  If you followed the Matrix Multiplication path you have already staged one small tile by hand. This is that same pattern at production scale, and it arrives with a trap: tile this N x N kernel and the wall clock will barely move. At these sizes the L2 cache is already absorbing the redundant reads for you, so elapsed time and SM throughput look nearly identical before and after. The SASS global-sector counter is what grades you here, because it measures the traffic your kernel actually issues rather than the traffic today's cache happens to hide. You will also meet the case a small hand-worked tile never shows you: the last partial tile, when N is not a multiple of the tile size.

• Coalescing: Shared-Memory Matrix Transpose

  This kernel transposes an N x N matrix correctly, but it is slow for a reason that has nothing to do with the math: its writes to global memory are uncoalesced. Neighbouring threads write to rows N apart, so each warp's store turns into many separate memory transactions. Stage a tile in shared memory so both the read and the write hit contiguous addresses.

• Streams: Overlap Copies with Compute

  This is a host-side lab: you write the orchestration, not a kernel. The serial version copies the whole array to the device, runs the compute kernel, and copies it back - three phases that each wait for the last. Split the work into chunks and drive them across several CUDA streams so a chunk's copy overlaps the previous chunk's compute. The harness times your version against the serial baseline and gates on the speedup.

• SAXPY: Make Adjacent Lanes Access Adjacent Data

  The kernel computes SAXPY correctly, but its thread-to-element mapping makes neighboring lanes touch elements far apart in memory. Keep the same arithmetic and fix only the indexing so each warp accesses contiguous values. GPUFlight's SASS coalescing efficiency is the judge, making this a direct lesson in how a source-level index expression becomes memory transactions.

• 1D Stencil: Reuse Neighbor Values With A Halo Tile

  The naive three-point stencil is correct, but every output reloads values its neighboring threads also requested. Stage one block of input plus a one-element halo on each side in shared memory, then reuse those values for the stencil. GPUFlight judges the reduction in issued global-memory sectors rather than cache-dependent wall time.

• Shared Memory Banks: Eliminate a 32-Way Conflict

  Each lane reads a different shared-memory word, but the starter spaces those words 32 apart. For 32-bit values, bank = word index mod 32, so all 32 addresses land in one bank and the warp request is serialized. Remap each warp to 32 adjacent words so lane 0 uses bank 0, lane 1 uses bank 1, and so on.

• Matrix Transpose: Pad Away Shared-Memory Bank Conflicts

  This transpose already coalesces both global-memory paths through a shared tile. The remaining bottleneck is inside shared memory: reading a column from tile[32][32] makes a warp touch 32 distinct words with the same bank index. Add one padding column so successive rows rotate through successive banks without changing the logical 32 by 32 tile.

• FP16 Is Not Enough: Activate Tensor Cores

  The starter already stores A and B as half precision, but each lane performs ordinary scalar multiply-add work, so the Tensor pipe remains idle. Make the full warp cooperate on 16 by 16 WMMA fragments and issue mma_sync while keeping float accumulation and the same matrix result.

• Tensor Cores: Keep the Accumulator in the Fragment

  The starter issues real WMMA instructions, but after every 16-element K slice it stores the partial C tile to global memory and loads it again for the next slice. Keep the accumulator fragment live across the entire K loop and store the completed tile once.

• Read-Only Cache: Keep a Lookup Page Hot in L1

  Experimental — correctness is graded; the performance target awaits GPU calibration. The profile table has 128 pages: 4 MiB fits in the GPU's shared L2 cache, but not in one SM's L1/TEX cache. The starter assigns round-robin logical queries directly to threads, so every 256-thread block cycles through all 128 pages. Remap workers so each block completes the 256 queries for one 32 KiB page, then write each result back to its original logical output slot.

• Constant Memory: Broadcast Hot Coefficients

  Every lane in a warp repeatedly requests the same entry from a small read-only coefficient table, but the starter uses a general global-memory path. Choose the CUDA memory path that matches a small, warp-uniform lookup and preserve the index pattern.

• Asynchronous Prefetch: Double-Buffer Shared Tiles

  Experimental — correctness is graded; the performance target awaits GPU calibration. The starter waits for each global-to-shared tile load before doing any arithmetic. Build a two-stage cuda::pipeline so tile n+1 moves into the unused shared buffer while the block computes tile n.

• Tensor Memory Accelerator: Load a 2D Tile

  Experimental — correctness is graded; the performance target awaits GPU calibration. The starter spends many thread instructions calculating addresses and copying a 32 by 32 tile into shared memory. On Hopper or newer, let one elected thread issue a descriptor-driven TMA transfer and let the block wait on a transaction barrier.

• Transfer Churn: Stage the Data Once, Not Per Row

  The starter copies 64 CPU rows to CUDA one at a time, computes each row sum, and returns a CUDA tensor. Stage the full input once and perform the equivalent row reductions on the device.

• Batch the Work: One Tensor Op Beats a Python Loop

  This function applies the same scale-and-bias to every row of a matrix, but it does it with a Python for-loop that touches one row at a time. Each tiny row operation becomes its own GPU kernel launch, so the work fans out into hundreds of launches plus per-iteration Python overhead. Keep the exact same result and rewrite solve() as a vectorized expression over the full tensor. GPUFlight counts the kernel launches inside the judged region, so it rewards collapsing the loop into one batched operation.

• Hidden Sync: Stop Reading Scalars Off the GPU Every Step

  For a CUDA tensor, .item() makes the calling host wait until the scalar is available. Repeated reads can interrupt enqueueing and reduce overlap. Other independent GPU work may continue. Accumulate on the device and return one Python float after the final read.

• Stage the Batch: One Copy, One Matmul

  This function scores a batch of CPU feature rows against a weight matrix on the GPU - logits = data @ weight. The starter copies one row to the device inside a loop and runs a tiny matmul per row, so both the PCIe bus and the launch queue set the pace instead of the math. Keep the same result and stage the whole batch to the GPU once, then do a single batched matmul. GPUFlight counts host-to-device transfers AND kernel launches in the judged region, so it rewards staging the input once and batching the matmul.

• CUB BlockReduce: Match Hand-Tuned Performance In Five Lines

  You have already hand-built the block reduction: stage values in shared memory, fold them in log2(blockDim) barrier steps, then issue one atomic per block. Production CUDA code rarely writes that by hand. NVIDIA ships the same algorithm, tuned per architecture, as cub::BlockReduce in the CUB library (part of the CUDA toolkit). Replace the contended atomic starter with a BlockReduce call and match the hand-tuned kernel's throughput with a fraction of the code.

• Matmul Ladder I: Give Each Thread A Column Strip

  Shared-memory tiling fixed the global-memory traffic. The next ceiling is shared memory itself: a one-output-per-thread kernel reads shared memory twice for every multiply-add it performs. This lab moves to the register-blocking harness - each block owns a fixed 64x64 tile of C, delivered as a flat 1-D grid - and asks each thread to carry a strip of 16 outputs in registers, so one shared-memory read of B is reused 16 times.

• Reduction 5: Stable Softmax, Two Row Reductions

  The tree you built in stages 3 and 4 was never really about addition. Softmax needs the same tree twice per row, once with fmaxf and once with +, and the max one is there to keep exp from overflowing rather than to make anything faster. This is the reduction pattern arriving at a real Transformer kernel.

• Reduction 6: Online Softmax, One Read Instead of Three

  Your stage 5 kernel is correct, stable, and fully parallel. It also reads every row from global memory three times: once for the max, once for the sum, once to normalize. This stage removes two of those traversals without changing a single answer, using the identity modern attention kernels are built on.

• LayerNorm: One Block Owns One Row

  LayerNorm sits before or after every attention and MLP block in a transformer: normalize each row to zero mean and unit variance, guarded by a small epsilon. The starter assigns one THREAD per row - with this launch that leaves nearly every thread idle and each working thread walking 1024 columns serially, three times. Rewrite it so one BLOCK owns one row and cooperates on both reductions.

• Welford LayerNorm: Mean And Variance In One Read

  Experimental — correctness is graded; the performance target awaits GPU calibration. The two-pass LayerNorm reads every row twice before it can normalize - and a third time to do it. Welford's algorithm (1962) maintains a running (count, mean, M2) triple that yields BOTH the mean and the variance in a single pass, with the numerical robustness the naive E[x^2] shortcut lacks. Chan's merge formula combines partial triples, so the block can tree-combine per-thread results. Add a shared row cache and the row touches global memory once in, once out - this is how production LayerNorms (PyTorch/Apex) are written.

• RMSNorm: Fuse The Scale Into One Pass

  Experimental — correctness is graded; the performance target awaits GPU calibration. RMSNorm is LayerNorm minus the mean-centering: y = x / rms(x) * gamma, where rms is the root-mean-square of the row. LLaMA-family models use it everywhere because it is cheaper and just as effective. The starter computes the sum of squares in one traversal, then re-reads the row from global memory to normalize and scale. Stage the row in shared memory during the statistics pass and the second global read disappears.

• Fusion Warm-Up: Remove the Scratch Round Trip

  Stage 5 fixed how a warp reaches memory. Now reduce both the traffic and the storage used by an unnecessary intermediate. Fuse the two loops in solution.cu, then remove the N-sized scratch allocation from pipeline.cu.

• SwiGLU: Apply Fusion to a Transformer MLP

  Stage 6 removed a scratch-buffer round trip from SAXPY followed by scaling. Now apply the same optimization inside a Transformer MLP: fuse the two passes in solution.cu, then remove the N-sized scratch allocation from pipeline.cu.

• Quantized GEMV: Decode Is A Bandwidth Problem

  Experimental — correctness is graded; performance awaits GPU calibration. This single-vector GEMV reads FP32 weights. Replace them with the supplied int8 values and per-row scales, then compare issued sectors and elapsed time. Test weights are constructed from that quantized representation; arbitrary model weights would introduce accuracy loss.

• Embed GPUFlight: Own the Profiling Lifecycle

  The starter already has a complete CUDA vector-add and its result check. Your job is to wire in the GPUFlight profiling lifecycle yourself: initialize the SDK before the GPU work, capture kernel activity with the Trace engine, print a session report, and shut the SDK down cleanly. Follow the three TODO markers in main().

• Inference server: batch the requests

  A decoder serves many requests, but the naive step advances the sequences one at a time, so each tiny kernel uses a sliver of the GPU and the per-step latency is paid once per sequence. Profile the run, then advance the whole batch in one launch to fill the GPU and raise throughput without changing the output.

• Inference server: cache the keys and values

  A batched attention decoder generates tokens for many sequences at once. The naive step recomputes every past token's key and value projection, for every sequence, on every step, so the work grows with the sequence length. Cache the projections and reuse them to avoid repeated K/V projections; attention still grows with context length.

• Matrix Multiplication 1: CPU Baseline

  Build a correct CPU matrix multiplication baseline before moving the same computation to CUDA. For each output element C[row, col], multiply one row of A by one column of B and sum the products.

• Matrix Multiplication 2: First CUDA Kernel

  You have a real your selected GPU. Take the same N = 4096 matrix multiplication you built on one CPU core and launch it across thousands of CUDA threads. Each thread owns one output element of C. Write the kernel, run it on the GPU, and watch your first CUDA computation come to life.

• Matrix Multiplication 3: Share One 4 x 4 Tile

  Every output C[row, col] needs one row of A and one column of B. That means neighboring output threads request many of the same A and B values. Use shared memory as a small workspace for one 4 x 4 matrix multiply: let the 16 threads load A and B together, synchronize, then reuse those tiles to calculate C.

• Matrix Multiplication 4: Reuse a Shared-Memory Tile

  Apply the shared-memory tile pattern to N = 4096 on your selected GPU. Load each 16 x 16 A and B tile once, let the block reuse it, then compare the profile with the naive kernel from Stage 2.

• Matrix Multiplication 5: Compare Your Kernel with cuBLAS

  If you have not finished Stage 4, we strongly recommend doing that first. This stage uses its tiled matmul as the baseline. If shared-memory tiling already makes sense to you, you can start here. Your task is to replace that matmul with a cuBLAS call and compare the results.

• Elementwise 1: CPU Vector Add Baseline

  Add two arrays element by element on one CPU core, and record how long it takes. This is the same shape a Transformer uses when it adds an attention or MLP output back into the residual stream.

• Elementwise 2: The Same Add on the GPU

  Move the CPU loop from stage 1 onto the GPU. The arithmetic does not change; what changes is that the loop counter becomes a thread index and many elements are computed in parallel.

• Reduction 1: CPU Sum Baseline

  Add 4,194,304 floats into a single number on one CPU core, and record how long it takes. Every attention softmax, every LayerNorm, and every loss value in a Transformer starts with a reduction like this one.

• Reduction 2: One Atomic Per Value

  Move stage 1's sum onto the GPU with the smallest change that is genuinely correct: one thread per element, and one atomicAdd each into a single output. It gives the right answer. Then look at what it costs.

• Decoder Block 1: The CPU Reference

  Write the complete mathematical reference for one modern pre-norm decoder block: RMSNorm, packed QKV, RoPE, grouped-query causal attention, and a SwiGLU feed-forward, for a four-token sequence on one CPU core. This is the block recipe shared by the Llama, Mistral and Qwen model families, and every CUDA stage in this path is judged against what you build here.

• Decoder Block 2: The CUDA Pipeline

  Optimize a correct fourteen-launch decoder baseline. The starter implements all six kernel bodies; the trusted driver checks nine intermediate tensors before you optimize individual operators.

• Decoder Block 3: Packed QKV Projection

  Fuse the three QKV projection launches into one T x 512 x 768 GEMM over the packed weight slab, then tile it. Packing alone buys about 2x; the gate asks for 3.5x against the same-run three-launch baseline, which is where shared-memory tiling comes in.

• Decoder Block 4: RoPE and GQA Are Layout, Not Copies

  Rotate Q's eight heads and K's two heads in place inside the packed (T, 768) tensor from stage 3, using precomputed per-position tables. V is never rotated, nothing is ever copied, and every mistake is reported as the exact head it landed in.

• Decoder Block 5: Causal GQA Attention, Correct First

  One block per (query row, query head): compute scaled scores against the mapped KV head, mask the future, run a stable softmax, and weight V. The input contains planted rows whose scores overflow float exp, so stability is graded, not advised. No speed gate: the algorithm comes first.

• Decoder Block 6: Online Attention, the Capstone

  Replace stage 5's materialized score row with tiled online attention inside the full decoder block: per-row running max, denominator and output accumulator, K/V tiles staged once per QUERY TILE instead of once per query row. Timed against stage 5's exact kernel in the same run; the gate asks for 1.8x.

• 2D Stencil: Pay the Halo Ring Once

  A five-point stencil replaces each interior point of an N x N grid with a weighted average of itself and its four neighbours, and copies the border through unchanged. The starter is correct but every thread fetches all five inputs straight from global-memory request, so neighbouring threads keep re-requesting the same points. Stage the block's tile plus a one-cell halo ring in shared memory and take every read from there. GPUFlight judges issued global-memory sectors, which count the requests before caches can hide them.

• 2D Stencil: Roll the Row Window Through Registers

  Same five-point stencil, different ownership: each block now owns a tall strip of rows and marches down it, so the row window (y-1, y, y+1) shifts by one every step. The starter re-reads all three rows from global-memory request at every step; the fix keeps two of them in registers and reads only the new row. GPUFlight judges issued global-memory sectors.

• 3D Stencil: Stage the Plane, Pay for the Depth

  A seven-point stencil over an N x N x N cube reads each point's six face neighbours: four in the same z plane, one behind, one ahead. The starter fetches all seven values from global-memory request per output; the fix stages each z plane in a shared tile as the block sweeps, which serves five of the seven reads on-chip and honestly cannot serve the other two. GPUFlight judges issued global-memory sectors.

• 3D Stencil: Rotate the Depth Window Through Registers

  The starter is the previous lab's accepted solution, byte for byte: plane staged in shared memory, both depth reads still global-memory request. The fix is 2.5D blocking, the technique production stencil code actually uses: the plane just used and the plane coming next ride in registers and rotate one step per iteration, so every input plane is fetched from global-memory request exactly once. GPUFlight judges issued global-memory sectors.

• Telemetry 1: Collect Slow Kernel Events

  Keep every kernel event whose duration meets the threshold. The starter is correct, but every kept event reserves its output position through the same global counter — about 2.1 million contended atomics at the judged size. Change the reservation strategy so each block performs one global atomic without dropping or duplicating events. Output order does not matter in this lab.

• Telemetry 2: Preserve the Timeline Order

  Keep the slow events in their original timeline order. The starter is correct but serial: one thread walks the full timeline while the rest of the block is idle. For [800, 4200, 6100, 900, 7300] with threshold 4000, the output must be [4200, 6100, 7300]. Replace the serial loop with a one-block exclusive prefix scan that assigns stable output positions without atomics.

• Telemetry 3: Build a Kernel-Latency Distribution

  Count kernel durations into five latency bins. The starter is correct, but all 4,194,304 events update only five global counters, and a short-kernel-heavy workload funnels most updates into bin 0. Give each block a private shared-memory histogram, then merge each private bin into the global result once.

• Telemetry 4: Filter the Events In Place

  Compact the kept events into the front of the same array, preserve their order, and report the new count. The starter is correct but serial: one thread owns the read/write order, so it cannot overwrite an unread event. Replace it with a block-parallel scan while preserving that safety rule. For [800, 4200, 6100, 900, 7300] with threshold 4000, the valid prefix becomes [4200, 6100, 7300] and filteredCount is 3.

• Can RTX 6000 Match H100?

  Test a hypothesis: independent FP32 Horner chains may exploit RTX 6000 Ada well enough to reach 80% of a measured H100 reference. Optimize, measure native kernel time, then profile both GPUs to explain what changed.

• Why Is My GPU Waiting for the Next Batch?

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. Start with eight pinned CPU batches and a CUDA projection.

• I Added torch.compile. Why Is It Still Slow?

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. The starter compiles a preprocessing chain with an intentional scalar extraction.

• Why Did My Attention Layer Allocate So Much Memory?

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. Replace explicit scores, causal masking, softmax and values with scaled_dot_product_attention.

• My Model Recompiles When the Batch Size Changes

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. The request batches are [7,16,9,32,7,24], totaling 95 rows.

• Faster Serving, Slower Users?

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. Choose a maximum batch size and maximum wait for a deterministic sixteen-request arrival trace.

• A Harmless contiguous() Call Copied My Whole Tensor

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. A transpose changes shape/strides without copying data.

• Same Kernels, Fewer Python Submissions: CUDA Graphs

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. Capture sixteen static-shape inference steps after warmup on a side stream.

• Inference: Batch Size, Memory and Precision

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. Process exactly 256 examples through a two-layer projection.

• Training: Accumulate Gradients with Equal Work

  Experimental performance lab. Correctness is graded; speed and memory observations await GPU calibration. Compute the weight gradient of mean squared error over all 256 examples.