# 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](performance_guide_porting.md#perf-guide-port-label) 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](performance_guide_finetuning.md#perf-guide-finetuning-label). ## Best Idiom Search 1. **One datapoint per axis.** For each idiom axis your kernel touches, measure its variants in isolation and note the winner. 2. **Construct the hypothesis.** The per-axis winners give the combination you expect to be best. 3. **Measure the hypothesis.** Combinations frequently do not behave like the sum of their parts. 4. **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): ```bash cmake -DBUILD_EXAMPLE=ON \ -DCUBLASDX_EXAMPLES_ENABLE_PERFORMANCE_GUIDE=ON \ -DCUBLASDX_CUDA_ARCHITECTURES= \ -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 `/example/cublasdx/performance_guide//` 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](blas_performance_guide.md#perf-guide-label)). ## 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_`. 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. * `.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( )` 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`. | * 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` | 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](performance_guide_problem_shape.md#perf-guide-relaxable-axes-label) | 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](performance_guide_problem_shape.md#perf-guide-shape-label)). ### Tile configurator A first tile before measured searching — run this Python configurator explicitly rather than copying the descriptor from an example: [nvMatmulHeuristics](https://github.com/NVIDIA/CUDALibrarySamples/tree/master/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: ```bash 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](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label) 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](blas_performance_guide.md#perf-guide-measurement-label); 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](performance_guide_finetuning.md#perf-guide-finetuning-label).