GPU Kernel Performance Engineering Lab

Four lines change, the timer drops, and the kernel looks faster. Then one output check finds stale values. The optimization didn't win anything. It made the program wrong more quickly.

This lab turns that failure into a repeatable engineering habit. You'll keep one square matrix, one data type, one launch geometry, and one timing harness while changing only the kernel's memory path. Every candidate must pass exact output checks and CUDA correctness tools before its timing counts.

Three prerequisites supply the vocabulary. CUDA for ML Training established grids, blocks, warps, asynchronous launches, and memory hierarchy. Accelerator Architecture Field Guide connected arithmetic intensity and resource ceilings to workload bounds. GPU Profiling, Correctness, and Benchmarking separated host timelines, kernel counters, and correctness evidence. Here those ideas control one kernel from first measurement through recovery.

Keep one experiment stable

The workload transposes an $N \times N$ row-major float32 matrix. If input element $A[r,c]$ lives at input[r * N + c], correct output satisfies:

$$B[c,r] = A[r,c]$$

No floating-point arithmetic changes a value. A correct result must therefore match bit for bit, so max_abs_error should be 0. The default matrix gives every location a distinct integer-valued float, which keeps a misplaced element from hiding behind repeated test data. That exact contract is stronger than a loose tolerance and simpler than checking only a checksum.

The supplied harness contains five variants:

copy isn't a transpose baseline. It performs a different operation with the same logical read and write volume. Treat it as context for the memory path, not as a denominator for a claimed transpose speedup.

Every iteration follows the same loop:

Record hypothesis before editing. Otherwise every surprising counter becomes a story invented after result.

Build reproducible lab

You need an NVIDIA GPU, a CUDA Toolkit with nvcc and Compute Sanitizer, and Nsight Compute's ncu command. Tool versions and GPU identity belong beside every result because compiler decisions, architecture, clocks, and profiler metric availability can change evidence.

Save transpose lab source as transpose_lab.cu. Also save recovery record template for deliberate failure later.

First create one artifact directory and capture environment. These commands produce text files you can compare without relying on terminal scrollback:

Compile optimized code with line information for profiler and sanitizer attribution. -Xptxas=-v also prints per-kernel register and static shared-memory use:

Expected build evidence is structural, not one universal register count:

• compiler names each instantiated kernel;
• transpose_tiled<0> reserves $32 \times 32 \times 4 = 4096$ bytes of static shared memory;
• transpose_tiled<1> reserves $32 \times 33 \times 4 = 4224$ bytes;
• register count is recorded instead of assumed;
• no compilation error is present.

Run fixed $4096 \times 4096$ matrix for 100 measured launches per normal variant:

Harness prints CSV columns variant, correct, mismatches, max_abs_error, average_ms, and effective_gb_s. Device and launch shape go to standard error, so tee keeps CSV machine-readable while terminal still names hardware. Pass conditions are:

Don't rank variants yet. A single process run can be affected by clocks, temperature, competing work, or startup state. Repeat separate processes under same conditions, alternate order by round, and preserve every row:

Compare distributions, not only best row. If one candidate wins once and overlaps baseline in later runs, evidence doesn't support keeping it.

Read naive memory path

CUDA coalesces a warp's global-memory requests into transactions that cover requested addresses. Adjacent float32 words need fewer transactions than words spread across matrix rows.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/} The naive kernel reads input with good spatial order, then loses that order on output.

Here is complete indexing logic from transpose_naive. A block has 32 threads in x and 8 in y; each thread handles four rows separated by kBlockRows:

Within one warp, threadIdx.x runs from 0 through 31 while threadIdx.y stays fixed. Input indexes therefore differ by one float. Output indexes differ by $N$ floats because x became output row. Correct index mapping creates inefficient stores.

Transpose does almost no useful floating-point arithmetic. Its logical data volume per launch is one read and one write for every element:

$$\text{logical bytes} = 2N^2 \cdot \operatorname{sizeof}(\text{float})$$

For $N=4096$, that's $134{,}217{,}728$ bytes, or 128 MiB. Harness reports:

$$\text{effective GB/s} = \frac{2N^2 \cdot 4}{\text{average time in seconds} \cdot 10^9}$$

Effective bandwidth isn't physical DRAM traffic. Cache hits, excess transactions, error-correction traffic, and architecture-specific paths can make hardware traffic differ. Use it as a stable workload-normalized rate for comparing same operation on same setup.

A $4096 \times 4096$ transpose averages $0.250$ ms. What effective rate should the ledger record, and what hardware claim remains unsupported?

Answer

The logical volume is $2 \times 4096^2 \times 4 = 134{,}217{,}728$ bytes. Dividing by $0.000250$ seconds and then by $10^9$ gives about $536.9$ GB/s. That is a workload-normalized effective rate. It doesn't establish physical DRAM traffic or percentage of peak bandwidth without profiler counters and the matching hardware conditions.

Profile before editing

Nsight Compute groups kernel evidence into named sections. LaunchStats reports grid, block, registers, and shared memory. SpeedOfLight compares compute and memory resource throughput with device ceilings. MemoryWorkloadAnalysis traces traffic through memory hierarchy. Occupancy reports active-warps capacity and its resource limiters.^{[2]Reference 2Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Collect one naive launch. Full set may replay kernel several times to gather counters, so don't use profiler run as timing result:

Read report in this order:

1. Confirm kernel name, matrix size, block 32 x 8, and one profiled launch.
2. Check whether memory throughput is closer to its ceiling than compute throughput.
3. Inspect global load and store requests, sectors, and reported access inefficiencies.
4. Check static shared memory is zero for naive variant.
5. Read occupancy limiter, but don't optimize it unless evidence connects limiter to latency or throughput.

Roofline reasoning predicts a memory-bound transpose because arithmetic intensity is near zero. Profiler still matters: it distinguishes useful memory traffic from inefficient transactions and verifies that observed kernel is one you meant to measure. Roofline is a ceiling model, not a promise that every low-intensity kernel reaches peak bandwidth.^{[3]Reference 3Roofline: An Insightful Visual Performance Model for Multicore Architectureshttps://doi.org/10.1145/1498765.1498785}

Write one bound statement before editing:

If report doesn't support memory-path hypothesis, stop. A shared-memory rewrite would be guesswork.

Iteration one: coalesce both global directions

Shared memory is on-chip storage shared by threads in one block. Tiled transpose first copies input tile into shared memory with adjacent global reads. After block-wide barrier, threads read tile with swapped indexes and write adjacent output words. CUDA's official guidance uses shared memory for this exact reordering pattern.^{[4]Reference 4CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/}

The core kernel below accepts Padding as compile-time tile-width change. Variant tiled instantiates Padding = 0:

Barrier is a correctness boundary. Every thread must finish writes into tile before any thread reads values written by another warp. Stream order doesn't provide that block-internal guarantee.

Run correctness and event timing first, then collect profiler report:

Expected evidence has two parts:

• global reads and writes now use adjacent lane addresses;
• shared-memory column read maps many lanes onto same bank.

Shared memory has 32 banks for the access model used here, and successive 32-bit words map to successive banks. Different words requested from same bank in one warp are serialized, except supported broadcast case where threads request same word.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

For unpadded tile, column element for lane $i$ has word offset $32i + c$. Its bank is:

$$\operatorname{bank}(i) = (32i + c) \bmod 32 = c$$

All 32 lanes address different words in bank $c$. Global traffic improved, but shared-memory conflict is now visible.

Profiler versions and architectures can expose bank evidence with different metric names. Use report's Shared Memory tables and source correlation first. If you need raw name for scripted collection, query installed tool instead of copying a metric from another GPU:

Record metric name, value, and source line. A screenshot without workload identity or source version isn't enough.

Iteration two: pad shared tile

Candidate change is one token in tile width:

Column element for lane $i$ now has word offset $33i + c$:

$$\operatorname{bank}(i) = (33i + c) \bmod 32 = (i + c) \bmod 32$$

Lanes 0 through 31 map to 32 distinct banks. Padding doesn't change matrix shape, output, or logical global bytes. It changes address stride inside shared memory. NVIDIA's best-practices transpose example uses same extra column to remove bank conflict.^{[4]Reference 4CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/}

Would float tile[32][34] remove the column conflict as completely as float tile[32][33]? Derive the bank pattern before answering.

Answer

No. A 34-word row gives $\operatorname{bank}(i) = (34i + c) \bmod 32 = (2i + c) \bmod 32$. Lanes 0 through 15 reach 16 distinct banks, then lanes 16 through 31 repeat those banks. Each bank in that set receives two different words, leaving a two-way conflict. A 33-word row works because its stride is one modulo 32.

Check padded result before timing or profiling it:

Compare same evidence fields across tiled and padded:

Don't write “padding is X times faster” without your hardware, software, shape, precision, timing method, baseline, and correctness record. Architecture and toolchain decide actual size of benefit, so a universal speedup would be false.

A candidate is 6% faster, but its edit both pads the tile and changes the block from 32 x 8 to 32 x 16. What can the evidence ledger conclude, and how should the experiment recover?

Answer

The run shows that the combined candidate differs, but it can't attribute the change to padding or launch geometry. Restore 32 x 8 and measure padding alone against the fixed baseline. Then hold tile width fixed and test 32 x 16 separately. Each branch still needs the same exact-output, profiler, sanitizer, and repeated-timing gates.

Occupancy is constraint, not score

Occupancy is ratio of active warps on a streaming multiprocessor (SM) to hardware maximum. More resident warps can hide latency, but highest occupancy doesn't automatically deliver highest performance. Registers, shared memory, threads per block, and barriers can limit resident blocks, while forcing resource use down can introduce spills or extra instructions.^{[4]Reference 4CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/[2]Reference 2Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Lab block has 32 x 8 = 256 threads, or eight 32-thread warps. Resource comparison stays concrete:

Small 128-byte increase may leave occupancy unchanged, or cross allocation boundary on some architecture or a larger tile design. Read actual report. Don't infer resident blocks from source alone.

Use these diagnoses:

Keep padded variant when it removes measured conflict and improves repeat timing without correctness or resource regression. Revert when counter improves but end-to-end kernel timing doesn't. Counter movement isn't product outcome by itself.

Deliberately remove barrier

Performance edits often break synchronization before they break memory bounds. Variant broken uses padded tile but omits __syncthreads() between block writes and reads. Some warps can read locations before producer warp writes them.

Run smaller matrix under tools. NVIDIA recommends memcheck before racecheck because racecheck doesn't check invalid memory addresses.^{[5]Reference 5Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/}

Expected diagnostic split:

• memcheck may report no illegal address because every index stays in allocation bounds;
• racecheck should report shared-memory read-after-write hazards around missing barrier;
• synccheck may report no misuse because kernel omitted barrier instead of executing divergent or invalid barrier;
• output checker may fail, but accidental correct output on one run doesn't clear race.

That last case is important. Scheduler timing can hide data race during ordinary run. Dynamic race evidence invalidates kernel even when sampled output happens to match.

Recover by restoring block-wide barrier, then rerun exact output, memcheck, and racecheck on padded variant:

Recovery record must connect symptom to invariant:

Fill recovery-template.md and keep it with raw reports. That artifact is stronger than “added synchronization” because reviewer can trace failure, repair, and post-recovery cost.

Decide between transpose, reduction, and scan

Shared-memory tile is right for layout reordering. It isn't universal answer for every memory-bound kernel. Ask what output dependency requires before choosing primitive.

Consider moving correctness check from CPU back to GPU. If you need only total mismatch count and maximum absolute error, use reductions. Both outputs collapse $N^2$ comparisons into one or few scalars. A scan would do extra work by producing prefix result for every comparison.

Now consider compacting indexes of mismatched elements. Each failing element needs destination position based on number of failures before it. That is prefix dependency, so scan computes offsets. A reduction can count total failures but can't assign each failure unique compacted slot.

Decision rule stays mechanical:

1. Same number of values, different layout: transpose.
2. Fewer summary values: reduction.
3. Same number of prefix states or compaction offsets: scan.

Tiling can appear inside all three implementations, but dependency graph picks algorithm. Don't start by copying transpose tile into reduction or scan kernel.

Close evidence ledger

Complete comparison only after all variants share workload and validity contract:

One row isn't complete if its “correctness” cell says only “looks right.” Keep raw CSV, .ncu-rep files, sanitizer logs, compiler resource output, environment files, source hash, and filled recovery record.

You have finished when you can defend all six claims:

• workload and timing method stayed fixed;
• baseline bound came from profiler evidence;
• each candidate changed one causal mechanism;
• output check matched transpose exactly;
• sanitizer cleared accepted kernel;
• reported performance includes conditions and distribution, not unsupported headline speedup.