GEMM From Naive CUDA to Tensor Cores

A matrix multiplication can return every right number and still waste 97% of a GPU. When you write the textbook three-loop dot product in naive CUDA, each thread spends its entire existence waiting on High Bandwidth Memory (HBM). Swapping thread coordinates fixes memory coalescing, packing scattered bytes into wide bus transactions. Tiling inputs into shared memory raises arithmetic intensity by orders of magnitude, amortizing DRAM fetches across hundreds of threads. 2D register blocking then insulates shared memory itself, feeding multiple accumulators per thread from register files. Finally, Tensor Cores replace scalar arithmetic instructions with warp-synchronous matrix multiply-accumulate operations, shifting compute ceilings into hundreds of TFLOP/s.

The previous GPU compiler lesson separated source language, generated code, and target architecture. In the kernel engineering lab, memory access, occupancy, and synchronization became tuning decisions. Finally, the profiling lesson required correctness before speed. Keep all three contracts active here.

We build one operation, $C = AB$, through five successive kernel stages. This represents the $\alpha=1$, $\beta=0$ case of general matrix-matrix multiplication (GEMM), $C \leftarrow \alpha AB + \beta C$, where we don't read the old output buffer. Matrix $A$ has shape $[M,K]$, $B$ has shape $[K,N]$, and $C$ has shape $[M,N]$. All lab buffers use row-major storage, so neighboring columns occupy neighboring memory addresses.

The complete CUDA source is available as gemm_lab.cu. Its kernels are intentionally educational. A production library has more schedules, layout transforms, epilogues, and hardware-specific paths than one chapter should hide inside a code listing.

You'll trace warp addresses, eliminate shared-memory bank conflicts with padding and swizzling, explain both shared-memory barriers, check a padded WMMA layout, and produce a correctness-gated measurement receipt. The host C++ checks and CPU index model run on any machine without a GPU. Compiling CUDA, executing kernels, running sanitizers, and collecting hardware timings remain required on your target device.

Compute one output before launching a warp

Start with matrices small enough to multiply on paper:

$$A = \begin{bmatrix}1 & 2 & 0 \\ -1 & 3 & 2\end{bmatrix}, \qquad B = \begin{bmatrix}2 & -1 \\ 1 & 4 \\ 3 & 0\end{bmatrix}$$

$A$ is $2 \times 3$ and $B$ is $3 \times 2$, so $C$ must be $2 \times 2$. Entry $C_{1,0}$ takes row 1 of $A$ and column 0 of $B$:

$$C_{1,0} = (-1)(2) + (3)(1) + (2)(3) = -2 + 3 + 6 = 7$$

Applying the same dot product to each output gives:

$$C = \begin{bmatrix}4 & 7 \\ 7 & 13\end{bmatrix}$$

GEMM performs $MNK$ multiplications and about the same number of additions. Performance tools conventionally count $2MNK$ floating-point operations (FLOPs), including one multiply and one add for each inner-loop step. Our tiny product therefore counts $2 \cdot 2 \cdot 2 \cdot 3 = 24$ FLOPs.

In naive scalar code, each thread computes a single output $C_{i,j} = \sum_{k=0}^{K-1} A_{i,k} B_{k,j}$. For each step along $k$, the thread fetches one float from $A$ (4 bytes) and one float from $B$ (4 bytes), performing one fused multiply-add (2 FLOPs). That yields an arithmetic intensity of:

$$I_{\text{naive}} = \frac{2\ \text{FLOPs}}{4\ \text{bytes} + 4\ \text{bytes}} = 0.25\ \text{FLOP/B}$$

Consider an NVIDIA A100 GPU with 2,039 GB/s HBM2e memory bandwidth and 19.5 TFLOP/s peak FP32 CUDA core throughput. At 0.25 FLOP/B, attainable performance caps out at:

$$\text{Throughput}_{\text{naive}} = 2{,}039\ \text{GB/s} \times 0.25\ \text{FLOP/B} \approx 509.8\ \text{GFLOP/s} \approx 0.51\ \text{TFLOP/s}$$

That's just 2.6% of the GPU's 19.5 TFLOP/s compute capability. Over 97% of the arithmetic hardware sits idle, starved for bytes.

The script below reproduces the hand calculation and compares two traffic models. Its first model assumes every output reloads two input values for every $k$. An optimistic second model counts each input and output element once, giving a lower bound for a cold, standalone operation rather than a claim about a real cache.

Those intensities are small because the matrices are small. For a square FP32 GEMM with $M=N=K=n$, the unique-byte model becomes:

$$I_{\text{ideal}} = \frac{2n^3}{4(n^2+n^2+n^2)} = \frac{n}{6}\ \text{FLOP/B}$$

When $n=4096$, $I_{\text{ideal}} \approx 682.7$ FLOP/B. Large square GEMMs contain abundant reuse to become compute-bound, but only if the kernel captures that reuse on chip. The roofline model bounds attainable compute rate by the lower of peak arithmetic throughput and memory bandwidth multiplied by arithmetic intensity.^{[1]Reference 1Roofline: An Insightful Visual Performance Model for Multicore Architectureshttps://doi.org/10.1145/1498765.1498785} It tells us whether reducing bytes can still move the ceiling.

Why is 0.375 FLOP/B an optimistic bound for the tiny product rather than a measured intensity?

Answer

It counts every A, B, and C element exactly once. A real kernel may fetch cache lines it only partly uses, reload evicted data, or move padding and metadata. A warm cache can also reduce DRAM traffic for one launch, so measure device traffic and record cache conditions.

A correct thread can still request the wrong addresses

The first CUDA kernel assigns one thread to one output. Each thread runs the same $K$-step dot product:

CUDA linearizes threadIdx.x first across warp lanes 0 through 31. In this strided kernel with a $16 \times 16$ block, lanes 0 through 15 own rows 0 through 15 of column 0, while lanes 16 through 31 own rows 0 through 15 of column 1. Because $A$ is stored row-major, row $r+1$ begins $K$ elements ($4K$ bytes) after row $r$.

When $K=64$ floats (256 bytes), lane 0 requests byte 0, lane 1 requests byte 256, and lane 2 requests byte 512. A hardware DRAM cache sector is 32 bytes wide (holding 8 contiguous FP32 floats). Because adjacent lanes touch addresses 256 bytes apart, each lane's request lands in a completely separate 32-byte sector. Lanes 0 through 15 force 16 distinct sector transactions for a single instruction. Stores to $C$ suffer the exact same $N$-stride penalty. The numbers are mathematically correct, but memory hardware chokes on scattered transactions.

Swap index assignments so threadIdx.x selects col:

With the coalesced kernel's $32 \times 8$ block, a full warp of 32 lanes owns one row and 32 consecutive columns. For matrix $A$, every thread in the warp reads A[row * K + k]. All 32 lanes request the exact same address, which the SM services with a single 32-byte sector broadcast. For matrix $B$, lanes 0 through 31 request 32 consecutive columns ($32 \times 4 = 128$ bytes), fitting neatly into 4 contiguous 32-byte sectors. Stores to $C$ similarly coalesce into 4 contiguous sectors.^{[2]Reference 2CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/}

Coalescing cuts bus transactions by 73% for the first warp iteration. Yet its arithmetic intensity remains stuck at 0.25 FLOP/B. Packing memory requests tightly prevents bus transaction replays, but threads still reload inputs from global memory on every step. Shared memory tackles data reuse next.

Turn global loads into shared tiles

Pick a $16 \times 16$ output tile. One threadblock loads a $16 \times 16$ tile of $A$ and a $16 \times 16$ tile of $B$ into fast, on-chip shared memory, computes 256 partial dot products, then advances 16 positions along $K$. Every loaded $A$ value feeds 16 output columns, and every loaded $B$ value feeds 16 output rows.

Two block-wide barriers form correctness boundaries:

1. Read-After-Write (RAW) barrier: Prevents any thread from reading a shared tile before all threads finish loading it from global memory.
2. Write-After-Read (WAR) barrier: Prevents fast threads from looping around and overwriting shared memory before slower threads finish consuming the current tile.

NVIDIA's classic GEMM pattern relies on this cooperative staging to slash global DRAM transactions.^{[2]Reference 2CUDA C++ Best Practices Guidehttps://docs.nvidia.com/cuda/cuda-c-best-practices-guide/}

Edge handling preserves barrier participation. Each cooperative load checks its own coordinate and writes zero when out of range. A thread outside $M$ or $N$ still loads an input needed by its valid neighbors and must participate in __syncthreads(). Early exits before barriers cause undefined behavior and GPU hangs. Only the final global store to $C$ is predicated.

For each $16 \times 16 \times 16$ tile step, the block performs $2 \cdot 16^3 = 8{,}192$ FLOPs. It loads two 256-element FP32 tiles ($2 \cdot 256 \cdot 4 = 2{,}048$ bytes) from global memory. Arithmetic intensity at the global-to-shared boundary jumps to:

$$I_{\text{shared}} = \frac{8{,}192\ \text{FLOPs}}{2{,}048\ \text{bytes}} = 4.0\ \text{FLOP/B}$$

This represents a $16\times$ reduction in global memory traffic compared to scalar code. On our A100, 4.0 FLOP/B raises the bandwidth ceiling from 0.51 TFLOP/s to $2{,}039 \times 4.0 \approx 8.16$ TFLOP/s.

An edge thread has row >= M. Why must it still reach both barriers inside the K loop?

Answer

Other threads in the block may need its B load and its barrier participation. The edge thread still performs both predicated loads, zero-fills only invalid source coordinates, reaches each block-wide barrier, and skips its final out-of-range store.

Eliminate shared-memory bank conflicts with padding and swizzling

Shared memory isn't an unconstrained cache. On modern NVIDIA architectures, shared memory is divided into 32 independent memory banks, interleaved cyclically in 4-byte (32-bit) words:

$$\text{bank} = \left(\frac{\text{byte address}}{4}\right) \bmod 32$$

Each bank can service only one 32-bit word per clock cycle. When multiple threads in a warp access different words within the same bank simultaneously, the hardware serializes the requests. A 16-way or 32-way bank conflict stalls the execution pipeline for 16 or 32 consecutive clock cycles.

Examine the inner loop access pattern in gemm_shared:

For $B$, adjacent threads in a warp have adjacent threadIdx.x values. At fixed $k$, lane $j$ reads column $j$, accessing bank $(k \cdot 16 + j) \bmod 32$. Across all 32 lanes, requests distribute across banks without conflicts.

For $A$, thread $(ty, tx)$ reads As[threadIdx.y][k]. Notice that threads with different $ty$ share the same column $k$. In a $32 \times 32$ tile where stride is 32 floats:

$$\text{bank} = (\text{row} \cdot 32 + k) \bmod 32 = k \bmod 32$$

Every thread in the warp hits the exact same bank $k \bmod 32$ with different row addresses. That triggers a catastrophic 32-way bank conflict, destroying shared-memory throughput.

Two techniques eliminate this serialization:

1. Stride padding

Add a padding float to the shared memory array declaration:

With stride 33, the bank calculation becomes:

$$\text{bank} = (\text{row} \cdot 33 + k) \bmod 32 = (\text{row} + k) \bmod 32$$

Because $\gcd(33, 32) = 1$, each consecutive row shifts the bank assignment by 1. All 32 threads land on 32 distinct banks. Bank conflicts drop to zero.

2. Bitwise XOR swizzling

While padding is simple, it wastes shared-memory capacity and breaks the 16-byte alignment required for 128-bit vector loads (float4 or uint4). Production kernels in libraries like CUTLASS use bitwise XOR swizzling instead.^{[3]Reference 3Efficient GEMM in CUDAhttps://docs.nvidia.com/cutlass/latest/media/docs/cpp/efficient_gemm.html}

Swizzling permutes the column index by XORing it with bits from the row index:

By XORing row bits into column bits, each row's columns map to a distinct permutation of banks. No memory is wasted on padding, 128-bit memory alignment is preserved, and bank conflicts vanish.

Let each thread keep more than one answer

In gemm_shared, each thread computes a single output value. For every multiply-accumulate step, the thread reads one float from As and one float from Bs: 8 bytes read from shared memory for 2 FLOPs. That gives an internal shared-memory arithmetic intensity of just 0.25 FLOP/B. The SM's shared-memory load pipelines and register ports become saturated.

Register blocking gives each thread a 2D tile of $R_M \times R_N$ outputs held in local registers. In our lab kernel, each thread computes 4 horizontal outputs ($1 \times 4$ tile):

A single float loaded from As into register a now feeds 4 independent multiply-adds. In an $8 \times 8$ register tile ($R_M = 8, R_N = 8$), a thread loads 8 values from $A$ and 8 values from $B$ (16 floats = 64 bytes), then executes $8 \times 8 = 64$ FMAs (128 FLOPs):

$$I_{\text{smem}} = \frac{2 \cdot R_M \cdot R_N}{4(R_M + R_N)} = \frac{128\ \text{FLOPs}}{64\ \text{bytes}} = 2.0\ \text{FLOP/B}$$

That represents an $8\times$ reduction in shared-memory load traffic. Registers insulate shared memory just as shared memory insulates global DRAM.

More registers aren't free. Each thread must allocate $R_M \times R_N$ accumulator registers plus temporary storage for inputs. An SM holds 65,536 32-bit registers. If each thread consumes 64 registers, an SM can host at most $65{,}536 / 64 = 1{,}024$ active threads (capping occupancy at 50% on Ampere). If register allocations spill into local memory (backed by DRAM), high latencies and cache thrashing erase all benefits.^{[3]Reference 3Efficient GEMM in CUDAhttps://docs.nvidia.com/cutlass/latest/media/docs/cpp/efficient_gemm.html} Tuning register tile sizes balances arithmetic reuse against warp occupancy.

Hand a tile to Tensor Cores

Tensor Cores are specialized execution units designed for matrix multiply-accumulate operations. NVIDIA introduced Volta Tensor Cores through the Warp Matrix Multiply Accumulate (WMMA) API in CUDA 9. Rather than having individual threads execute scalar FMAs, all 32 lanes in a warp cooperate to multiply fragments of matrices:

$$D = A \cdot B + C$$

For FP16 inputs, a standard fragment represents a $16 \times 16 \times 16$ matrix tile, accumulating into FP32 accumulators.^{[4]Reference 4Programming Tensor Cores in CUDA 9https://developer.nvidia.com/blog/programming-tensor-cores-cuda-9/}

Three rules prevent bugs when programming WMMA:

1. Full warp cooperation: All 32 lanes must participate with matching matrix shapes and types. A diverged or partially active warp produces undefined results.
2. Strict pointer and stride alignment: For FP16 inputs, load_matrix_sync requires a 32-byte-aligned base pointer and a leading dimension divisible by 8 elements (16 bytes). The FP32 output store requires 32-byte pointer alignment and a leading dimension divisible by 4 elements (16 bytes).
3. Padded fragment boundaries: Tensor Core instructions can't load partial fragments. Edges must be padded to multiples of 16 in all three dimensions ($M_p, N_p, K_p$), zero-filled, and cropped after computation.^{[5]Reference 5CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

The lab kernel uses row-major FP16 inputs, FP32 accumulation, and one warp per $16 \times 16$ tile:

Because host setup pads all dimensions to multiples of 16 ($M_p, N_p, K_p$), the grid launches only complete tiles.

Padding adds explicit arithmetic overhead. For default shapes $M=257, N=263, K=251$, WMMA executes on $272 \times 272 \times 256$:

$$\text{Overhead} = \frac{2 \cdot 272 \cdot 272 \cdot 256}{2 \cdot 257 \cdot 263 \cdot 251} - 1 = \frac{37{,}879{,}808}{33{,}930{,}682} - 1 \approx 11.6\%$$

Reported TFLOP/s must always use logical operations ($2MNK$), measuring useful work completed rather than padded instruction waste.

Layout is part of the type

A[row * K + k] assumes row-major storage. B[k * N + col] assumes row-major storage. Changing a fragment layout tag from wmma::row_major to wmma::col_major without physically transposing memory doesn't create a faster kernel; it computes a completely wrong matrix product.

Always specify and verify:

• Logical shapes $[M,K]$, $[K,N]$, $[M,N]$;
• Storage order (row-major or column-major);
• Leading dimensions and memory strides;
• Memory alignment and padding multiples;
• Input, accumulator, and output data types.

Why mixed precision demands FP32 accumulation

Tensor Cores achieve extreme throughput by multiplying reduced-precision inputs (FP16 or BF16) and accumulating sums in full FP32. Understanding why FP16 accumulation fails reveals the numerical foundation of deep learning systems.

Consider the IEEE 754 half-precision (FP16) format:

• 1 sign bit, 5 exponent bits, 10 mantissa bits;
• Machine epsilon $\epsilon_{\text{FP16}} = 2^{-10} \approx 9.77 \times 10^{-4}$;
• Dynamic range: minimum subnormal $\approx 6 \times 10^{-8}$, maximum finite value $\approx 65{,}504$.

In a transformer model with hidden dimension $K = 4096$, calculating a single dot product sums 4,096 products. If inputs are normalized with variance $\sim 1$, partial sums quickly grow toward $10^2$ or $10^3$.

When an FP16 accumulator reaches $1{,}024$ ($2^{10}$), its least significant mantissa bit represents $1{,}024 \times 2^{-10} = 1.0$. Any subsequent product smaller than 0.5 completely disappears when added (swamping)! Over a 4,096-step reduction, hundreds of small gradients or activations round to zero, causing severe gradient underflow and network divergence.

In contrast, single precision (FP32) provides:

• 8 exponent bits, 23 mantissa bits;
• Machine epsilon $\epsilon_{\text{FP32}} = 2^{-23} \approx 1.19 \times 10^{-7}$;
• Dynamic range up to $3.4 \times 10^{38}$.

Even when the running sum reaches $1{,}024$, the FP32 resolution step is $1{,}024 \times 2^{-23} \approx 1.22 \times 10^{-4}$, easily preserving tiny numerical updates across long reductions.

The two-reference testing protocol

Validating a mixed-precision kernel requires two distinct CPU references:

1. Arithmetic correctness reference: Convert FP32 inputs to FP16, convert them back to FP32, and accumulate in FP64. This reference isolates kernel arithmetic and accumulation fidelity from input quantization loss.
2. Application fidelity reference: Accumulate original unrounded FP32 inputs in FP64. This measures total end-to-end numerical change introduced by quantizing inputs to FP16.

We check kernel correctness against Reference 1 using absolute and relative error bounds:

$$|C_{\text{gpu}} - C_{\text{ref}}| \le \text{atol} + \text{rtol} \cdot |C_{\text{ref}}|$$

For FP16 WMMA with FP32 accumulation, the lab enforces $\text{atol} = 5 \times 10^{-3}$ and $\text{rtol} = 2 \times 10^{-3}$. Absolute tolerance handles values near zero, while relative tolerance scales with large outputs.

The hardware roofline across kernel stages

We can now place every kernel stage onto the hardware roofline to understand how memory hierarchy optimizations shift the performance ceiling.

On an NVIDIA A100 SXM4 GPU:

• HBM2e bandwidth: $B_{\text{mem}} = 2{,}039$ GB/s.
• Peak FP32 CUDA core throughput: $P_{\text{FP32}} = 19.5$ TFLOP/s.
• Peak FP16 Tensor Core throughput: $P_{\text{TC}} = 312$ TFLOP/s.

The roofline knee defines the minimum arithmetic intensity required to hit peak compute:

$$\text{Knee}_{\text{FP32}} = \frac{19{,}500\ \text{GFLOP/s}}{2{,}039\ \text{GB/s}} \approx 9.56\ \text{FLOP/B}$$ $$\text{Knee}_{\text{TC}} = \frac{312{,}000\ \text{GFLOP/s}}{2{,}039\ \text{GB/s}} \approx 153.0\ \text{FLOP/B}$$

Notice the critical insight: writing a naive WMMA kernel without large shared-memory tiles and multi-stage pipelining leaves Tensor Cores 97% starved for data! Reaching 312 TFLOP/s requires clearing the 153 FLOP/B roofline knee, which demands large threadblock tiles ($128 \times 256$), deep register blocking, and asynchronous prefetching.

Run the lab and keep the receipt

Save the download as assets/gemm_lab.cu relative to your working directory. Check the installed toolkit and device first. Set GPU_ARCH to an architecture supported by both the device and compiler, such as sm_80 on an Ampere GPU:

WMMA hardware support began at compute capability 7.0, but toolkit support is a separate constraint. CUDA 13.0 removed offline compilation for architectures below 7.5; a Volta sm_70 exercise needs a compatible CUDA 12.x toolchain. A binary targeting sm_80 doesn't become a Hopper-specialized TMA kernel merely by running on newer hardware.^{[6]Reference 6CUDA Toolkit 13.0 Release Noteshttps://docs.nvidia.com/cuda/archive/13.0.0/cuda-toolkit-release-notes/index.html}

The CLI rejects nonpositive or malformed dimensions, padded buffers exceeding signed 32-bit index ranges, and unsupported grid sizes. Those checks don't guarantee enough device memory. The host FP64 reference also costs $O(MNK)$: start with small correctness shapes before attempting large matrices.

Exact latency and throughput require device measurements. This is the output schema when every kernel passes:

After correctness, the program performs five warmups, records CUDA events around repeated launches in one stream, synchronizes the stop event, and reports the mean event interval per launch. It excludes allocation, transfers, FP16 conversion, padding, reference computation, and cropping. Short kernels can include stream idle gaps between host launches. Reusing the same buffers also favors warm-cache behavior. For a serious receipt, repeat whole runs and report the distribution of their means.

Use a structured row so someone else can reproduce the comparison:

Don't write 4.2x faster unless the receipt names both kernels, same shape, same device state, same timing method, same precision contract, and equivalent correctness. Compare against a tuned library such as cuBLAS or a CUTLASS profiler kernel before calling a handwritten kernel competitive.

The lab's FP16 WMMA row and FP32 scalar rows have different input precision. Their timing ratio alone isn't a same-precision optimization result. Use rounded inputs and an explicitly matched library compute mode for a controlled comparison.

Run Compute Sanitizer to verify memory and race safety:

Nsight Compute can then explain timing changes. Start with a small metric set:

This command captures the first matching launch. The basic set is a starting capture. Add sections for memory sectors, occupancy, registers, shared memory, and tensor-pipe activity as needed.^{[7]Reference 7Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Check the index model without a GPU

Download gemm_index_check.py alongside the CUDA file. It executes the two scalar ownership maps, cooperative shared loads, four-output register layout, and zero-padding/crop algebra on six integer fixtures, including $17 \times 65 \times 19$ and $73 \times 61 \times 67$. It assumes both block barriers work; it can't detect real scheduling races, floating-point error, or WMMA instruction behavior.

The CUDA download also exposes a host-only build that checks its actual argument parser, index-overflow guards, reference multiplication, NaN rejection, and tolerance comparison. It excludes every CUDA declaration:

Where asynchronous copies and TMA take over

The shared and WMMA kernels issue a load, wait, compute, and repeat. A pipelined main loop instead prepares tile $t+1$ while computing tile $t$:

The two slots must not alias. Copy completion makes a tile readable; consumer completion makes its storage reusable. An asynchronous MMA can still read shared memory after issue, so releasing a stage at instruction issue is unsafe.

NVIDIA documents three distinct hardware boundaries:

In pre-Ampere code, moving data from global DRAM into shared memory required two instructions per 16 bytes: LDG (load from global to register) and STS (store from register to shared memory). This consumed register file bandwidth and SM issue slots just to route bytes into shared memory.

Ampere introduced cp.async (hardware instruction LDGSTS), copying data directly from global memory into shared memory without touching registers. This frees registers for accumulators and lets data movement proceed in the background while math units execute compute instructions.

Hopper introduces two further leaps:

1. Warpgroup MMA (wgmma): Four warps (128 threads) execute matrix multiply-accumulate cooperatively. Unlike Volta and Ampere WMMA, which required loading input matrices from shared memory into thread registers (ldmatrix), wgmma reads matrix inputs directly from shared memory. This cuts register allocation in half.
2. Tensor Memory Accelerator (TMA): A dedicated hardware unit that copies multidimensional tensor tiles between global memory and shared memory asynchronously. A single thread issues a TMA descriptor instruction; hardware handles address calculations, strides, and out-of-bounds boundary clipping automatically.^{[5]Reference 5CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

CUTLASS 3.x and its CuTe template engine wrap these hardware primitives into composable C++ layouts (Tensor, Layout, TiledCopy, TiledMMA), automating multi-stage pipelining and warp specialization.^{[8]Reference 8CUTLASS and CuTe Documentationhttps://docs.nvidia.com/cutlass/latest/overview.html} Use the CUTLASS profiler for verified library receipts:

An async-copy kernel returns correct values but matches the synchronous kernel's timing. Why isn't the API name proof that transfer overlapped arithmetic?

Answer

A high-level copy can fall back to a synchronous path when its alignment or size contract isn't met. Even with an asynchronous instruction, too little independent computation, too few pipeline stages, or waits in the wrong place can remove overlap. Inspect the emitted instruction path and stall metrics before changing the stage count.

Diagnose the first bad signal

Every row starts from a symptom, not a favorite optimization. If correctness fails, stop timing. If timing moves without the expected counter change, revisit the hypothesis.

Keep the artifact honest

A complete lab submission contains:

• Source and exact compile command;
• Device, driver, toolkit, and target architecture;
• Odd, tiled, and long-$K$ correctness shapes;
• Reference dtype plus atol and rtol;
• Warmup, timing, synchronization, and cache-state method;
• Logical and padded operation counts;
• One profiler capture tied to a named kernel;
• A cuBLAS or CUTLASS baseline with the same input and output contract;
• Conclusions limited to measured shapes.

The progression has a stable meaning: coalescing repairs warp transactions; shared tiling captures DRAM reuse; XOR swizzling eliminates bank conflicts; register blocking increases per-thread reuse and independent accumulation; WMMA moves cooperative tiles onto Tensor Cores; and asynchronous copies with TMA overlap data movement with compute.