# cuBLASDx

> Device-extension BLAS (GEMM, TRSM) meant to be fused into custom CUDA kernels rather than called from the
> host. Performance depends heavily on the kernel the library is embedded in, so the entry point is an
> empirical guide with matching runnable code, not a reference page.

cuBLASDx is a CUDA library from the MathDx SDK. Device extension libraries are fused into custom kernels and
are invocable from CUDA C++, NVIDIA Warp and Numba CUDA code. cuBLASDx exposes GEMM, TRSM and data
layouting/copying functionality at threadblock (both) and thread (TRSM only) level. A call has two parts:
compose a descriptor (the mathematical problem plus the available resources), then execute it through one of
several library idioms.

Execution performance depends on the chosen problem shape and on resource availability, both of which belong to
the surrounding kernel — so the reference documentation explains concepts but cannot answer user-dependent
questions. The guide below is built around experimentation instead: establish which resources are available and
which problem axes can move, pick a first approximation, profile against theoretical limits, and converge by
measuring.

One hard rule before any measurement: **never benchmark on zero-filled buffers** — zeros draw less power, so
clocks boost and measured throughput is overstated. Fill operands from a distribution resembling
the real ones. Full measurement protocol:
[Performance Measurement](markdown/blas_performance_guide.md#perf-guide-measurement-label).

## Start here

The performance guide below covers GEMM. TRSM is documented in [Using cuBLASDx TRSM](using_trsm.html).

- [BLAS Performance Guide](markdown/blas_performance_guide.md): choose the workflow for a new kernel or a port; also
  holds library setup, the Build Card, and the measurement protocol.
- [Porting an existing kernel](markdown/performance_guide_porting.md): identify application constraints and tunable
  axes, validate the intended instruction path, measure against the parity band, and stop at parity.
- API cards: [Hopper pipeline (wgmma)](markdown/performance_guide_api_hopper.md),
  [Blackwell pipeline (tcgen05)](markdown/performance_guide_api_blackwell.md), and
  [supermma](markdown/performance_guide_api_supermma.md). Read the card for the selected execution path.
- [Finding problem shape](markdown/performance_guide_problem_shape.md): describing the problem before choosing a
  configuration; which operators have defaults and which do not; [the axes that can be
  relaxed](markdown/performance_guide_problem_shape.md#perf-guide-relaxable-axes-label) — the highest-leverage decisions, made
  before any tuning.
- [Fitting into cuBLASDx idioms](markdown/performance_guide_idioms.md): map your kernel onto one or more of the 19
  idioms. Start from the "If your kernel..." table.
- [Finding the best coarse solution](markdown/performance_guide_coarse_solution.md): the search procedure, the solution
  variants for each idiom, and how to build and run them.
- [Finetuning your program](markdown/performance_guide_finetuning.md): closing the remaining gap once a coarse solution
  is in hand.

## Picking a tile

Tile shape, block dimension and pipeline depth dominate GEMM performance and interact, so a hand search is a
rebuild per point. `example/cublasdx/performance_guide/suggest_tile.py` (wrapping the optional nvMatmulHeuristics
package) prints a `Size<...>`, a `BlockDim<...>` and a stage count for `suggest_pipeline` from a performance
model in about a second; `--epilogue-registers N` shrinks the suggestion for fused kernels. Details and worked
numbers: [the coarse-solution chapter](markdown/performance_guide_coarse_solution.md#perf-guide-tile-configurator-label).

## The executable half

- [Building and running the variants](markdown/performance_guide_coarse_solution.md#perf-guide-build-label):
  one program per solution variant, under `example/cublasdx/performance_guide/` in the distribution. Most idioms
  have several competing implementations, measured against a cuBLASLt + thrust reference.

## Other documentation

- [Requirements](requirements_func.html): supported architectures, toolkits and types.
- [Installation](installation.html): obtaining MathDx and wiring it into a build.
- [Using cuBLASDx GEMM](using_cublasdx.html): the main narrative guide.
- [Using pipelined GEMM](using_pipelines.html): the pipeline API, which unlocks TMA and the Hopper/Blackwell
  MMA instructions.
- [Using cuBLASDx TRSM](using_trsm.html): triangular solve.
- [Achieving high performance](performance.html): the guide's advice as a flat checklist, without the procedure.
- [API reference](api/operators.html): the full reference, starting with operators.
- [Learning path through examples](examples_learning_path.html): the non-guide examples, ordered.
- [Python bindings](python_bindings.html)
- [Release notes](release_notes.html)