Finding Best Coarse Solution#

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. This chapter searches the axes the application leaves free, including tile M/N/K, block dimension and pipeline depth. For a port, identify the actual constraints using the porting workflow and sweep the remaining axes within its variant budget. The stopping criterion is parity with the reference.

Goal: a kernel within a modest factor of achievable performance, in as few compile-and-measure cycles as possible. The work is elimination — each measurement should rule out a family of configurations, not move a number by two percent; anything worth only a few percent belongs in Finetuning Your Program.

Best Idiom Search#

One datapoint per axis. For each idiom axis your kernel touches, measure its variants in isolation and note the winner.
Construct the hypothesis. The per-axis winners give the combination you expect to be best.
Measure the hypothesis. Combinations frequently do not behave like the sum of their parts.
Perturb one axis at a time, and read the results:
- Moved — the axis matters and its isolated winner was context-dependent: adopt the better value, re-perturb the axes it interacts with (tile with block dimension, depth with shared memory).
- Stable and good — settled.
- Stable and bad — back out half the choices and re-measure; bisecting isolates the responsible subset cheaply.

This is cheap because the variants exist as separate compilable programs (a datapoint is a build and a run), and while axes do interact — hence measuring the hypothesis — a bad result on one rarely invalidates the winners on others.

Building And Running The Variants#

Built by default with the examples (53 variant programs; an idiom’s reference_kernels.cu, where present, is compiled once and shared by its variants):

cmake -DBUILD_EXAMPLE=ON \
      -DCUBLASDX_EXAMPLES_ENABLE_PERFORMANCE_GUIDE=ON \
      -DCUBLASDX_CUDA_ARCHITECTURES=<your arch, e.g. 90a> \
      -DCMAKE_BUILD_TYPE=Release ..

make -j10 guide_decoupled_precision_inputs_gmem_to_smem   # one variant
make -j10 guide_01_decoupled_precision_inputs             # all variants of one idiom (the search granularity)
make -j10 cublasdx_performance_guide                      # everything

Binaries land under <build>/example/cublasdx/performance_guide/<idiom>/<variant> and run with no arguments. The sources live under example/cublasdx/performance_guide/ in the source tree and in the installed examples.

Warning

Always build these Release with -O3 -DNDEBUG — debug measurements tell you nothing (Agent BLAS Performance Guide).

The Variants#

One directory per idiom under example/cublasdx/performance_guide/ (part of the installed examples); one .cu file per solution variant, directly in the idiom directory; CMake target guide_<variant>. Each variant’s header comment states what it isolates. Do not write a kernel from scratch — copy the nearest variant and change it.

Every binary runs up to three modes, and every results row is keyed by one: batched (64^3, one block per problem — the shape of any many-identical-small-GEMMs workload), pipelined single (one 8192^3 GEMM tiled across the grid) and pipelined batched (256^3 per batch, tiled by the pipeline).

A mode an idiom cannot express reports a concrete reason instead of a number. The orderings usually transfer across architectures and eliminated memory passes are eliminated everywhere, but absolute times and shared-memory budgets do not — re-measure on yours.

<variant>.cu — the variant’s kernels, cuBLASDx calls and executable entry point.
scenario.hpp — the idiom’s BLAS descriptors and settings (sizes, constants, config aliases).
host.hpp — launch details: reference arms, cuBLASLt usage, validation and timing call sites, launch functions.
reference_kernels.cu — reference kernels shared by the variants of an idiom, where present.
common/guide_harness.hpp and common/guide_run_*.hpp — shared allocation, launch, validation and timing helpers.
configs/problem.hpp — shared workload sizes, tile sizes, data types and measurement settings.
configs/cublasdx_descriptor.hpp — shared GEMM descriptors; configs/config.hpp includes this file.

To add a variant, copy the nearest .cu file in the idiom directory, rename it, adapt its includes, and register it with add_guide_variant(<idiom_dir> <variant_name>) in the guide’s CMakeLists.txt.

Every variant is implemented and validates against a cuBLASLt + thrust reference computing the same arithmetic. The harness enforces no PASS, no time — a variant that fails validation prints no timing.

Pick The Execution Path#

Unsure: start on the register path — least setup, what the suggested layouts and accumulator are tuned around, and the cheapest to migrate later.

Path	Choose it when	What it costs
Shared memory output `execute(alpha, a, b, beta, c)`	`C` genuinely lives in shared memory — a second GEMM consumes it, or the epilogue needs cross-thread order.	A store, a barrier and a reload a register epilogue would not pay; shared memory for `C` on top of `A` and `B`.
Register output `execute(a, b)` returning an accumulator	The default. Epilogue is elementwise or a per-fragment reduction; `alpha` / `beta` trivial or foldable.	Register pressure grows with `M * N`, so a large tile can cost occupancy. Nothing else.
Pipeline API `suggest_pipeline(...)`	Operands come from global memory and `K` spans several tiles — most device-wide GEMMs.	The most setup: host-side pipeline object, device handle into the kernel, epilogue as a lambda. The only path to `TMA`, `wgmma`, `tcgen05`.

Path

Choose it when

What it costs

Shared memory output

execute(alpha, a, b, beta, c)

C genuinely lives in shared memory — a second GEMM consumes it, or the epilogue needs cross-thread order.

A store, a barrier and a reload a register epilogue would not pay; shared memory for C on top of A and B.

execute(a, b) returning an accumulator

The default. Epilogue is elementwise or a per-fragment reduction; alpha / beta trivial or foldable.

Pipeline API

suggest_pipeline(...)

Operands come from global memory and K spans several tiles — most device-wide GEMMs.

The most setup: host-side pipeline object, device handle into the kernel, epilogue as a lambda. The only path to TMA, wgmma, tcgen05.

The paths are exclusive per call, not per kernel: a pipelined GEMM still ends in an accumulator and still chooses a register or shared-memory epilogue.
Under WithPipeline(), asking the descriptor for an accumulator is a compile error — take it from the tile pipeline.
A pipeline on a single-K-tile problem is strictly overhead.

Sweep The Few Parameters That Matter#

Trust the tuning-database defaults first: omitted BlockDim (BLAS::suggested_block_dim), suggested layouts, the MMA choice, and suggest_pipeline without a depth were all measured — an early override replaces a measurement with a guess. The exception is operators that describe the problem: Arrangement must reflect how the data actually sits in memory; that default is not a tuned one.

Parameter	Range worth trying	Notes
Tile shape, `Size<M, N, K>`	Powers of two, `32`–`256` per dimension	Highest leverage. Bigger tiles: more reuse, more registers (cost grows with `M * N`).
`BlockDim`	`128`, `256` as starting points; measure larger blocks when the workload warrants them	Only if the surrounding kernel does not fix it. Sweep to confirm `suggested_block_dim`, not to discover.
Pipeline depth	`2`–`6`, capped at `16`	Stages cost shared memory linearly and stop paying once loads are hidden; there is usually a clear knee. A depth above the `K`-tile count is legal, but only the tiles that exist are prefetched, so extra stages cost shared memory without adding overlap.
Precision	Axes That Can Be Relaxed	An algorithmic decision, typically worth more than everything else combined. Measure the accuracy given up.

Choosing candidates#

Start from the tuning-database defaults, keep a ledger of every rebuilt variant, and stop when the remaining axes move the result by less than the measurement noise.
Tile, block dimension and pipeline depth interact: when the tile is free, retain the top few tile candidates and measure the small cross-product of coupled axes around each before rejecting a family.
If the output format constrains some tile dimensions, keep only those dimensions fixed. Sweep the remaining tile dimensions, block dimension and depth when a measurable gap justifies it.
Stop the port’s coarse search once the kernel is within the parity band of its reference. Do not re-scan axes after a change that cannot couple to them.

Epilogue resources#

An epilogue change couples to depth and grid mainly through resources: one that leaves shared-memory and register use alone normally moves neither, while one that adds staging or scratch shares the block’s shared-memory budget with the pipeline stages and can cost a stage or occupancy.
suggest_pipeline without a depth sizes the stages for the whole budget and knows nothing about your scratch, so once scratch exists pass an explicit depth and check the total against the limit.

Generating variants#

Tile and block dimension are template arguments, so a scan is a rebuild per point (tens of seconds each for a pipeline kernel) — generate the cases with a Python script and use a partial descriptor summed with the varying operator (Finding Problem Shape).

Tile configurator#

A first tile before measured searching — run this Python configurator explicitly rather than copying the descriptor from an example: nvMatmulHeuristics (optional, not part of cuBLASDx, needs pip network access — when offline, or when porting a kernel whose tile is fixed, skip it and keep the reference’s tile) answers from a performance model in about a second:

pip install nvidia-matmul-heuristics
python3 example/cublasdx/performance_guide/suggest_tile.py --m 8192 --n 8192 --k 8192 --gpu RTX_PRO_6000

Prints a Size, a BlockDim and a stage count for suggest_pipeline; --list-gpus / --list-precisions enumerate supported values.
On B200, check whether the model’s candidate assumes a 2CTA instruction. 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. Do not halve N. Also measure BlocksPerCluster<2, 2> and BlocksPerCluster<4, 2>: clusters are a tuning axis of their own (TMA multicast and the 2SM atom), and the operator is a hint, so unsupported shapes fall back cleanly.

For a new kernel, measure both arms and Size<128, 256, 32>; the Blackwell card lists the gates. For a port, keep the reference’s cluster structure and the porting workflow’s variant budget.
The knobs interact and decide most of the distance to the hardware: move tile, stage count and block dimension together rather than one in isolation. A fused kernel shifts the optimum, so repeat the search for it instead of reusing the plain GEMM’s answer.
--epilogue-registers N shrinks the suggestion for fused kernels (e.g. default Size<128, 256, 64> / BlockDim<256> becomes Size<64, 128, 128> / BlockDim<128> at N = 128). This is the only register-pressure control it exposes; cuBLASDx itself has none — the tile is the knob.
Model estimates for CUTLASS kernels, not cuBLASDx measurements: trust the ranking, confirm with a run. Extra fused operations can change the best user tile, so retain several model candidates and repeat the micro-parameter scan around each good measured tile, including once before and once after the tile scan.

Measure, Compare, Stop#

Measurement follows Performance Measurement; the guide binaries already embody it. Done when the kernel is within a modest factor of a reference point (roofline bound, cuBLASLt on the same shape, or the nearest measured guide variant) and further perturbations move it only a few percent — then switch to Finetuning Your Program.