CuTe Layout Visualizer

An interactive browser-based tool for visualizing CUTLASS CuTe layouts, thread-value (TV) layouts, and the core CuTe layout-algebra operations.

Try now: https://kainzhong.github.io/CuTe-Layout-Visualizer/

Know which address you are reading and tell if the load is vectorized / coalseced instantly!

Why not just print_latex?

CUTLASS ships cute::print_latex(...), which dumps a LaTeX snippet you then have to paste into a .tex file, compile with pdflatex, and open in a PDF viewer every time you want to look at a layout. And all it ever shows you is one static layout.

This tool runs in your browser and lets you:

• Visualize operations, not just layouts. composition, complement, logical_divide, zipped_divide, logical_product, and zipped_product each get their own tab that renders the inputs and the result side-by-side, with coloring that makes the relationship between them obvious.
• Edit inputs live. Change a shape or a stride and hit Render — no rebuild, no LaTeX toolchain, no PDF reader.
• Toggle display modes per cell between value (layout(i)), index (1D flat coord), and coord ((m,n)) to match whatever mental model you're working in.
• Share a URL to a specific visualization and have a colleague open it in one click.
• Open multiple tabs so you can compare layouts or operations side by side.

Live editing, link sharing, multi-tab support	Illustrative display, more than just print_latex

Features

Layouts

• Layout — Render any CuTe layout specified as shape:stride, including nested/hierarchical modes (e.g. ((2,4),(2,4)):((1,8),(2,16))). Auto-stride: omit the stride to default to column-major.

• TV Layout — Map a thread-value layout onto a tile to see how threads and values are distributed across a 2D grid. Supports two input methods: direct (TV_Layout, Tile) or derived-from (Thread_Layout, Value_Layout) via make_layout_tv. The tab opens pre-filled with a thr/val pair ((4,8):(8,1) × (2,2):(2,1)) and the TV layout / tile it produces, so the derivation is visible from the first screen rather than starting from a bare TV layout. Click a thread to isolate its cells. This is also where the memory-access checks live, as two fully independent collapsible sections, each with its own data layout and tensor_dtype (a GMEM tile and its SMEM staging buffer are the same tile with different strides, so they get separate inputs):

  • Check Coalesced Read (GMEM) — colors the grid by warp-wide memory issue (one color = all threads of a warp accessing the same block of vector-width consecutive vids) and labels every cell with its physical offset from the GMEM layout. The vector width is derived from the value layout and the data layout — the widest run of adjacent addresses each thread owns — so there is nothing to configure. A summary line counts the distinct 32-byte sectors each issue touches against the theoretical minimum, so "is this coalesced" gets a number, not a squint.
  • Check Bank Conflict (SMEM) — appends the 32-bank SMEM bank id to each cell, with an optional bank filter and a Swizzle<B, M, S> applied to the offset before the bank is computed.

  Both checks are properties of the (TV layout, data layout) pair rather than of any copy atom, which is why they live here rather than in the Copy tabs.

• Coordinate (TMA) layouts — The Layout tab also accepts CuTe's scaled-basis strides, k@i, which make a layout map a coordinate to a coordinate instead of a 1-D offset. That's what TMA and identity/predication tensors are built from, and it's why (3,4):(1@1,1@0) is a transpose while (4,5):(1@0,1@1) is the identity. You can paste a coordinate-tensor printout verbatim, origin and all — (2,2) o (4,4):(1@0,1@1) — and the origin offsets every cell. Cells show the output coordinate, coloured by output axis 0, so which logical mode feeds which output dimension is obvious at a glance.

• Swizzle — Visualize a CuTe Swizzle<B, M, S> as a before/after pair over a base layout. Top grid shows the raw base layout (cell = logical offset a); bottom grid shows the same coords with each cell labelled a → b where b = a ⊕ (((a >> (M+S)) & ((1<<B)-1)) << M). Bottom-cell colour is keyed to the swizzled offset b, so same colour = same post-swizzle address bucket — makes conflict-avoidance patterns visible at a glance.

Operations

Every operation tab shows the inputs and the result as linked visualizations, not just the algebra.

• Composition & Complement — Render A, B, B_complement, and composition(A, B). Toggle the complement view to see how complement(B, size(A)) fills in the remaining layout. Cells from the first tile in A are edge-highlighted in amber; the complement layout mirrors those highlights so the correspondence is immediate.
• Complement (standalone) — Render a layout and its complement against a given cotarget size. Useful for isolating what the complement operation actually produces before involving composition.
• Logical Divide — logical_divide(A, tiler) with full tile coloring: cells belonging to the same tile share a color across A and the result. Supports single-layout and multi-line (by-mode) tilers. For 2-mode tilers, the two axes use distinct accent colors (red for mode-0, deep blue for mode-1) so the row-axis and column-axis selections are visually separable.
• Zipped / Tiled / Flat Divide — zipped_divide(A, tiler) and its two reshape-only siblings tiled_divide (unpacks outer one level) and flat_divide (every mode flat). All three produce the same set of cells at the same positions; a dropdown picks which textual form to display while the visualization stays identical. Coloring matches Logical Divide so you can see that these are pure rearrangements of the same cells.
• Logical Product — logical_product(A, tiler) with slide-based coloring: A is the block being reproduced and gets color 0; each "slide" of A across the tiler produces the next tile, colored in the next shade. Supports single-layout and multi-line (by-mode) tilers.
• Zipped / Tiled / Flat Product — zipped_product(A, tiler) and its two reshape-only siblings tiled_product (unpacks outer one level) and flat_product (every mode flat). All three produce the same set of cells at the same positions; a dropdown lets you see the textual layout for each variant while the visualization stays identical. Each column is one copy of A → column k gets color k, column 0 matches A.
• local_tile — local_tile(A, tiler, coord), which is just zipped_divide followed by a slice: cut A into tiles, keep the tile modes whole, index into the "rest" with your coord. Since the point of it is picking, the visualization greys out everything the coord discarded and colours only what survives — one colour per surviving tile, so an _ in the coord reads as several tiles rather than one big selection. Also prints the base offset separately from the layout, because slicing produces both and a layout structurally cannot carry a constant. Two things it makes concrete: _ keeps a mode rather than picking ((None, None) selects no tile at all — it exposes the tile indices as modes, which is what the TMA flow does before tma_partition), and mode 0 is unpacked one level, not flattened — ((2,2),(4,2)) comes back as two modes that are each still tuples, and only a tile whose sub-modes are already scalars looks flat. Either way mode 0 stops being one mode, which is why TMA code has to call group_modes(x, 0, 2) to fold back what local_tile just unfolded.
• Blocked Product — blocked_product(A, tiler), the rank-preserving cousin of logical_product: each output axis carries (block_i, tile_i) merged, so copies of A are laid down as contiguous sub-blocks of a bigger matrix (grid is size(A_i) * size(tiler_i) per axis). This is what you want when building a matrix tile from a per-thread block and a thread layout — it's also the primitive behind tile_to_shape. A gets color 0, and each block-copy at tile position (t0, t1) gets color t0 + t1 * size(tiler[0]).
• Raked Product — raked_product(A, tiler), the interleaved twin of blocked_product. Same set of cells, same 2D shape, but the zip order is reversed — tile-mode first, block-mode second — so cells of a single copy of A are scattered across the output tile at stride size(tiler_i) along each axis instead of clumped into a contiguous sub-block. This is the primitive behind make_layout_tv: it's why each thread's values are spread across the tile for coalesced memory access. Same coloring scheme as Blocked Product — compare the two tabs with identical inputs to see the scattered-vs-clumped difference at a glance.

Workspace

• Scoped navigation — Tabs are grouped into scopes so the tab bar doesn't turn into a wall of buttons as more features are added. The current scopes are:

  • Basics (blue) — Layout, TV Layout, Swizzle.

  • Layout Operations (purple) — Composition, Complement, Logical Divide, Zipped / Tiled / Flat Divide, Logical Product, Zipped / Tiled / Flat Product, Blocked Product, Raked Product.

  • Copy (emerald) — the copy-construction pipeline, six tabs mirroring CuTe's own layering.

    • make_copy_atom — build one Copy_Atom. Pick the Op (CopyUniversalOp or cpasync.CopyG2SOp; neither takes constructor parameters), then give make_copy_atom its tensor_dtype and num_bits_per_copy. Shows the one-thread / N-contiguous-value shape a single instruction moves.
    • make_tiled_copy — the primitive. You supply layout_tv and Tiler_MN yourself. (They need not have matching shapes — layout_tv is (num_threads, num_values), Tiler_MN is the (M, N) tile.)
    • make_tiled_copy_tv — the derived form. Give it a thr_layout × val_layout and it runs raked_product → right_inverse, prints layout_mn / Tiler_MN / layout_tv, and hands them to the above.
    • partition_S / partition_D — hand that TiledCopy a tensor and one thread. tiled_copy.get_slice(t).partition_S(gS) is what a kernel actually calls, and it is where the three levels of repetition finally line up against real data. The tensor is not pre-divided — tidfrg_S runs zipped_divide(tensor, Tiler_MN) itself, and any tensor mode past the tiler's rank (a GEMM's k) simply rides along and multiplies Rest. The picture is three grids, one per level of repetition, which is lossless because a tile's partition doesn't depend on which tile it is: (1) one Tiler_MN tile with the (t, v) that owns each element — a thread's work inside the tiler; (2) one cell per tile over the (M, N) plane, each labelled with the coordinate you'd type at tSgS(_, i, j) and the tile's origin; and (3) one cell per tensor mode the tiler never touched — the k-loop and the pipeline stage — a strip when there's one (drawn as a row, since a k-loop reads horizontally), a grid when there are two, hidden when the tensor is exactly the plane. That last split matters because Rest really is two different things wearing one mode: the part zipped_divide produced from the tiled modes, and the untiled modes carried through verbatim. Each grid prints, in blue, the layout it actually draws — mode 0 of zipped_divide(tensor, Tiler_MN), then mode 1 split at the tiler's rank — so the three headers are a partition of that layout and reading them top to bottom is the derivation. A 128×256 tensor is 1024 + 32 cells this way instead of 32768 that repeat the same picture 32 times. Tile origins are evaluated, not assumed, so a Tiler mode with a stride (16:2) shows its tiles genuinely interleaved — starting one column over rather than a tile-width away. The tab takes a rank-2 tiler and at most 2 untiled modes; both are limits of the grids rather than of CuTe, and both refusals say so. For every Op offered here ValLayoutSrc == ValLayoutDst == ValLayoutRef, so S and D differ only in which tensor they are given; there is one tensor box and its label follows the toggle. A tiler that doesn't divide the tensor is refused: C++ catches it in shape_div at compile time, but CuTeDSL returns a layout that reads off the end (checked — the same layout this tool derives, covering 388 positions of a 384-element tensor).
    • make_tiled_tma_atom — the TMA path, which skips threads entirely. Give it a GMEM tensor, an SMEM layout, a swizzle (a picker, since CUtensorMapSwizzle is a closed enum — each option labelled with both the Sw<B,M,S> CuTe prints and the byte width you think in) and a CTA tiler; it derives the TMA box — you never specify it — and shows three things: the box laid over the CTA tile against its SMEM destination, the literal cuTensorMapEncodeTiled argument list, and the coordinate tensor the function returns.

    The first four tabs render as a SRC → DST pair: a SRC / DST / BOTH toggle above the diagram, each pane titled with its memory space (GMEM / SMEM / RMEM / TMEM) and a small → between them. In BOTH mode the two grids sit side by side at half width, same aspect ratio. The memory movement is a constrained picker in section 0, not a free choice: it lists only the pairings the selected Op supports, and the pane titles follow it. cpasync.CopyG2SOp offers exactly one (GMEM→SMEM, so the picker is disabled); CopyUniversalOp offers the six cross-space pairs over GMEM/SMEM/RMEM — TMEM is excluded because it isn't thread-addressable and needs a tcgen05 Op. For both current Ops ValLayoutSrc == ValLayoutDst, so the panes match; they diverge for shuffling atoms like ldmatrix.

    TMA is the odd one out and deliberately so: one thread issues the instruction with a logical coordinate and the TMA unit does address generation, bounds handling and the swizzled SMEM write, so there is no TV layout and no per-thread view. What replaces it is the box. Presets run from a 64-cell tile up to a full 64x64 GEMM stage; the small ones use wide element types on purpose, because a swizzle does nothing until the tile spans 8 rows of 128 B — 8x64 in half_t, but only 8x8 in uint128_t. The DST pane reports how many cells the swizzle actually moved, so "0 moved" tells you the tile is too small rather than leaving you to wonder. Two more things the tab makes visible that are otherwise hard to see: the swizzle triple CuTe prints (Sw<3,4,3>) is in bytes, so it is shown alongside its element-grid equivalent (Sw<3,3,3> for half_t) — the form the Swizzle tab takes; and the TMA constraints are host-side assert()s that a release build removes, so violations (majorness mismatch, box extent > 256, unaligned strides, a box row wider than the swizzle) are reported inline next to the picture that caused them rather than replacing it with an error.

    The two tiled tabs draw how the TiledCopy covers one tile — color per thread, brightness per atom invocation — and run the checks CuTe documents but never enforces: that layout_tv fills its tiler, that thr/val layouts are compact, and that the atom's values land on a stride-1 run of one tile axis. A stride-0 mode in layout_tv is recognised as deliberate broadcast (several threads reading one element, as make_tiled_copy_A produces) rather than flagged as overlap.

  • MMA (amber) — the matrix-multiply side, mirroring the same atom-then-tile-then-partition layering as Copy.

    • make_mma_atom — build one warp-level MMA Atom (cute.nvgpu.warp.MmaF16BF16Op / MmaTF32Op / MmaFP8Op) and draw its three operand fragments. This is the one structural difference from a Copy Atom: two TV layouts over one tile become three over three — A over (M,K), B over (N,K), C over (M,N) — so A and B sit side by side sharing the K axis while C spans the width below them. All three are bijections, so every cell is owned by exactly one lane. The tab is a lookup table rather than a derivation, because CUTLASS defines these as hand-written MMA_Traits specializations: an MMA atom is the PTX register signature transcribed into layout form. Every entry is diffed against CuTeDSL, and the corpus is exhaustive over the parameter domains — which is what pins the tab's claim that the layouts depend only on the Op family and K, never on ab_dtype or acc_dtype. A Highlight thread box dims all three grids at once, so one lane's 8 cells of A, 4 of B and 4 of C read as the single fact they are. An Alternative View toggle (shared with make_tiled_mma) re-lays the grids as the picture a matrix multiply is normally drawn as — empty top-left, A bottom-left, B rotated to K×N top-right, C bottom-right — so A's K axis lines up with B's and B and C share their N axis. B is only shown transposed; the TV mapping is the same one, checked cell for cell in the tests.
    • make_tiled_mma — replicate that atom across warps. Two rows of grids, because the two arguments do genuinely different things: atom_layout_mnk says how many warps there are along M, N and K (top row, one atom-sized area each), and permutation_mnk says what tile each mode is really over (bottom row, where only the first copy of the warp pattern is coloured and red lines mark where the copies meet). A Show TVs / Show Warps toggle switches the cells between make_mma_atom's T/V labels and warp names, and one focus box below it drives all six grids — labelled Warp id or Thread ID to match the toggle, since the question is always "which unit am I looking at" and only the unit changes. Leave it blank and every warp (or thread) that touches a cell is stacked in it, one per line; type an id and only that unit's region stays coloured — same hue in A, B and C — with everything else greyed and unlabelled. In the bottom row the copies that unit also lands in are drawn in its hue at reduced brightness rather than greyed, because the warp pattern really does repeat there. An id that isn't a whole number, or one past the last warp/thread, is reported as an error rather than ignored. The warp view is where the tiling's asymmetry becomes visible: an A cell carries W0 and W2 because A does not depend on N and both N-warps read it, a B cell carries W0/W1 for the mirror reason, and a C cell carries the K-warps because they are not readers at all — each holds a partial sum of the same accumulator, which a K-split TiledMMA must reduce afterwards. Past four warps a cell collapses to one compact line (ΣW0..W7 for C, plain W0..W7 for A and B, since nothing accumulates into those).
    • partition_A / B / C — hand that TiledMMA a tensor and one thread. tiled_mma.get_slice(t).partition_A(gA) is what a kernel calls, and it is the MMA twin of partition_S / partition_D: the same three grids, meaning the same things — a thread's work in one tile, that tile over the operand's plane (M×K for A, N×K for B, M×N for C), and that plane over the untiled modes. The three grids' blue headers are mode 0, modes 1–2 and modes 3+ of the returned layout and concatenate back to it, so grid 2 always shows the Rest the call actually returns. One thing genuinely differs from the copy side, and the tab is built around it: permutation_mnk folds into those Rest modes. A copy's Tiler_MN gives one Rest entry per tile; an MMA's permutation doesn't, because the warp pattern can repeat inside the permuted tile — so the unit grid 1 draws is not the TiledMMA's tile but the block one Rest position covers, gathered from the derivation rather than assumed. With permutation_mnk = (32,32,16) a B tile is 32×16 while a Rest block is 16×16, and grid 2 shows 8×2. A permuting layout like ((2,32):(32,1),32,16) does one more thing: it interleaves the rows, so Rest mode 0 comes back a tuple of (repeat inside the tile, blocks across the tensor) and the block's rows are no longer consecutive in the tensor — both reported in the headers. Every form is a preset and checked against CuTeDSL. These take the already partitioned tensor on purpose: the fragment copies the partition's mode order so that a linear walk of the registers is a monotone walk of the source and the load vectorizes. Switch the A tensor between row- and column-major and grid 4's strides flip from 24, 8 to 8, 16. make_fragment_C is the exception — an accumulator is never read in the order it was partitioned, so it is always plain compact.
    • partition_fragment_A / B / C — what that thread then holds in registers. thr_mma.partition_fragment_A(sA) is the call a kernel writes (sgemm_sm80.cu:160) and is defined as make_fragment_A(partition_A(sA)), so this tab is the previous one plus a step and shares its inputs. The fragment has the partition's shape and compact strides, so the register count is fixed — what isn't is the order, and that's the whole reason make_fragment_A takes an already-partitioned tensor rather than a shape: it copies the partition's mode order so a linear walk of the registers is a monotone walk of the source and the load vectorizes. Two grids: which register holds which slot, and the registers in index order coloured by run of consecutive source elements. Switch the A presets between row- and column-major and the strides go from 24, 8 to 8, 16 and the runs from 2 to 1 — the same instruction dropping from a 32-bit load to a scalar one. make_fragment_C is the exception and ignores the order entirely; CUTLASS says so twice, once by providing a static partition_fragment_C(mma, shapeMN) that needs only a shape while A and B carry "should not be used in a static context".

  Click a scope at the top of the nav card to swap in its tabs. The active scope has a color accent (left stripe + active-tab highlight) so you always know which section you're in. Deep-link URLs auto-flip to the right scope. Scopes are designed to be extended — MMA was added this way, and further groups can be too without cluttering the existing ones.

• Cmd+Enter to render — ⌘↵ on macOS, Ctrl+↵ elsewhere, renders whichever tab is visible without scrolling down to the button. Works from inside any input, including the multi-line tiler boxes. The hint under each Render button shows the right key for your platform.

• Multiple tabs — Open several independent workspaces side by side. Each tab is fully self-contained.

• Shareable URLs — Every operation has an "Export URL" button that copies a deep link to the current visualization. Paste it into chat or a doc and the recipient lands on the same view.

• Zoom — Click "Zoom in" on any panel to fit by the shortest side (useful for very wide or very tall layouts).

• Everything is drawn on load — every tab renders its default inputs when the page opens, so switching tabs shows a working example immediately rather than an empty box. Every shipped default is a valid configuration.

• Presets — Built-in examples per tab covering the common patterns.

• Rank warning — Layouts with outer rank > 2 still render (flattened to 2D), but the tool surfaces a warning so you know the structure is being collapsed.

Usage

Enter a CuTe layout string using the standard notation:

(shape_0, shape_1):(stride_0, stride_1)

Nested shapes and strides are supported:

((2,4),(2,4)):((1,8),(2,16))

Pick the tab that matches what you want to visualize, fill in the inputs, and click Render. For operation tabs, the result is computed client-side using a JavaScript port of python/pycute/layout.py — the same algebra CUTLASS itself uses, just rendered in your browser instead of compiled into a PDF.

URL format

The ?key=... query parameter deep-links to a specific visualization:

?key=layout-(10,10):(1,10)
?key=tv-1-(32,4):(1,32)-(8,16)
?key=tv-2-(2,3):(3,1)-(2,2):(2,1)
?key=composition-(4,4):(4,1)-(2,2):(1,2)
?key=complement-(2,2):(1,2)-(4,4):(1,4)
?key=logical_divide-(12,32):(32,1)-3:1\n8:1
?key=zipped_divide-(12,32):(32,1)-3:1\n8:1
?key=logical_product-(2,2):(1,2)-(2,2):(1,2)
?key=zipped_product-(2,2):(1,2)-(2,2):(1,2)
?key=blocked_product-(2,2):(1,2)-(3,3):(1,3)
?key=raked_product-(2,2):(1,2)-(3,3):(1,3)
?key=make_mma_atom-f16bf16-half_t-float-16
?key=make_tiled_mma-f16bf16-half_t-float-16-(2, 2, 1)-(32, 32, 16)

Local development

No build step, no server. Just open index.html in your browser — everything is plain HTML/CSS/JS with zero dependencies.

See CLAUDE.md for architecture notes (file layout, adding a new tab, input conventions, URL scheme).

TODO

Copy Atoms

CopyUniversalOp and cpasync.CopyG2SOp are covered by a single merged tab (their Copy_Traits are byte-identical, so one visualization suffices). Four more tabs, each with dropdown-driven variants, cover everything non-trivial that's left. Rough order by implementation complexity:

1. ldmatrix / stmatrix (warp copy) — single tab with:

   • Direction toggle: load (ldmatrix, src = shuffled smem) vs store (stmatrix, dst = shuffled smem).
   • Transpose picker: N (no transpose) / T (transpose).
   • Count picker: x1 / x2 / x4 / x8.
   • Dtype picker: u32 (SM75/SM90), u16 (SM75/SM90), u8 / sub-byte (SM100 additions).

   All SM75 LDSM, SM90 STSM, and SM100 LDSM/STSM variants fold into this one tab. Same pipeline as the existing Copy tabs — direct extension, easiest to build first.

2. TMA bulk tensor (cpasync.CopyBulk*) — partly built: the make_tiled_tma_atom tab covers CopyBulkTensorTileG2SOp (plain .tile load, num_multicast = 1) over a flat rank-2 tensor. Still to come, each a clean extension of the same pipeline:

   • Variant picker: LOAD_MULTICAST (num_multicast > 1 truncates the box), STORE, REDUCE_ADD. Optionally the non-tensor BULK_COPY_G2S / BULK_COPY_S2G.
   • CTA picker: 1-CTA (SM90) vs 2-CTA (SM100 SM100_TMA_2SM_LOAD*; ThrID = 2).
   • im2col mode, gather4 / scatter4, internal_type recasts, rank > 2 tensors.
   • tma_partition, which splits a tile into (TMA, REST) when the box doesn't cover it in one instruction (the tab already reports the instruction count).

3. tcgen05 TMEM load / store — single tab collapses all ~160 SM100_TMEM_LOAD_* / SM100_TMEM_STORE_* variants:

   • Direction toggle: load (TMEM → regs) vs store (regs → TMEM).
   • DP picker: 16dp / 32dp (rows per warp).
   • Width picker: 64b / 128b / 256b.
   • Repeat picker: 1 / 2 / 4 / 8 / 16 / 32 / 64.
   • _16b packing toggle.

   Most combinatorial space of the four. Needs the shared SVG builder (buildColoredLayoutSVG and the tile-lookup helpers) to generalize beyond the rank-2 (1, elements) atom-val-layout assumption — TMEM atoms have rank ≥ 3 atom layouts like (tid, (dp, bit, rep)).

4. UTCCP (tcgen05 multicast into TMEM) — single tab, smallest family:

   • Shape picker: 128dp{256,128}bit, 4dp256bit, 4x32dp128bit, 2x64dp128bitlw{0213,0123} (4-5 options).
   • 1-CTA / 2-CTA toggle.

   Structurally distinct from TMA and from TMEM load/store, so a separate tab rather than merging.

Coverage rationale: the Copy_Traits families with trivial layouts (ThrID = Layout<_1>, SrcLayout = DstLayout = Layout<Shape<_1, bits>>) — all four SM80_CP_ASYNC_* variants, SM75_U32x1_MOVM_T, SM100_LOAD/STORE_256bit_CACHE_NOALLOCATION — are pixel-identical to UniversalCopy and don't need their own tabs. They fold into the existing CopyUniversalOp / cpasync.CopyG2SOp coverage.

MMA Atoms

make_mma_atom and make_tiled_mma cover the three dense cute.nvgpu.warp WarpMmaOp subclasses. What is left, roughly in order:

1. make_tiled_copy_A / _B / _C — where the two scopes finally meet: a TiledMMA's tv_layout_A is exactly what these hand a copy as its layout_tv, and the broadcast make_tiled_mma draws is what makes size(layout_tv) a multiple of size(Tiler_MN) there.
2. MmaF16BF16SparseOp — traces fine, but MmaAtom exposes no metadata (E) layout and its A layout is over the logical (M,K) tile, so the 2:4 compression that makes it sparse would be invisible in the only thing the tab draws. Needs a metadata view first.
3. Hopper wgmma / Blackwell tcgen05 — a different shape of problem, not a bigger version of this one: wgmma reads A/B from SMEM through descriptors and tcgen05 accumulates in TMEM, which is not thread-addressable, so the "three register fragments" reading stops holding.
4. The block-scaled Ops (MmaMXF4Op, MmaMXF8Op, MmaMXF4NVF4Op, MmaMXF8F6F4Op) subclass MmaOp directly rather than WarpMmaOp and are sm_120-only — absent rather than untested.