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!

Open visualization tabs are saved automatically in browser storage, including their order, custom names, active tool, inputs and constructor/view choices. They restore after a hard refresh, including unfinished edits. Opening a shared URL adds its visualization alongside your saved tabs; refreshing that URL preserves subsequent edits. Variables remain stored separately.

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.
• Layout views share one visualization component, including the value / index / coord buttons and standalone Zoom in / Zoom out control.
• 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.
• Decode or edit inputs. Every pasteboard decode input has a folded Manual construction section. Decoding fills its editable fields; Render and Export URL use your current adjustments. MMAAtom and CopyAtom sections accept make_mma_atom and make_copy_atom URLs, including atom inputs nested inside TiledMMA/TiledCopy construction.
• Open multiple tabs so you can compare layouts or operations side by side. Double-click a tab label to rename it; your name survives rendering.
• Show Variable Table opens a persistent right-side library of named objects loaded from the pasteboard. Layout/tensor results copy as plain shape:stride text; Swizzled Layout results copy composed Sw<B,M,S> o shape:stride text; MMAAtom, CopyAtom, TiledMMA, TiledCopy, TV Layout, and Swizzle copy their encoding URLs. Preview shows result text below its button; Visualize opens a new tab named after the variable. URLs with several outputs let you choose which result to save. GEMM has type None, because it writes into D rather than returning an object.
• Share the entire Variable Table. Export table copies a compressed CVT1.… string containing all names, types, and values, and reports its exact character count. Import table from Pasteboard restores the variables alongside existing entries, skips identical entries, and adds numeric suffixes to conflicting names. Invalid imports leave the table untouched. Length depends on the data: one simple layout measured 84 characters; a sample of ten mixed variables measured 505. Sharing uses the browser's native gzip support and works without a server.
• Reuse TV Layouts in make_tiled_copy_tv with Decode from Pasteboard. URLs exported from thread/value inputs preserve those inputs exactly. Direct TV maps are accepted when a supported thread/value factorization matches every point; other maps belong in primitive make_tiled_copy.

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

Features

UTILITY

• make_smem_layout_a / b / epi — Hopper CuTeDSL shared-memory layout helpers. Choose a LayoutEnum, static tile shape, dtype, and stage count; inspect the returned composed layout, which includes its swizzle and base shape/strides. The epilogue helper also exposes optional smem_trg_shape and smem_order. The grid offers combinable value, index, and coord toggles, matching BASICS → Layout. Value shows the swizzled physical element offset; index and coord describe the selected stage’s logical grid. Register its exported URL in the Variable Table to save a Swizzled Layout, preserving the base layout and swizzle together, including the identity swizzle. Hopper is enabled; the Blackwell architecture option is disabled pending SM100 validation.

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.

• Index — Slice a layout with CuTe-style indices such as (_, _, _, 0) or ((_,1),_). Shows the exact returned layout, base offset, and original tensor offsets for the selected elements. Supports None as an alias for _, nested modes, and selecting a single element.

• 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.

gemm (MMA)

Decode a make_tiled_mma URL directly from the pasteboard, enter per-thread RMEM fragment layouts for A, B and C, and see how cute.gemm sweeps them. D is disabled and follows C. Constructor and fragment inputs use the same collapsible sections as the other MMA tabs. Below the decode button, expand Manual construction to adjust the Op, dtypes, atom layout, and permutation; decoding fills these fields. The three coverage grids show MK / NK / MN, with every cell labeled TiledMMA. Alternative View draws MK / KN / MN, with B above C/D. All K blocks are visible; the K picker appears only for rules [3]/[5], which have an explicit K mode, and highlights the A/B blocks contributing to C/D in one accumulation step. The five documented dispatch rules are listed, with rejected rule [1] struck through, the applied rule highlighted, and current warp backend support explained. The text result distinguishes the TiledMMA dimensions in scalar elements, the GEMM dimensions in TiledMMA units, and the final matmul dimensions in scalar elements. The grids show fragment repetitions, including nested modes, and do not promise a compiler issue order.

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, Index.

  • 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, eight 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.

    • make_tiled_copy_A / B — derive a TiledCopy from the selected MMA operand layout and a SMEM → RMEM CopyAtom. The Compare picker switches between SRC/DST and MMA/COPY. The left pane shows the copy address-supplying lanes or the MMA operand register slots; the right always shows the TiledCopy destination over the same logical operand tile. Show TVs / Show Warps and a shared focus box apply to both panes. Decode from Pasteboard in each constructor section imports a copied make_tiled_mma or make_copy_atom URL (or its query string). Supports the existing ldmatrix Ops and SIMT loads; dtype follows the MMA operand and unpack_bits=None. It visualizes logical ownership, without a tensor's physical SMEM addresses or access-legality analysis. Both TiledMMA and CopyAtom have a folded Manual construction section below their decode buttons; decoding fills the editable fields.

    • retile — view how a TiledCopy regroups an existing RMEM tensor. Enter its tiled TV layout, tiler and CopyAtom value count, or decode a make_tiled_copy, make_tiled_copy_tv, or make_tiled_copy_A/B URL. Enter the input register layout or decode a make_fragment_A/B/C URL. Input and result use the same scalar cells, T/V labels, logical R# words, colors, zoom and fullscreen as make_fragment. The returned layout and intermediate grouping are shown; no data moves and no new fragment is allocated. Supports two-dimensional tilers, 8/16/32-bit scalar values, and shareable retile URLs.

    • partition_S / partition_D — decode an exported make_tiled_copy, make_tiled_copy_tv, or make_tiled_copy_A/B URL, then hand that TiledCopy a tensor and one thread. Use Decode from Pasteboard, or expand the default-folded Manual construction section to choose and edit a primitive, thread/value-derived, or MMA-derived TiledCopy. Decoding fills those controls; the constructor, tiler, tiled/SRC/DST layouts, and thread/value/instruction counts remain visible as a derived summary. 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. SIMT SRC/DST layouts coincide, while ldmatrix SRC and DST retain their distinct CuTe ownership maps, including aliases on ignored address lanes. There is one tensor box and its label follows the toggle. A tiler that doesn't divide the tensor produces CuTeDSL's returned layout with a warning: 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). The result box below Render shows the complete partition_S or partition_D layout, including its strides.

    • 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.

    make_copy_atom, make_tiled_copy, make_tiled_copy_tv, and make_tiled_tma_atom 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. The badges use blue for GMEM, green for SMEM, amber for RMEM, and violet for TMEM; partition_S / partition_D shows the selected tensor's space on all three grid titles. 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. The A/B/C visualization titles in the two current warp MMA tabs show RMEM, because these Ops use register fragments. The partition_A/B/C tab accepts a tensor layout without a memory-space choice, so it does not assign a space to that tensor.

    • make_mma_atom — build one warp MMA (warp.MmaF16BF16Op, MmaTF32Op, MmaFP8Op) or Hopper WGMMA (warpgroup.MmaF16BF16Op, MmaF8Op, MmaI8Op). Register operands show thread/value ownership across 32 or 128 threads; shared-memory operands show a single-color logical tile with an SMEM badge. WGMMA supports an editable N input (illegal values are rejected), register/shared-memory A for floating point, K-major/M-major A and K-major/N-major B where supported, and mixed FP8/INT8 A/B types. INT8 currently accepts SMEM A only because CuTeDSL 4.8 rejects register A during compilation. The default view follows major modes; Alternative View always arranges A as M×K, B as K×N and C as M×N. Highlighting affects register operands only. Presets mix warp MMA and WGMMA, and Export URL / Decode from Pasteboard preserve all constructor parameters. SMEM cells are unlabeled by default; value shows logical offsets, not a particular physical address or swizzle layout.
    • make_tiled_mma — replicate warp or WGMMA atoms along M/N/K, then apply permutation_mnk. A WGMMA atom is one 128-thread warp group: (2,1,1) uses two groups, eight warps, and 256 threads. Show TVs, Show Warps, and WGMMA's Show Warp Groups expose register slots, physical warps, and cooperative groups respectively; focus follows the selected unit. SMEM operands remain uniform-color descriptor tiles. The default view follows M-major A / N-major B; Alternative View always draws MK × KN → MN. The top row shows atom replication, the bottom adds permutation/repetition boundaries. A larger permutation tile repeats work using the same groups. K replication creates partial accumulators and requires a separate reduction. Decode an MMAAtom in section 1 or a complete TiledMMA URL in section 2; editable constructors, exports, deep links and named-variable previews preserve WGMMA parameters. Downstream Copy/partition/fragment/GEMM tabs currently remain warp-only.
    • 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. The separate make_fragment tab takes this returned layout. For A/B, its register order follows the partition's strides, so switching the source tensor between row- and column-major changes the register mapping. C uses compact order regardless of those strides. The result box below Render shows the complete partition_A, partition_B, or partition_C layout, including its strides. The single TiledMMA input section contains the MMA Op and tiling controls under Manual construction, folded by default. Decode from Pasteboard imports a make_tiled_mma URL while keeping the operand, thread, tensor, and fragment selection.
    • make_fragment_A / B / C — paste an already partitioned layout, or press make_fragment below Export URL on the partition tab to pass a selected slice and thread across. The Modes to select for make_fragment box accepts CuTe coordinates such as (_, _, _, 0) or (None, None, None, 0); blank keeps the first three modes and selects index 0 from later modes. Use _ to explicitly keep all modes. The preview shows the selected layout, value count, and source offset. The returned CuTeDSL layout maps each value coordinate to a scalar register slot. The diagram uses the same Rest-position colors and coordinates (i,j) as partition_A/B/C's second grid. Each colored block expands into one cell per atom value, labeled T#V# and with its 32-bit register R# below. Repeated register labels show values packed into the same word; the diagram cannot show runtime numerical values. The thread ID is a label supplied separately because the returned layout is identical for every thread. A/B use the partition's stride order for their outer modes, while C uses compact order. This view supports the current warp MMA register Ops; Hopper SMEM descriptors and Blackwell TMEM need separate fragment renderers.

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)
?key=make_tiled_copy_ab-B-f16bf16-half_t-float-16-(2,2,1)-(32,32,16)-ldmatrix-4-0-16-warp

Local development

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

See AGENTS.md for the working conventions and docs/architecture.md for detailed architecture notes.

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. Remaining warp matrix copy Ops — make_copy_atom already includes warp.LdMatrix8x8x16bOp, warp.LdMatrix16x8x8bOp, and warp.LdMatrix16x16x8bOp, with transpose, matrix-count, and dtype controls; unpack_bits is fixed to None because it does not change the drawn layouts. LdMatrix8x16x8bOp and the StMatrix* store family remain.

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 covers three dense warp families and three dense Hopper warpgroup families. make_tiled_mma supports both scopes; downstream partition/copy/fragment/GEMM tabs currently support warp atoms only. What is left, roughly in order:

1. make_tiled_copy_C — A/B are now visualized. C needs an accumulator-to-copy view for epilogue instructions.
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. WGMMA partitioning and Blackwell tcgen05 — WGMMA atoms and tiled constructors are visualized, but descriptor fragments are not yet supported in downstream tabs. tcgen05 accumulates in TMEM and remains outside the current MMA tabs.
4. The block-scaled Ops (MmaMXF4Op, MmaMXF8Op, MmaMXF4NVF4Op, MmaMXF8F6F4Op) subclass MmaOp directly rather than WarpMmaOp and are sm_120-only — absent rather than untested.