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: 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, 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 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, 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<Description, Arch> 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 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); 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.

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, utcmma every gate in the Blackwell card (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 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, is the list of ways). Re-ask Axes That Can Be Relaxed — 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 — 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.