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; 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:

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 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:

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, and so are the lifetime, stream and launch rules (Pipeline Lifetime, Streams And Launch): 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:

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