# Finetuning Your Program #### NOTE This chapter is part of the agent-targeted Agent BLAS Performance Guide and is also published as Markdown for agent consumption, in the `markdown/` folder of the HTML documentation tree. **Scope.** For a port, walk the Low-Hanging Fruit table once as a checklist (it needs no measurement to apply), then return to the [porting workflow](performance_guide_porting.md#perf-guide-port-label): check launch geometry, measure against the parity band, and stop there. Profile only if a gap above the band remains after the geometry check, and read profiler output through `grep`, never as whole report pages. The remaining gap is usually one of: occupancy, shared memory traffic, instruction mix, or the surrounding code. Do not guess which — at this stage intuition and measurement disagree often; every section starts from something observed (a counter, a SASS listing, an occupancy number). The same rule applies to *explanations*: a hypothesis naming the cause of a gap (a technique the reference has that you lack, a feature the library cannot express) is worthless until it survives a falsification run — disable the technique on the reference and re-measure. If the reference keeps its margin without it, the hypothesis was wrong, and acting on it — or stopping because of it — would have ended the search one step before the real cause. ## Low-Hanging Fruit Pass Run this checklist immediately after [Fitting Into cuBLASDx Idioms](performance_guide_idioms.md#perf-guide-idioms-label), before the coarse parameter search. It is mechanical, it needs no measurement to *apply*, and it catches the cases where a kernel is slow simply because it is not using the API the library expects. This ordering is mandatory: searching parameters around a hand-written copy loop, a shared-memory round trip, or a generic-architecture descriptor only finds the least bad version of the wrong implementation. The “Achieving High Performance” page of the HTML documentation lists the same advice in checklist form, plus a few items on layouts and precision selection that are worth a second pass after tuning. | Check | What to look for | |----------------------------------------------------------------------------|--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| | Is this a release build? | Compile the measured artifact with both `-O3` and `-DNDEBUG` (normally via
`CMAKE_BUILD_TYPE=Release`). Record the exact compile command. Missing `-DNDEBUG` can leave
device-side checking paths in generated kernels, so an optimized build without it is not an acceptable
performance result. | | Is the accelerated cubin actually present? | For Blackwell architecture-specific descriptors, inspect the binary rather than trusting an architecture
shorthand. If `-arch=sm_100a` is lowered to generic `sm_100` by the active toolchain, use the explicit
`--generate-code=arch=compute_100a,code=sm_100a` spelling and confirm an `sm_100a` cubin before
benchmarking. | | Is the right API used? | The timed mathematical contractions must call tensor-based `execute()` rather than the pointer-based one.
The tensor API is what gives you custom layouts, suggested layouts and the register accumulator. Calling
only `cublasdx::copy` or tensor/layout helpers while implementing the contractions as handwritten FMA
loops does **not** satisfy a request to implement the computation with cuBLASDx; reject that result before
benchmarking it. | | Is `cublasdx::copy` used wherever possible? | Hand-written copy loops rarely vectorize as well. Remember `cublasdx::copy_wait()` after it. | | If not copying, could a virtual tensor be used? | An operand generated from indices (optionally combined with a small side vector) can be evaluated
straight into shared memory, so a full matrix never has to exist in global memory. | | Are suggested layouts used? | `suggest_layout_smem_a()` and friends, rather than plain `get_layout_smem_*()`.

Setting the `LeadingDimension` operator silently disables them. This is a mandatory fork in the search, not permission to skip the
suggested-layout path: measure direct leading dimensions against explicit packing into suggested layouts,
including pack/unpack work in timing. If both are legal, reporting only `get_layout_smem_*()` leaves the
layout axis open and the result is not tuned. | | Is the register API used? | The accumulator-returning `execute()` overload keeps the output in registers. | | Does each accumulator-owning thread produce multiple outputs? | Make this a shape-screening requirement for performance candidates. A tile that gives active compute
threads only one output element rarely amortizes pipeline, callback, and epilogue overhead. Scan `M`
and `N` for several output elements per accumulator-owning thread, then verify the resulting register
pressure rather than assuming that the largest tile wins.

Especially important for trivial `alpha` / `beta` cases, where a shared-memory round trip buys nothing. | | Did the tile search start from the configurator, then adapt its CTA model? | Optional, and only for a kernel whose tile is free: run `python3 example/cublasdx/performance_guide/suggest_tile.py` (needs the configurator
installed) before hand-selecting tiles and retain several of its top candidates. Its CUTLASS
model can propose a **2CTA** tile.

Translate its combined tile to a per-CTA `Size`: `256x256x64` becomes `Size<128, 256, 64>`
with `BlocksPerCluster<2, 1>` for the 2SM arm, or the same `Size` without a cluster for the 1CTA
arm; `BlocksPerCluster<2, 2>` and `BlocksPerCluster<4, 2>` are worth the same comparison.

For a new kernel, measure both and `Size<128, 256, 32>` before moving away from that family; see the
[Blackwell card](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label) for the gates.

This is a starting point, not a fixed recommendation. Fused transforms, reductions, callbacks, and
long-lived values change register and shared-memory pressure, so the best user tile can differ. Run
micro-parameter scans both before tile search and around every promising tile. | | On SM100-family, was a cluster arm measured? | `BlocksPerCluster` is a tuning axis of its own on SM 1000/1030/1070/1100: an eligible cluster enables
TMA multicast and the 2SM UTCMMA atom. Compare the 1CTA arm against `BlocksPerCluster<2, 1>`,
`BlocksPerCluster<2, 2>` and `BlocksPerCluster<4, 2>` (the operator is a hint: unsupported shapes fall
back to 1SM atoms, and pipeline creation validates cluster/grid divisibility). Keep the per-CTA
`Size` fixed across arms so only the cluster shape changes; see
`example/cublasdx/22_persistent_cluster_gemm`. | | Are bulk postprocessing functions used? | `process_and_store()`, `reduce_and_store()` and `axpby()` instead of per-element loops.

Especially important on Blackwell, where these route through the staged `TMEM` path. | | How is gating employed? | Check every `is_thread_active()` guard. One place it must **not** appear is around
`finish_accumulation()`. | | Internal or reusable accumulator? | Internal accumulation is the more portable speed-of-light path; a reusable accumulator costs an extra
synchronization point and register pressure on architectures without TMEM. | | Are the general performance operators applied? | `MaxAlignment`, `StaticBlockDim()`, `EnableInputStreaming()`. Each has conditions attached — in
particular `StaticBlockDim()` is silently ignored unless the flat block size is a multiple of 32 and the
GEMM’s thread layout spans every thread of `BlockDim`.

Treat input streaming as a measured fast path, not an unconditional tag. It forbids A/B load operations;
also validate epilogues that scatter an accumulator to shared memory, persistent/multiple-accumulator
forms, and cooperative scratch placed next to the pipeline buffer. Keep a non-streaming descriptor as a
control. On Blackwell UTCMMA the streaming descriptor needs exactly 128 threads. | | Is the pipeline thread count instruction-compatible? | Treat this as a mandatory eligibility check before the general block-size scan. On Blackwell, an
eligible UTCMMA input-streaming pipeline requires exactly **128 threads**; using 256 can
silently select a much slower implementation family.

Do not extend that rule to load-operation, reusable-accumulator, or auxiliary-shared-memory
pipelines: keep 256 as their control and validate 128 separately. On Hopper, test **128 and 256
threads**, the two supported high-performance pipeline choices.

Re-scan this axis after changing the tile because fragment ownership and register pressure change
with `M` and `N`. | | Can a leading dimension be employed? | Where the data genuinely has one. Weigh it against losing the suggested layouts. | | Are the expected instructions generated? | Check the SASS. Confirm the intended MMA and `LDSM` / `STSM` instructions actually appear (expect
`LDSM` only on `supermma` GEMMs and `STSM` only when an accumulator is stored to shared memory);
if they do not, the cause is usually a layout or accumulator mismatch rather than a tuning parameter.

If the architecture-specific instructions are missing entirely, check that the `SM` operator carries
the `arch_specific` modifier — `SM<1000, arch_specific>()` is required for `sm_100a` features
such as `tcgen05`. Depending on the toolkit, the disassembler may spell the corresponding
instructions `UTMALDG` and `UTCHMMA`, so include both families in automated SASS searches. | ## Profile First Run Nsight Compute on the kernel before changing anything — but after the launch-geometry check in the [porting workflow](performance_guide_porting.md#perf-guide-port-label), which explains most uniform gaps without a profiler. The first question is only which of the four following sections applies, and the roofline plus the warp-stall breakdown answer it: * **Stalled on \`\`MIO Throttle\`\` or \`\`Short Scoreboard\`\`** — shared memory is the bottleneck. Go to **Shared Memory Traffic**. * **Stalled on \`\`Long Scoreboard\`\`** — waiting on global memory. Go to **Overlap And Latency Hiding**. * **Low achieved occupancy with no memory stall** — go to **Occupancy And Register Pressure**. * **None of the above, and well below the roofline** — the wrong instructions are probably being emitted. Go to **Instruction Mix**. If the GEMM turns out to be a small fraction of the kernel’s time, stop and go to **Tuning The Surrounding Code** instead. ## Occupancy And Register Pressure Register use grows with `M * N`, so the register accumulator is what usually costs occupancy. In order: * shrink the tile; * drop the block dimension (each thread owns more work, fewer blocks resident); * move a large epilogue out of registers into shared memory. A low occupancy number is not by itself a finding — GEMMs hide latency with async instructions, not warp swaps. Act on it only when a stall reason points here. ## Shared Memory Traffic This section is bank conflicts (plain, unswizzled layouts can cause many times more shared-load conflicts than the suggested ones). * Confirm `suggest_layout_smem_*()` is in use *and* returned a swizzled layout — it falls back silently under `LeadingDimension` or `cublasdx::experimental::Tile`. * A custom layout beats the suggested one only when the access pattern is not the GEMM’s — a cooperative pre-pass reading by row, or one buffer serving two operand slots (idioms 04 and 12). * Suggested layouts only rearrange within the same shared-memory footprint as the default layout (or a smaller one under clusters) — they never pad. Where conflicts need padding to go away, such as a K-major FP32 GEMM, try `suggested_leading_dimension_of` instead. ## Instruction Mix Dump the SASS and confirm the intended MMA appears, along with `LDSM` / `STSM` where you expect them. Expect them only where they can occur at all: `LDSM` is generated only for `supermma` GEMMs — the `wgmma` and `tcgen05` paths feed the MMA from shared memory through descriptors, not through register fragments — and `STSM` only when an accumulator is stored to shared memory. Their absence elsewhere is not a finding. When they are missing where they should appear, the cause is almost never a tuning parameter: * No architecture-specific MMA (`wgmma`, `tcgen05`) — the `SM` operator is missing its `arch_specific` modifier, or another silent `wgmma` gate listed in the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label) is: `WithPipeline()`, `EnableInputStreaming()`, `SM == 900` exactly, 16-byte `A` / `B` / `C` alignment, a `BlockDim` of 128, 256, 512 or 1024 threads, a supported `Precision`, a `Size` divisible by the warpgroup tiling, TN for tf32 / int8 / uint8 / fp8, and global tensors of exactly the descriptor’s value types — miss any one and the library quietly leaves the `wgmma` path. * No `LDSM` / `STSM` — a suggested layout was paired with a default accumulator, or a non-identity load op changed the instruction selection on the `wgmma` or `tcgen05` path. `supermma` supports load functors with both APIs, subject to the `EnableInputStreaming()` restriction. **Plain layouts are not this cause**: they keep `LDSM` and cost bank conflicts instead, so absence of `LDSM` is not evidence of plain layouts (see **Shared Memory Traffic**). * No `TMA` — check that the kernel uses the pipeline API and that its operand layouts and alignments support TMA. A pipeline can use cooperative copies instead, including when its computation uses `supermma`. ## Arrangement And The MMA Instruction `Arrangement` is not only a statement about how your data sits in memory — it decides which MMA the library can select, and each instruction class wants something different. In the `N` / `T` shorthand used below the two letters are `op(A)` and `op(B)` against column-major storage, so `T` is row-major and `N` is column-major: `TN` is a row-major `A` with a column-major `B`. | Instruction class | Wants | Why | |---------------------------------|----------------------------------------------------------------------------------|-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| | `supermma` | `TN` — row-major `A`, column-major `B` | The default preference. Any other arrangement has to be transposed somewhere, and that cost lands
either in a load op or in an extra shared-memory round trip. | | `supermma`, **16-bit** operands | anything | `ldmatrix.trans` transposes during the shared-to-register load, so the arrangement becomes free.
Here, and for 16-bit operands on `gmma` below, the arrangement can be chosen purely to suit the
surrounding kernel. | | `supermma`, **tf32** | `TN` — row-major `A`, column-major `B` | There is no `.trans` form for 32-bit, so `ldmatrix` serves only K-major operands; the other
arrangements compile and load with plain vector loads. | | FMA, **fp32** | `NT` — column-major `A`, row-major `B` | `float` inputs never reach tensor cores ([supermma card](performance_guide_api_supermma.md#perf-guide-api-card-supermma-label));
the FMA path vectorizes its shared loads along `M` for a column-major `A` and along `N` for a row-major
`B`. | | `gmma` — Hopper `wgmma` | anything for **16-bit** operands; `TN` **only** for tf32 / int8 / uint8 / fp8 | For 16-bit operands the `wgmma` shared-memory descriptor takes either majorness. For tf32, int8,
uint8 and fp8 it is a restriction rather than a preference: an arrangement `wgmma` cannot express does
not fail to compile; the library selects an older instruction and says nothing. | | `utcmma` — Blackwell `tcgen05` | anything for 16-bit, `int8` / `uint8` and `fp8` operands; `TN` **only** for tf32 | The `tcgen05` shared-memory descriptor takes either majorness except for tf32, which is a restriction
of the same silent kind as on `gmma`. See the [Blackwell card](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label). | Which class you actually get is resolved silently, so read this table for the class your descriptor *resolves to*, not the one your architecture advertises: `gmma` requires every gate listed in the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label), `utcmma` every gate in the [Blackwell card](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label) (the pipeline path on SM 1000/1030/1070/1100 with `arch_specific` and exactly 128 threads); everything else resolves to `supermma`, or to FMA for `float` inputs. The practical consequence: tf32 wants `TN` on every path (`gmma`, `utcmma` and `supermma`’s `ldmatrix` all need K-major 32-bit operands), and `float` off the pipeline path wants `NT`. Two things follow. If your operands are 16-bit on `supermma`, `gmma` or `utcmma`, arrangement is a free variable and should be chosen to make the *surrounding* code coalesce. If not, arrangement is a real constraint — transposing in global memory beforehand is sometimes cheaper than fighting it, though that is a full extra pass and therefore something to measure rather than assume. ## Overlap And Latency Hiding Increase pipeline depth until the `Long Scoreboard` stalls stop falling, then stop. Each stage costs shared memory linearly, and past the knee it buys occupancy loss and nothing else. A depth above the number of `K` tiles is legal: the pipeline prefills only the tiles that exist, so extra stages cost shared memory without adding overlap. Tune depth per workload mode: a long single GEMM can profit from more stages than a batched GEMM with only a few `K` tiles, even with the same descriptor. After a material win, re-scan only the axes the change can couple to; an axis whose first scan showed under two percent sensitivity stays where it was. If you are not on the pipeline path at all and the kernel is global-memory bound, that is the finding: go back to [Finding Best Coarse Solution](performance_guide_coarse_solution.md#perf-guide-coarse-label) and change paths rather than tuning this one. ## Tuning The Surrounding Code Once the GEMM is tuned, most of the kernel’s time is often elsewhere. Check: the global access pattern feeding the GEMM; grid shape and wave quantization; whether a neighbouring kernel could be fused in (the idiom catalogue, [Fitting Into cuBLASDx Idioms](performance_guide_idioms.md#perf-guide-idioms-label), is the list of ways). Re-ask [Axes That Can Be Relaxed](performance_guide_problem_shape.md#perf-guide-relaxable-axes-label) — an axis fixed at the start may not be fixed any more. ## Checklist Before calling the kernel finished: 1. **Low-Hanging Fruit Pass** has been walked, top to bottom. 2. The expected MMA and `LDSM` / `STSM` instructions appear in the SASS. 3. The measurement follows the protocol in [Agent BLAS Performance Guide](blas_performance_guide.md#perf-guide-label) — warm-up runs discarded, enough measured runs to be stable, clocks fixed if the number needs to be reproducible. 4. The result has been compared against a reference point: a roofline bound, cuBLASLt on the same shape (best of its heuristic candidates), or the naive implementation you started from. 5. Any relaxed axis has its accuracy cost written down. A precision relaxation whose error was never measured is not a result.