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

# API Card: Blackwell Pipeline (tcgen05)

#### 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 Blackwell `tcgen05` (`utcmma`) pipeline path on SM 1000, 1030, 1070
and 1100 — `SM<1000, arch_specific>` built for `sm_100a`, and the `1030` / `sm_103a`,
`1070` / `sm_107a` and `1100` / `sm_110a` equivalents (`1070` is
Rubin SM107, which uses the SM 1030 instruction set: TMA and floating-point UTCMMA, no integer
`tcgen05`).

SM 1200-class parts are not on this path and resolve to `supermma`. The host and device calls match
the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label); this card repeats them so that it reads alone and states every difference.

The shipped example whose descriptor passes the gates below when compiled for a Blackwell architecture
is `example/cublasdx/19_gemm_batched/batched_gemm_pipeline.cu` (128 threads, `EnableInputStreaming()`, 16-byte alignment, `get_tile` once and
`reset_tile` per iteration). Its Blackwell build needs CUDA Toolkit 13.1 or newer and `100a`
in `CUBLASDX_CUDA_ARCHITECTURES`.

**Descriptor.** The `utcmma` instruction class has silent gates. Missing any one of them leaves the
`tcgen05` path with no error; the result is the older `supermma` class.

* `WithPipeline()`, `EnableInputStreaming()` and `SM<1000, arch_specific>()` (or `1030` / `1070` / `1100`);
  `family_specific` is not accepted.
* `BlockDim<128>` **exactly**. This differs from Hopper: 256 or 512 threads silently leave the path.
* 16-byte alignment of `A`, `B` and `C` (`Alignment<16, 16, 16>` = `MaxAlignment()`); a smaller `C`
  alignment deselects the path.
* `Precision` from the supported real combinations: `half` inputs with a `half` or `float`
  accumulator, `bfloat16` with `float`, `tf32` with `float`, `int8` / `uint8`
  with `int32` (not on SM 1030 or 1070, which have no integer `tcgen05` instruction), and
  `fp8` (`e4m3` or `e5m2`, any pairing) with `half` or
  `float`.

  `double` and complex `Type` are off the path.
* `Size`: `M` a multiple of 64 up to 256 (`fp8` with CUTLASS older than
  4.5.1: a multiple of 128, or exactly 64);

  `N` a multiple of 8 between 8 and 256 for 16-bit and `tf32` inputs (for
  `int8` / `uint8` either 8 with a K-major `B` or a multiple of 16; for `fp8` a multiple of 8
  with a K-major `B` and of 16 otherwise);

  `K` a multiple of 16 for 16-bit inputs, of 8 for `tf32`, of 32 for 8-bit inputs,
  and at most 256 (a larger `M` or `K` leaves the path for `supermma`
  instead of failing; TMA box modes cap at 256).

  The accumulator must also fit tensor memory: `M / 128` blocks for `M` a multiple of
  128, otherwise `M / 64` blocks, times `N` columns (one 32-bit column per accumulator element, `half`
  included) must not exceed 512 columns.

  With `float` accumulation `256 x 256` and `192 x 128` fit and `192 x 256` does
  not.
* `Arrangement` only for `tf32`, which needs K-major `A` and `B` (`row_major` `A`, `col_major`
  `B`, that is TN). `half`, `bfloat16`, `int8` / `uint8` and `fp8` reach `tcgen05` in either
  majorness; this differs from Hopper, where the 8-bit types need TN.

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<128>()                                            // exactly 128
    + cublasdx::Alignment<16, 16, 16>()                                                        // = MaxAlignment()
    + cublasdx::WithPipeline() + cublasdx::EnableInputStreaming()
    + cublasdx::SM<1000, cublasdx::arch_specific>());                                          // or 1030, 1070, 1100
```

Confirm the path in the SASS, per kernel, as the [porting workflow](performance_guide_porting.md#perf-guide-port-label) describes (`/UTC[A-Z]*MMA/`).

As on Hopper, 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` (a same-size `cutlass::bfloat16_t*` passes every other
check), both silently run `supermma` and show no match.

A build that lands on plain `sm_100` instead of `sm_100a` is rejected by `ptxas`:
the `tcgen05` instructions exist only on the `a` target.

**Host side.** Identical to the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label):

```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())              // stages + barriers + TMEM handle
                  .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());
```

* The shape checks, the error codes (including `alignment_unfulfilled`), the free `Depth` and the compile-time
  checks are those of the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label), and so are the lifetime, stream and
  launch rules ([Pipeline Lifetime, Streams And Launch](performance_guide_api_hopper.md#perf-guide-pipeline-lifetime-label)): one handle serves the whole grid, it is passed by value,
  and a pipeline with `RequiredMantissaBits` allocates and slices on the stream passed at creation.
* `get_block_dim()` returns **192** threads for `BlockDim<128>`: the 128 epilogue threads plus one warp
  that issues `tcgen05.mma` and one that issues TMA. Use it, not the descriptor’s thread count, and
  declare the kernel with `__launch_bounds__(DevicePipeline::max_threads_per_block, 1)`.

  Passing `cublasdx::fixed_blocksize` as the template argument after the result-storage mode
  (`suggest_pipeline<Depth, BLAS, cublasdx::internal_accumulator, cublasdx::fixed_blocksize>`) gives a 128-thread block
  instead, in which warps 0 and 1 issue the MMA and TMA work in addition to their share of the epilogue;
  it stays on the `tcgen05` path and is the slower configuration.
* The pipeline allocates all 512 columns of tensor memory when the tile pipeline is constructed and frees them
  when it is destroyed, which is one more reason the kernel runs one CTA per SM. Do not allocate tensor memory
  in the same kernel.
* 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, 48 KiB
for `Size<128, 256, 64>` in bf16, as on Hopper;

`buffer_size()` adds 8 bytes per barrier, two per stage plus two per tensor-memory stage, and a 4-byte
tensor-memory handle: 196,708 B at depth 4 for that tile.

The opt-in limit on a B200 is the H100’s 232,448 B for everything in the block; compute `buffer_size()`
for a candidate depth on the host and compare it with `cudaDevAttrMaxSharedMemoryPerBlockOptin` instead of launching.

On SM 1070 (Rubin SM107) the depth `suggest_pipeline` picks without an explicit depth is budgeted to 227
KiB;

a tile whose single stage plus fixed storage does not fit that but fits below 327 KiB gets depth 1, and
such a launch needs `cudaFuncSetAttribute(kernel, cudaFuncAttributeSharedMemoryMode, cudaSharedMemoryModeAllowOversizedSharedMemory)` instead of the `cudaFuncAttributeMaxDynamicSharedMemorySize` call, which is ignored once a
shared-memory mode is set.

Check the requested size against `cudaDevAttrOversizedSharedMemoryPerBlock` (CUDA runtime 13.4 or newer; the examples handle both
paths in `example::set_max_dynamic_shared_memory_size` from `example/cublasdx/common/common.hpp`). Explicit depths may use more than 227 KiB.

Tensor memory is staged separately: the accumulator gets `512 / (columns per stage)` stages (at most 16), two for `128 x 256`
in `float`, so the epilogue of one tile overlaps the MMAs of the next without any shared memory
of its own.

**Device side.** Use the loop from the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label):

```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 128 epilogue 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
}
```

* `get_tile` once per CTA, `reset_tile` for every further tile, a host-computed bound and a uniform exit
  across the whole block, MMA and TMA warps included: as on Hopper.
* The epilogue functor runs on the 128 threads of `BlockDim`, and every one of them owns part of the result:
  `acc.is_thread_active()` is true for all 128 and exists for code shared with the other paths. The MMA and
  TMA warps never enter the functor; `tile.epilogue_sync()` is the barrier among the 128 and
  `__syncthreads()` inside the functor deadlocks against the other two warps.
* Epilogue accessors: `reduce_and_store(out, lambda)`, `process_and_store(out, op)`, `partition_and_store(out)`, `acc.axpby(alpha, beta, out)` and
  `acc.get_results()` have the meaning given in the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label), with two differences.

  Each of them acquires the tile’s tensor-memory stage, reads it and releases it, so **each of the 128
  threads makes exactly one such call per tile**, unconditionally: a second call waits for the next
  tile’s accumulation, and a thread that skips the call leaves the stage barrier one arrival short, so
  the kernel hangs.

  Read the fragment once with `get_results()` and keep the register copy when the values are needed more
  than once; `acc.map_fragment_index(i)` and `acc.is_index_in_bounds(i)` do not touch tensor memory and may be called freely.

  The output of `partition_and_store` must be a memory-backed tensor (shared or global); `process_and_store`,
  `reduce_and_store` and `axpby` take memory or a register fragment of exactly the accumulator’s
  fragment size, and any other destination is a compile error (`partition_and_store` into a register fragment
  points you to `process_and_store` with `identity{}`).
* With `cublasdx::reusable_accumulator` as the template argument after `BLAS` of `suggest_pipeline`, call `tile.get_accumulator()` once
  before the loop, `tile.execute(acc)` per tile, then `acc.finish_accumulation()` from **every** thread of the block,
  the MMA and TMA warps included and outside any `is_thread_active()` guard: the commit it performs is issued
  by the MMA warp, which that guard excludes.

  Then make the one consuming call above. Do not construct a second accumulator inside the loop.
* The pipeline owns named barriers 1 and 3 (the Hopper pipeline owns 1 and 2); named-barrier ids of your own
  must avoid them.
* Issue no `tcgen05.ld` / `tcgen05.st` and no tensor-memory fences of your own; the accessors above carry
  them. Asynchronous stores you issue in the epilogue (`cp.async.bulk`) must be waited on before the CTA
  reuses their source or exits.
* A reference that uses `cta_group::2` (2-CTA UMMA) or TMA multicast across a cluster is matched with
  `BlocksPerCluster<M, N>()`, described in the [Hopper card](performance_guide_api_hopper.md#perf-guide-api-card-label). On SM 1000, 1030, 1070 and 1100 an even cluster
  `M` requests the 2-CTA atom, which covers a pair of per-CTA `Size` tiles;

  `B` is then multicast over cluster `M / 2`, so `BlocksPerCluster<2, 1>` is one pair with no
  multicast and `<4, 1>` multicasts `B` across two pairs, while `A` is
  still multicast along `N`.

  If the shape cannot take it (`N` not a multiple of 16; of 32 for `int8` /
  `uint8`, and for `fp8` with an MN-major `B`), the pipeline silently uses
  the 1-CTA atom with no diagnostic; measure before assuming the pair is active.

  The pair needs the same `EnableInputStreaming()` descriptor as any `tcgen05` path, and the grid must be
  whole clusters of per-CTA tiles (`cluster_grid_divisibility`). A CUTLASS `256 x N` 2SM tile corresponds to
  `Size<128, N, K>` with an even cluster `M`, not to a `Size<256, ...>` tile.
* `UTCHMMA` (16-bit inputs) and `UTCIMMA` (8-bit integers) are the SASS spellings seen for the check above, so
  match the `/UTC[A-Z]*MMA/` pattern rather than a name.