# Finding Problem Shape #### 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. Describe the problem in the terms the library reasons about: what is multiplied, how much, where it lives, how many independent instances at once. Output of this chapter: an operator answer per row below, plus the fixed/free record from **Axes That Can Be Relaxed**. ## Describing A GEMM Only three operators have no default; the rest resolve silently to something. | Operator | Required | What you have to decide | |-----------------------------------|---------------------|--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| | `Size` | **Yes** | The per-tile problem, not the whole matrix. All three must be greater than 0. | | `Function` | **Yes** | That this is a GEMM. Note the spelling — there is no `function::gemm`. | | `SM` | **Yes** | The target architecture, e.g. `SM<800>` for `sm_80`. An optional second parameter takes
`arch_specific` or `family_specific` (the `a` / `f` suffixes), which are only legal for
`Arch >= 900` and are what unlock warp specialization and the newest MMA instructions.

CMake and descriptor spellings differ on three-digit architectures: CMake/nvcc `100a` maps to
`SM<1000, arch_specific>()`, not `SM<100>`. Keep the build architecture and descriptor synchronized;
the performance-guide CMake does this through `GUIDE_ARCH`. | | `Block()` | **Yes**, to execute | GEMM is block-level only — `Thread()` is TRSM-only. Without `Block()` you have a description you can
query, but not run. | | `Precision` | Defaults | Input and accumulation precision. One-, two- and three-argument forms are all legal.

**Defaults to** `Precision` — so a descriptor that never mentions precision is an
fp32 one, whatever your data is. | | `Type` | Defaults | Real or complex. **Defaults to** `type::real`. | | `Arrangement` | Defaults | Row- or column-majorness of each matrix, written with `cublasdx::row_major` / `col_major`. This is the
basis for layout construction and instruction selection, so for good performance it should reflect how
the data actually sits in global memory.

**Defaults to** ``. | | `BlockDim` | Defaults | How many of the block’s threads do the GEMM. Flat size must be within `[32, 1024]`, and the launch rule is unchecked.

**Defaults to** `BLAS::suggested_block_dim`, the tuning database’s choice. | | `BlocksPerCluster` | Optional | Cluster extents in CTAs; takes effect with `WithPipeline()`. `Size` remains the per-CTA tile.

`N` and `K` default to 1; `K` must be 1 and `M * N` must not exceed 16. Without the operator,
the extents are `1 x 1 x 1` and clustered execution is disabled.

On supported paths, clustering can enable TMA multicast and 2-CTA MMA. The tile grid and launch must
match the selected cluster; see the [Hopper](performance_guide_api_hopper.md#perf-guide-api-card-label) and
[Blackwell](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label) cards. | | `LeadingDimension` | Optional | Only if your data is strided. `LDA` and `LDB` are both required when the operator is used; only
`LDC` is optional.

#### WARNING
Setting this **disables the suggested-layout database and the suggested accumulator** — both fall
back silently. Prefer passing the leading dimension to `get_layout_smem_*()` at the call site. | Tuning-only operators include `Alignment`, `StaticBlockDim`, `EnableInputStreaming`, `WithPipeline`, and `RequiredMantissaBits`. TRSM descriptors have additional operators and defaults that select the solve. ## Composing The Descriptor One operator per answer, summed with `+`, captured with `decltype`. Order does not matter; what matters is that every line is a decision you made rather than a default you inherited. The result is a type — everything downstream hangs off it. ```cpp #include #include using namespace cublasdx; // operators live in namespace cublasdx using GEMM = decltype(Size<128, 128, 32>() // per-tile M, N, K + Precision<__half, __half, float>() // A, B in half; accumulate in float + Type() // real, not complex + Function() // this is a GEMM + Arrangement() + SM<900, arch_specific>() // sm_90a + Block() // block-level operation + BlockDim<256>()); // threads spent on the GEMM ``` Partial descriptors are legal and reusable — this is how one axis is swept: ```cpp using problem = decltype(Size<128, 128, 32>() + Precision<__half, __half, float>() + Type() + Function() + Arrangement() + Block()); using on_ampere = decltype(problem() + SM<800>() + BlockDim<256>()); using on_hopper = decltype(problem() + SM<900, arch_specific>() + BlockDim<256>()); ``` ## Axes That Can Be Relaxed Algorithmic-level relaxations are usually worth more than anything in the later chapters, and only you can make them — they trade accuracy, memory or interface constraints for speed. Record each axis as **fixed** by the application or **free** to experiment with; that record is the search space every later chapter operates in. | Axis | Question to settle | |-----------------------|----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| | Computation precision | Can it be lowered? Inputs may be `float` while the multiply runs in `TF32` or `FP16`.

The single highest-leverage question here. Experiment, and **report the precision loss** — a
relaxation that silently degrades accuracy is not a win. | | Thread count | Is it fixed by the surrounding kernel, or can it be swept? | | Tile shape | Is the tile fixed by the surrounding problem, or can it be swept? | | Emulation | Can an emulated-precision scheme be applied with `RequiredMantissaBits`? See [Emulated Precision (RequiredMantissaBits)](#perf-guide-emulation-label). | | Shared memory budget | How much shared memory is free for the GEMM? Lowering precision lets more pipeline stages or a bigger
tile fit in the same budget. | | Register budget | Is the surrounding kernel register-heavy? Accumulator registers grow with `M * N`, so a smaller tile
may be necessary. | ## Emulated Precision (RequiredMantissaBits) `RequiredMantissaBits` `` turns a `float` or `double` pipelined GEMM into an integer-slice emulation: A and B are split into `ceil((X + 1) / 8)` `int8` slices under a shared exponent per row of A and per column of B, the slice products run on the integer tensor cores into an `int32` accumulator, and the result is recombined into the descriptor’s floating-point type. It is an accuracy/throughput trade: `X` is the number of mantissa bits you ask the result to preserve, so measure the error against a native reference before adopting it. The descriptor constraints enforced by this library (each is a `static_assert` in `blas_description.hpp`): * exactly one `RequiredMantissaBits` operator; * `WithPipeline()` and `Block()` are present, and `Function()`; * `Type()`; * `Precision()` or `Precision()` (all three equal, nothing mixed). At pipeline creation (`suggest_emulation_pipeline`, reached through `suggest_pipeline` / `make_pipeline`): * `A` and `B` must be rank-2 tensors whose element type equals the descriptor’s precision (both accumulator modes work; the slice loop drives its internal integer pipeline in reusable mode either way); * global `K` is bounded: `pipeline_error_code::emulation_k_dimension_too_large` is returned when `K` exceeds 33025, the largest `K` for which the `int32` accumulation cannot overflow; split `K` across several pipelines and add the partial results for larger problems; * the divisibility rules of [Shapes That Do Not Divide The Tile](#perf-guide-partial-tiles-label) apply to the sliced tensors as well. The operator is a no-op request on some architectures: where the native `float` / `double` pipeline already carries `X` bits, `suggest_pipeline` keeps it (see the operator reference). Runnable starting points are `example/cublasdx/16_dgemm_emulation` (the block-level building blocks), `20_dgemm_emulation_pipeline` and `21_sgemm_emulation_pipeline`. Lifetime and stream rules of the emulation pipeline (it allocates and launches slicing kernels when it is created) are in [Pipeline Lifetime, Streams And Launch](performance_guide_api_hopper.md#perf-guide-pipeline-lifetime-label). ## Shapes That Do Not Divide The Tile A pipeline has no partial-tile path. `suggest_pipeline` / `make_pipeline` return `pipeline_error_code::tile_divisibility` (`global dimensions are not divisible by the BLAS tile`) unless `M`, `N` and `K` are each a whole multiple of the `Size` tile, and `pipeline_error_code::cluster_grid_divisibility` unless the tile grid is a whole number of clusters. The block-level API has no global shape at all: `execute` works on shared-memory (or register) tensors whose shapes must be static and equal to `Size` (`All layout shapes must be static, only strides can be dynamic`), so a ragged edge is the caller’s tile loop to handle. Three strategies: | Strategy | Use when | Cost | |---------------------------------|-------------------------------------------------------------------------------------------------------------------------------------|-----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| | Pad the allocation | You own the allocation (or can copy once), the extra rows and columns of `C` may be written, and padded
`K` can be zero-filled. | Simplest: no extra code and no extra compile time. Wasted work is the padded fraction of the problem. | | Pad the inputs, mask the store | Only the write of `C` must stay inside the logical extent, but `A` and `B` can be padded (zeros in the
`K` padding). | One extra shared-memory hop in the epilogue: write the accumulator to a shared `C` tile
(`partition_and_store`), synchronize, then copy out with a per-element bounds check. There is no
dynamic-shape `C` on the accumulator API of this library, so the check is yours. | | Bulk plus remainder descriptors | Inputs cannot be padded, or the waste is unacceptable. | General, but one more descriptor and kernel instantiation per distinct remainder, so more compile time.
The remainder `Size` is a compile-time value, and it must itself be legal for the instruction path
(for example `M` a multiple of 64 for `wgmma`); a remainder that is not, or that is only known at
run time, falls back to padding. | Decision rule: pad if you can; mask the store when the output extent is contractual; split only when padding the inputs is impossible and the remainder is a few fixed, instruction-legal sizes. Sketch of the split for `M` (`N` is symmetric; for `K`, either zero-pad or run the tail pipeline with `beta = 1` into the same `C`): ```cpp // Base: your descriptor without Size (precision, Arrangement, Block, BlockDim, // WithPipeline, SM, ...) constexpr int tile_m = 128, tail_m = 64; // tail_m == M % tile_m, known at compile time using Bulk = decltype(Base() + cublasdx::Size()); using Tail = decltype(Base() + cublasdx::Size()); int m_bulk = M - tail_m; // a whole number of tile_m rows auto b = cublasdx::make_gmem_tensor(b_ptr, K, N, ldb); auto a0 = cublasdx::make_gmem_tensor(a_ptr, m_bulk, K, lda); auto p0 = cublasdx::suggest_pipeline(a0, b); // check p0, as always kernel<<get_block_dim(), smem0, stream>>>(p0->get_device_handle(), c_ptr, ldc); auto a1 = cublasdx::make_gmem_tensor(a_ptr + size_t(m_bulk) * lda, tail_m, K, lda); auto p1 = cublasdx::suggest_pipeline(a1, b); kernel<<get_block_dim(), smem1, stream>>>(p1->get_device_handle(), c_ptr + m_bulk, ldc); // col-major C ``` Each launch is an ordinary divisible problem, so everything else in the guide applies to it unchanged; the sub-tensor base pointers must still satisfy the declared `Alignment`.