Day 84Module 9

in-technical-review

CUTLASS and CuTe layouts

Two template arguments separate this page's two builds of one file. One says cutlass::arch::Sm75, a 16 by 8 by 8 instruction shape and two pipeline stages. The other says cutlass::arch::Sm80, 16 by 8 by 16 and three.

The tile shape, data types, epilogue, and problem stay fixed.

The sm_75 binary runs on Turing. The sm_80 binary needs an opcode present in NVIDIA's Ampere instruction table but absent from Turing. This lesson builds both CUTLASS targets, disassembles their tensor-core instructions, and checks the architecture gate.

A layout is a shape and a stride

CuTe is the layer CUTLASS 3.0 put underneath itself, and it has one type worth learning. A Layout is a pair: a Shape and a Stride. That is all of it.

Semantically it "implements a mapping from any coordinate within the Shape to an index via the Stride" (https://github.com/NVIDIA/cutlass/blob/v4.7.1/media/docs/cpp/cute/01_layout.md , checked 2026-09-01). A row-major 256 by 128 matrix is the layout (256,128):(128,1), because coordinate (row, col) maps to row * 128 + col. Column major is the same shape with the strides swapped.

Nothing is stored and nothing is copied; a Layout is a function from coordinates to integers, and CuTe writes it Shape:Stride.

Because layouts are functions, you can compose them. A o B means A(B(c)), and the result is another layout: two index maps become one, with its own shape and stride. That closure is why CUTLASS can build a whole addressing scheme out of small pieces.

The operation that matters most for a GEMM is division. logical_divide(A, B) splits a layout into two modes: "in the first mode are all elements pointed to by B and in the second mode are all elements not pointed to by B" (https://github.com/NVIDIA/cutlass/blob/v4.7.1/media/docs/cpp/cute/02_layout_algebra.md , checked 2026-09-01). Read B as the tile.

The first mode addresses inside one tile, the second mode picks which tile. zipped_divide is the same thing with the modes gathered so that the tile is mode 0 and the tile grid is mode 1, which is the form you index with.

Apply it to a row-major 256 by 128 matrix with a 16 by 16 tile and the answer is ((16,16),(16,8)):((128,1),(2048,16)). Every number in it is derivable: 16 by 16 elements inside a tile at the matrix's own strides, 16 by 8 tiles, and a tile step of 2048 down and 16 across. Evaluate it at a coordinate and you get tileRow * 2048 + rowInTile * 128 + tileCol * 16 + colInTile, which is the address day 16 computed by hand with four multiplications and three adds.

Same arithmetic. One is four lines and carries its own shape.

Diagram: one row-major matrix, split by an algebra rather than by hand. A 256 by 128 grid on the left with its layout printed under it, an arrow across the middle labelled zipped_divide(mA, (16,16)), and on the right the same elements regrouped into a tile picture and a tile-grid picture. Band 1, the matrix: the whole grid, one element highlighted at (37, 44). Caption "(256,128):(128,1). Element (37,44) sits at index 4780." Band 2, mode 0, one tile: a 16 by 16 block with the highlighted element at (5, 12), the other 255 greyed. Caption "layout<0> = (16,16):(128,1). Inside a tile, 256 elements, at the matrix's own strides." Band 3, mode 1, the tile grid: a 16 by 8 grid of tiles with tile (2, 2) highlighted. Caption "layout<1> = (16,8):(2048,16). 128 tiles, one index each, and 2 * 2048 + 2 * 16 + 5 * 128 + 12 = 4780." Alt text: "A row-major matrix divided into sixteen-by-sixteen tiles by CuTe. The divide splits one index map into two: two hundred and fifty-six elements inside a tile, and one hundred and twenty-eight tiles, whose strides add back to the same address."

The arch tag is not one instruction

The intuition to put down is that arch::Sm80 swaps one instruction for a newer one. It changes three things you can count in the disassembly, and only one of them is the multiply.

The first is the mma shape. Turing's half-precision tensor-core instruction is 16 by 8 by 8; Ampere's is 16 by 8 by 16, twice as deep in K. Hold the warp tile at 64 by 64 by 32 and the arithmetic is forced: (64/16) * (64/8) * (32/8) is 128 mma instructions per warp per K slice on Turing, and (64/16) * (64/8) * (32/16) is 64 on Ampere.

Same work, half the instructions, because each one is twice as deep.

The second is the fill. A three-stage mainloop moves tiles from global to shared memory with cp.async, which needs compute capability 8.0, so the Turing build cannot have three stages and does not get one. CUTLASS's dispatch policy header names the two mainloops in its own comments: MainloopSm80CpAsync is "n-buffer in smem (cp.async), pipelined with registers, with predicated gmem loads", and MainloopSm70TwoStage is "2 stage pipeline through 1 stage in smem, 1 in rmem, with predicated gmem loads" (https://github.com/NVIDIA/cutlass/blob/v4.7.1/include/cutlass/gemm/dispatch_policy.hpp , checked 2026-09-01).

Through a register, or not through a register.

Day 74 wrote that difference by hand.

The third is the shared-memory bill, stage count times tile, which is why the two are not independent knobs. CUTLASS's default for Sm80 deepens the K slice to 64 and asks for three stages of a 128 by 256 tile: 147,456 bytes per block, which an A100 has and an RTX 30 or 40, capped at 99 KB per block, does not. This page holds the tile at 128 by 128 by 32 in both builds for that reason, 32 KiB at two stages and 48 KiB at three.

These differences do not establish which build is faster. They show which kernel an architecture can execute, and CUTLASS reports that limit with a status code.

Two programs: the algebra, then the GEMM

Full programs in code/day84-cutlass/cute_layouts.cu and code/day84-cutlass/cutlass_gemm.cu. CUTLASS is header only: no package, no -l flag, clone the pinned tag and point -I at its include/ directory. The README has the clone and the three build lines.

The layout program is written so its gates can fail. Two check CuTe against an answer key nobody here wrote: NVIDIA's layout algebra document publishes twelve indices for the composition of (6,2):(8,2) with (4,3):(3,1), and the program reproduces them or exits nonzero.

    auto docA = make_layout(make_shape(Int<6>{}, Int<2>{}),
                            make_stride(Int<8>{}, Int<2>{}));
    auto docB = make_layout(make_shape(Int<4>{}, Int<3>{}),
                            make_stride(Int<3>{}, Int<1>{}));
    auto docR = composition(docA, docB);

The gate that earns the day is the third. Day 16's tile address is one expression:

static int day16Index(int tileRow, int tileCol, int rowInTile, int colInTile) {
    const int row = tileRow * kTileDim + rowInTile;
    const int col = tileCol * kTileDim + colInTile;
    return row * kK + col;
}

And here is the same map, built by the algebra instead:

    auto mA = make_layout(make_shape(Int<kM>{}, Int<kK>{}),
                          make_stride(Int<kK>{}, _1{}));
    auto tiler = make_shape(Int<kTileDim>{}, Int<kTileDim>{});
    auto zd = zipped_divide(mA, tiler);

The program evaluates both at all 32,768 elements and refuses to finish if they ever disagree. Not a tautology: the two sides were written from different places, one from the CuTe document and one from day 16's kernel. A kernel then gathers the matrix through the same Layout object on the device, which costs nothing to pass because static shapes and strides make the type empty.

The GEMM program is the CUTLASS 2.x device API, which predates CuTe and still reaches furthest down the hardware matrix. Fifteen template arguments, each of them a decision the library would otherwise make for you:

    cutlass::TensorRef<ElementInput const, LayoutInputA> refA(d_a,
                                                              LayoutInputA(kK));
    cutlass::TensorRef<ElementInput const, LayoutInputB> refB(d_b,
                                                              LayoutInputB(kK));
    cutlass::TensorRef<ElementOutput const, LayoutOutput> refC(
        d_c, LayoutOutput(kN));
    cutlass::TensorRef<ElementOutput, LayoutOutput> refD(d_d, LayoutOutput(kN));

    typename Gemm::Arguments args(cutlass::gemm::GemmCoord(kM, kN, kK), refA,
                                  refB, refC, refD,
                                  {ElementCompute(1.0f), ElementCompute(0.0f)},
                                  1);  // split-k slices
    Gemm gemmOp;

The dtypes are pinned in the source, not described here: FP16 in, FP32 accumulate, FP32 out, alpha 1 and beta 0. The CPU reference reads the same half values the GPU read and sums them in double, under day 66's K-scaled tolerance, so accumulation order is the only difference the gate can be measuring. The program times nothing: day 81 is the cuBLAS comparison and owns that number, and a timing here would invite a CUTLASS-versus-cuBLAS claim from a program that never calls cuBLAS.

Results

Verified 2026-09-02 on a Tesla T4, driver 580.173.02, CUDA 12.6 V12.6.85, against CUTLASS v4.7.1 at commit cb4247394dd82148787aed73e5dc7cef33cbf862. All three builds completed. The CuTe program passed all four gates and the sm_75 GEMM passed its numerical gate. As the written contract requires, the sm_80 binary built and produced SASS, then refused to execute on the CC 7.5 T4 with exit code 1 before any allocation.

Actual sm_80 execution is outside this T4 contract, not missing evidence for it.

The exact contract was re-verified under CUDA 13.0 V13.0.88 on the same T4. All three CUTLASS v4.7.1 builds succeeded, every CuTe and sm_75 correctness gate reproduced, and the sm_80 binary again exited 1 at the expected T4 capability gate. The fresh SASS counts and shapes were identical: sm_75 had HMMA 128, LDSM 20, LDGSTS 0 and HMMA.1688.F32; sm_80 had HMMA 64, LDSM 24, LDGSTS 24 and HMMA.16816.F32.

Capture	sm_75 build	sm_80 build
builds under CUDA 12.6	pass (11.18 s)	pass (11.37 s)
runs on the T4	pass	expected refusal, exit 1
`HMMA` count in the SASS	128	64
`LDSM` count	20	24
`LDGSTS` count	0	24
`HMMA` shape	`HMMA.1688.F32`	`HMMA.16816.F32`
The third build, `cute_layouts`, also passed under CUDA 12.6 in 12.28 seconds.
Its gates reported 12/12 composition indices, 32,768/32,768 tiled indices,
256/256 divide-identity indices, and 32,768/32,768 device-gather elements.
The sm_75 GEMM placed all 262,144 outputs inside tolerance; its worst element
used 0.194 of the error budget.

Five bets, each of which a single grep can kill.

VERIFIED: LDGSTS appears in the Ampere dump and not in the Turing one. The measured counts were 24 and 0. It is in NVIDIA's Ampere and Ada instruction table as "Asynchronous Global to Shared Memcopy" and it is absent from the Turing table entirely (https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html , tables 6 and 7, checked 2026-09-01), so a nonzero count in sass-sm75.txt would mean the opcode outlived the table that documents it.
VERIFIED: HMMA and LDSM appear in both. The counts were 128 and 20 for sm_75, and 64 and 24 for sm_80. This is because Turing has tensor cores and ldmatrix, and both opcodes sit in table 6. A build with zero HMMA is a build that quietly took the SIMT path, which is the failure this dump exists to catch.
VERIFIED: the Turing dump carries twice the HMMA count of the Ampere dump. The measured ratio was exactly 128/64 = 2. The shape histograms identified HMMA.1688.F32 and HMMA.16816.F32, matching the instruction arithmetic above: 128 mma per warp per K slice against 64. If the ratio is not near 2, ptxas is emitting more than one HMMA per PTX mma, and the shape histogram in the capture says how many.
VERIFIED: cute_layouts prints zd = ((_16,_16),(_16,_8)):((_128,_1),(_2048,_16)) and passes all four gates: 12 of 12 composition indices, 32,768 of 32,768 tiled indices, 256 of 256 on the divide identity, and 32,768 of 32,768 on the device gather. The leading _ characters mark compile-time integers, so a bare 128 in that string means a stride went dynamic and the addressing moved to runtime. All four measured gates matched these counts.
VERIFIED: cutlass_gemm_sm80 on the T4 exits nonzero at the capability gate, before it allocates anything. The measured exit code was 1, and cutlass_gemm_sm75 landed its worst element at 0.194 of the tolerance budget.

Use the opcode search as the portable check. Instruction totals change with the tile shape and CUTLASS release, but architecture support for an opcode does not.

Run it yourself

Build and run cute_layouts and cutlass_gemm_sm75 on compute capability 7.5 or newer. cuobjdump can produce both disassemblies without an NVIDIA GPU. The complete programs and pinned CUTLASS dependency do not fit Compiler Explorer's single-file environment.

Running cutlass_gemm_sm80 requires compute capability 8.0 or newer. If that hardware is unavailable, you can still build and disassemble the target. GPUs with the required feature include RTX 30, 40, and 50 models, L4, A100, and H100.

Start at /setup/learn-cuda-without-a-gpu.

The recorded provider fact sheet listed an H100 SXM5 at $0.001097 per second (FACT-SHEET.md section 4, checked 2026-08-29). This price is evidence from the recorded run, not a current recommendation.

Exercise

Change ThreadblockShape in cutlass_gemm.cu from 128 by 128 by 32 to 128 by 64 by 32, leave everything else alone, rebuild the sm_75 target and dump its SASS. Before you look, write down what happens to the HMMA count in the kernel and to the shared memory the block asks for.

Time: 30 to 45 minutes, including reading the changed SASS. Submit: your prediction, the two HMMA counts, and one sentence naming which of the two numbers you got wrong and why.

Check: answered from the dump, so no GPU is needed for any of it. The build fails loudly if the shape is illegal, because CUTLASS's shape constraints are static_asserts rather than runtime checks, and the program prints its threadblock shape, warp shape and shared-memory figure in its banner so the dump and the binary cannot be confused.

Hint 1

The warp tile did not move. Work out how many warps the new block tile holds before you work out anything about instructions.

Hint 2

A kernel's `HMMA` count is per thread, not per block. Ask what each warp is now responsible for, then ask what the block as a whole is.

Solution

Shared memory halves: one operand tile went from 128 by 32 to 64 by 32, so the pair falls from 16 KiB to 12 KiB per stage. The instruction count per thread does not move at all, and that is the answer most people get wrong. The block tile shrank from 2 by 2 warps to 2 by 1, so there are 64 threads instead of 128, but each warp still owns a 64 by 64 by 32 slab and still issues its 128 mma instructions.

Fewer warps doing the same work each, not the same warps doing less.

The obvious wrong answer, that halving the tile halves the instructions, comes from reading the threadblock shape as the unit of work. It is not: the warp shape is, and the threadblock shape only decides how many warps there are. That is the whole reason CUTLASS makes you write both.

Pitfalls

The GEMM returns before it runs and prints Error Architecture Mismatch. CUTLASS checks the arch tag your kernel was compiled with against the device it was handed, and kErrorArchMismatch is documented in its own header as "CUTLASS runs on a device that it was not compiled for" (https://github.com/NVIDIA/cutlass/blob/v4.7.1/include/cutlass/cutlass.h , checked 2026-09-01). The fix is to move the -arch flag and the arch::Sm* template argument together; either one alone is a build that lies.

Error Misaligned Operand, on a problem size that looks fine. A half-precision tensor-core GEMM asks for 128-bit aligned operands, which is 128 / sizeof_bits<half_t>, eight elements. A leading dimension that is not a multiple of eight fails the check at can_implement, before any kernel launches. Pad the matrix or take the alignment down and lose the vectorized load.

You built at -arch=sm_80 and left the tag at Sm75, and it worked. Of course it did: Turing kernels run on Ampere. You measured the two-stage mainloop on hardware that has cp.async and never used it. Check the disassembly before believing any number that has an architecture in its label, which is day 46's whole point.

You counted the mma instructions in the PTX and got a different answer from the SASS. PTX is the portable half, and ptxas is free to expand one mma into more than one machine instruction or to schedule several together. Count in the SASS, the way day 46 does.

A CUTLASS 3.x tutorial refuses to compile for the target architecture. The CollectiveBuilder path, the one that picks a mainloop for you from a few tags, has specializations only for sm90 and the Blackwell families in v4.7.1: the builders directory holds sm90_gmma_builder.inl, sm100_*, sm103_* and sm120_* and nothing older (https://github.com/NVIDIA/cutlass/tree/v4.7.1/include/cutlass/gemm/collective/builders , checked 2026-09-01). Older architectures still have CuTe-native mainloops, sm70_mma_twostage.hpp and sm80_mma_multistage.hpp, but you name the dispatch policy yourself. On a T4 the 2.x device API is the path that works, and this page uses it for that reason.

Go deeper

CuTe layouts, the source for Shape:Stride: https://github.com/NVIDIA/cutlass/blob/v4.7.1/media/docs/cpp/cute/01_layout.md (checked 2026-09-01)
CuTe layout algebra, composition, complement and the divides: https://github.com/NVIDIA/cutlass/blob/v4.7.1/media/docs/cpp/cute/02_layout_algebra.md (checked 2026-09-01)
CUTLASS, "Efficient GEMM in CUDA", the hierarchy the tile shapes name: https://github.com/NVIDIA/cutlass/blob/v4.7.1/media/docs/cpp/efficient_gemm.md (checked 2026-09-01)
The dispatch policies, one struct per mainloop with its architecture on it: https://github.com/NVIDIA/cutlass/blob/v4.7.1/include/cutlass/gemm/dispatch_policy.hpp (checked 2026-09-01)
CUDA Binary Utilities, section 4.1 "Turing Instruction Set" and 4.2 "NVIDIA Ampere GPU and Ada Instruction Set", which is where an unfamiliar opcode gets looked up: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html (checked 2026-09-01)

Day 85 leaves C++ for a week: one kernel through cuda.core, CuPy, Numba and PyCUDA. CuTe has a Python front end of its own, the CuTe DSL, which the CUTLASS README calls a public beta; it is this same algebra with a different surface. Day 86 is Triton, which takes the opposite bet.

You describe the tile and it writes the swizzle you just watched CuTe write out, and it wants compute capability 8.0 for the same generation of reasons this day did.