<a id="perf-guide-api-card-label"></a>

# API Card: Hopper Pipeline (wgmma)

#### 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.

This card covers the Hopper `wgmma` pipeline path; the [Blackwell card](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label) covers the
`tcgen05` pipeline path and the [supermma card](performance_guide_api_supermma.md#perf-guide-api-card-supermma-label) covers warp-level `mma.sync` with either API.

Two shipped examples show the scaffolding around it: `example/cublasdx/01_gemm_introduction/introduction_pipeline.cu` (host setup, dynamic shared memory,
launch, an `axpby` epilogue functor — its descriptor omits `EnableInputStreaming()`, so as shipped it
does not take the `wgmma` path) and `example/cublasdx/19_gemm_batched/batched_gemm_pipeline.cu` (`EnableInputStreaming()` included, `get_tile`
once, `reset_tile` per iteration).

The epilogue accessors below supplement those examples. The runnable
`example/cublasdx/performance_guide/` variants referred to in later chapters ship with the installed examples,
and the optional `nvMatmulHeuristics` configurator needs network access; neither is required for a port.

**Descriptor.** The `wgmma` (`gmma`) instruction class has silent gates. Missing any one of them leaves
the `wgmma` path with no error; the usual result is the older `supermma` class.

* `WithPipeline()`, `EnableInputStreaming()` and `SM<900, arch_specific>()`.
* 16-byte alignment of `A`, `B` and `C` (`Alignment<16, 16, 16>` = `MaxAlignment()`); a smaller `C`
  alignment deselects the path.
* `BlockDim` of 128, 256, 512 or 1024 threads (a power of two made of whole warpgroups).
* `Precision` from the supported real combinations: `half` inputs with a `half` or
  `float` accumulator, `bfloat16` with `float`, `tf32` with
  `float`, `int8` / `uint8` with `int32`, and `fp8`
  (`e4m3` or `e5m2`, any pairing) with `half` or `float`. Any
  other real combination and complex `Type` are off the path.

  (A `bfloat16` accumulator is not a silent gate: the descriptor rejects it at compile time, since
  accumulators must be `__half`, `float`, `double` or a 32- or 64-bit integer.)
* `Size` divisible by the warpgroup tiling: `M` by 64 with 128 threads, by 128 with 256 or 512, by 256 with
  1024; `N` by 8 with 128 or 256 threads, by 16 with 512 or 1024; `K` by 16 for 16-bit inputs, by 8 for
  `tf32`, by 32 for 8-bit inputs.
* `Arrangement` only for `tf32`, `int8` / `uint8` and `fp8`, which need K-major `A` and `B`
  (`row_major` `A`, `col_major` `B`, that is TN); `half` and `bfloat16` reach `wgmma` in either
  majorness.

The descriptor below satisfies all of them:

```cpp
using BLAS = decltype(cublasdx::Size<128, 256, 64>()
    + cublasdx::Precision<__nv_bfloat16, __nv_bfloat16, float>()
    + cublasdx::Type<cublasdx::type::real>() + cublasdx::Function<cublasdx::function::MM>()
    + cublasdx::Arrangement<cublasdx::row_major, cublasdx::col_major, cublasdx::row_major>()  // TN; free for bf16
    + cublasdx::Block() + cublasdx::BlockDim<256>()                                            // 128 or 256
    + cublasdx::Alignment<16, 16, 16>()                                                        // = MaxAlignment()
    + cublasdx::WithPipeline() + cublasdx::EnableInputStreaming()
    + cublasdx::SM<900, cublasdx::arch_specific>());
```

Confirm the path in the SASS, per kernel, as the [porting workflow](performance_guide_porting.md#perf-guide-port-label) describes (`/[HIQ]GMMA/`). A
descriptor that misses a gate above, and a launch whose global tensors’ element types are not exactly
`BLAS::a_value_type` and `BLAS::b_value_type`, both leave `wgmma` silently and show no `GMMA`
match.

The second is why the SASS check matters: the pipeline takes the suggested configuration only for those
exact types, and a `cutlass::bfloat16_t*` or `cute::bfloat16_t*` (`__nv_bfloat16` here) has the same size and
alignment, passes every check, and silently runs `supermma` while still using the pipeline API.

Start with 128 or 256 compute threads on Hopper; measure larger blocks only when the workload benefits
from them.

**Host side.**

```cpp
auto A = cublasdx::make_gmem_tensor<cublasdx::row_major>(a_ptr, M, K, lda);   // a_ptr: BLAS::a_value_type*
auto B = cublasdx::make_gmem_tensor<cublasdx::col_major>(b_ptr, K, N, ldb);
auto p = cublasdx::suggest_pipeline<Depth, BLAS>(A, B);                       // Depth: 1..16, explicit
if (!p) { /* cublasdx::pipeline_error_string(p.error().code) */ }
size_t smem = cublasdx::make_shared_storage_calculator()
                  .add(p->buffer_alignment(), p->buffer_size())              // pipeline stages + barriers
                  .add(scratch_alignment, scratch_bytes)                     // your epilogue / schedule scratch
                  .get();
cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, int(smem));
kernel<<<grid, p->get_block_dim(), smem, stream>>>(..., p->get_device_handle());
```

* `suggest_pipeline` encodes the TMA descriptors at construction, so give it the device pointers the kernel
  will read; the addresses are baked into the descriptors.

  It checks the shape and returns an error instead of a pipeline when `M`, `N`
  or `K` is not a whole multiple of the `Size` tile (`global dimensions are not divisible by the BLAS tile`), when the
  `K` extents of `A` and `B` differ (`incompatible K dimensions`), or, with
  `BlocksPerCluster`, when the tile grid is not whole clusters (`cluster_grid_divisibility`).

  There is no partial-tile path. `Depth` may exceed `K / tile_k`: the pipeline prefills only
  the K tiles that exist, so a shallower depth is only a shared-memory saving, not a requirement.
* `Alignment` is a promise and pipeline creation verifies it: the base pointers of `A`
  and `B` and the byte stride of a dynamic leading dimension must be multiples of the
  alignment the selected copies need, or `suggest_pipeline` / `make_pipeline` returns `pipeline_error_code::alignment_unfulfilled`.

  The alignment checked is the effective copy alignment (16 bytes when the operand takes TMA).

  A *static* stride is known at compile time: the effective alignment is the gcd of the declared
  alignment and the stride’s byte size, so an odd static leading dimension silently narrows the copy
  vector width and can make the operand ineligible for TMA (an error on the [Blackwell path](performance_guide_api_blackwell.md#perf-guide-api-card-blackwell-label)).

  Check `p.error().code` rather than assuming the launch is valid.
* Complex operands take TMA in the pipeline: the descriptor is built over a same-size real proxy type (the bytes
  moved are identical), so a complex `Type` does not prevent a TMA copy. `wgmma` and
  `tcgen05` are real-only; complex resolves to `supermma`.
* Compile-time checks on the pipeline call: the global tensors’ arrangement must equal the descriptor’s
  `Arrangement` (`static_assert`: `GMEM arrangement for A must match SMEM arrangement provided by Arrangement<>`),
  and `suggest_pipeline` without a depth stops with `The tile is too large: not even one pipeline stage fits in
  shared memory` when no stage fits; shrink the tile or lower the precision rather than overriding.
* There is no call to re-point a built pipeline at new pointers. Without `RequiredMantissaBits`, constructing one is cheap
  host-only work; build it per call and do not cache it for speed. The lifetime and stream rules are
  collected in [Pipeline Lifetime, Streams And Launch](#perf-guide-pipeline-lifetime-label) below.
* `get_block_dim()` includes the producer warpgroup: with `BlockDim<256>` the launch has 384 threads. Use
  it, not the descriptor’s thread count, and declare the kernel with
  `__launch_bounds__(DevicePipeline::max_threads_per_block, 1)`.
* Pass the handle as `__grid_constant__ DevicePipeline const device_pipeline`.

**Depth and shared memory.** One stage holds one A tile and one B tile: `(M*K + N*K) * sizeof(input)` bytes, for
example 48 KiB for `Size<128, 256, 64>` in bf16, giving a `buffer_size()` of 196,672 B at depth 4.

The H100 opt-in limit is 232,448 B for everything in the block, including your scratch; at most four
such stages fit, and `cudaFuncSetAttribute` failing with `invalid argument` means the total does not.

Compute `buffer_size()` for a candidate depth on the host and compare it with the limit instead of
launching.

Past the depth that hides the load latency — typically three or four stages for large 16-bit tiles —
extra stages change little, so sweep depth only if a gap above the parity band remains after the
geometry check.

**Device side.**

```cpp
extern __shared__ __align__(device_pipeline.buffer_alignment()) cublasdx::byte smem[];
auto tile = device_pipeline.get_tile(smem, row, col);          // ONCE per CTA
for (;;) {
    tile.execute([&](auto& acc) {                              // runs on the compute threads only
        auto out = cublasdx::get_tile(C, BLAS::c_shape, row, col);
        acc.partition_and_store(out);
    });
    /* advance (row, col) from a host-computed tile count; break uniformly for all threads */
    device_pipeline.reset_tile(tile, row, col);                // re-point; no barrier re-init
}
```

* Call `get_tile` once per CTA and `reset_tile` for every further tile. Constructing a new tile
  pipeline inside the loop re-initializes the shared barriers while the producer may still be issuing
  loads; the kernel hangs on the second tile.

  Bound the loop by a host-computed count, never by scanning shared memory for a sentinel, and keep the
  loop and its exit uniform across the whole block, producer warpgroup included.
* The tile pipeline and a pipeline-bound accumulator own cleanup state: they are not copyable or
  movable, so pass them by reference (never by value into a lambda or helper), declare them once outside
  the persistent loop and declare the accumulator after its tile pipeline.

  Every `execute` must end in an epilogue or `discard()`; destruction neither finishes
  accumulation nor discards an unfinished result (`docs/using_pipelines.rst`, “Pipeline and accumulator lifetime”).
* The epilogue functor runs on the compute warps only. Inside it use `tile.epilogue_sync()` for a barrier
  among those threads — `__syncthreads()` there deadlocks against the producer warpgroup.
  `acc.is_thread_active()` gates threads that own no fragment.
* Epilogue accessors: `reduce_and_store(out, lambda)` calls `lambda(result, old_output, coord)` with the **tile-local** `(m, n)`
  coordinate, so add the tile origin for a global index; `process_and_store(out, op)` for pure element-wise;

  `acc.get_results()` returns a register **copy** of the fp32 fragment (writes to it do not reach the
  accumulator), and `acc.map_fragment_index(i)` gives fragment
  element `i`’s tile-local `(m, n)` — derive any per-row grouping from it rather
  than by probing with `printf`.
* **Read the accumulator once.** `get_results()`, `process_and_store`, `reduce_and_store`, `partition_and_store` and
  `axpby` are each one reading operation, allowed once per ready accumulator.

  A second read hangs or returns the wrong stage on a tensor-memory accumulator (Blackwell) while a
  register accumulator tolerates it, so code that breaks the rule can pass on Hopper and hang on
  Blackwell. To apply several transforms, read `get_results()` once and work on the fragment.

  Calls that do not read (`map_fragment_index`, `is_index_in_bounds`, `is_coord_in_bounds`, `partition_and_copy`) are free
  (`docs/api/other_tensors.rst`, “An epilogue operation that reads the accumulator may be performed only once”).
* `process_and_store`, `reduce_and_store` and `axpby` also accept a **register fragment** as destination, on
  every backend, if its size equals the accumulator’s fragment size exactly; `reduce_and_store` then also visits
  padded slots, so guard coordinate-indexed lambdas with `acc.is_coord_in_bounds(coord)`.
  `partition_and_store` takes memory only. With a memory destination, `axpby` with `beta == 0` never reads
  `C`, so `C` may be uninitialized.
* Asynchronous stores you issue yourself in the epilogue (`cp.async.bulk`) must be waited on before the CTA
  reuses their source or exits; a CTA barrier does not wait for them.
* The pipeline supports CTA clusters and TMA multicast. To match a reference that multicasts an operand
  across a cluster, add `cublasdx::BlocksPerCluster<M, N>()` to the descriptor next to `WithPipeline()`: extents are in CTAs,
  `K` must be 1, `M * N` at most 16.

  `A` is multicast along cluster `N` and `B` along cluster
  `M`, so `BlocksPerCluster<2, 1>` multicasts `B`. The operator is a hint: non-TMA copies
  ignore it, and on Hopper multicast also needs a `BlockDim` that is a multiple of 128.
* Check what was selected with the static members `uses_cluster` and `uses_multicast` of the pipeline
  type (`std::decay_t<decltype(*p)>::uses_cluster`). A cluster can be active while multicast is silently demoted (a
  non-power-of-two extent, or a split that leaves less than whole swizzle rows), so test `uses_multicast`
  for multicast.

  When `uses_cluster` is true, launch with `p->get_cluster_dim()` through `cudaLaunchKernelEx` and `cudaLaunchAttributeClusterDimension`
  instead of `<<<>>>`, use a grid of whole clusters (CTAs are laid out `M`-fastest,
  so `rank_m = blockIdx.x % M`), and set `cudaFuncAttributeNonPortableClusterSizeAllowed` above 8 CTAs.

  The per-CTA tile grid must contain whole clusters (`M / tile_m` divisible by cluster
  `M` and `N / tile_n` by cluster `N`) or pipeline creation fails with
  `cluster_grid_divisibility`. Descriptor traits: `has_blocks_per_cluster_v`, `uses_blocks_per_cluster_v` (false without `WithPipeline` or
  for `BlocksPerCluster<1>`) and `blocks_per_cluster_of_v_m` / `_n` / `_k`.
* `example/cublasdx/22_persistent_cluster_gemm/persistent_cluster_gemm.cu` and `example::launch` in
  `example/cublasdx/common/cluster.hpp` show all of it: a persistent grid of whole clusters sized by
  `example::suggest_persistent_cluster_grid` (`cudaOccupancyMaxActiveClusters`, since clusters cannot straddle
  GPCs and `num_sms / cluster_size` can overestimate residency) and a windowed rasterizer for L2 locality. That
  schedule is the example’s own; a port [keeps the reference’s schedule](performance_guide_porting.md#perf-guide-port-integrate-label)
  and takes from the example only the launch mechanics.

<a id="perf-guide-pipeline-lifetime-label"></a>

## Pipeline Lifetime, Streams And Launch

* `suggest_pipeline` / `make_pipeline` for a descriptor without `RequiredMantissaBits` do no GPU work: no
  allocation, no kernel, no stream use (the optional stream argument is ignored). They only validate the shapes and
  alignments and encode descriptors on the host.
* With `RequiredMantissaBits` the opposite holds. Creation allocates temporary device memory (`cudaMallocAsync` or a
  memory pool) and launches the max-reduction and slicing kernels that read the global `A` and `B`, all on the stream
  passed as the last argument, and the host pipeline object frees the memory (`cudaFreeAsync`) on that same stream
  when it is destroyed. Therefore: pass the stream you launch the GEMM kernel on, so that the slicing is ordered
  before the kernel; keep `A` and `B` unchanged until that stream has consumed them; keep the host pipeline
  object alive until the kernel is enqueued (taking the device handle from a temporary host pipeline is deleted at
  compile time for emulation); and do not use the handle on another stream without your own synchronization.
* `get_device_handle()` returns a value type. Pass it to the kernel by value as `__grid_constant__
  DevicePipeline const`; copies are cheap and interchangeable. The handle holds the global-tensor descriptors (TMA
  descriptors or the tensors’ base pointers and strides, so it is bound to those addresses and extents) and the shared
  layouts, not any device allocation of its own.
* One pipeline serves the whole grid: every block receives the same handle and runs its own tile loop with its own
  shared-memory stages; do not create one per block or per tile. Create a new one when pointers, shapes or leading
  dimensions change.
* Launch with `get_block_dim()` (it includes the producer threads) and, for a cluster descriptor, with
  `get_cluster_dim()` through `cudaLaunchKernelEx` and `cudaLaunchAttributeClusterDimension`.
* Reusing one handle across launches is fine as long as the tensors it was built from stay valid and unchanged in
  shape; a second launch of the same kernel with the same handle needs no new pipeline.