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 covers the tcgen05 pipeline path and the supermma card 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:

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

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

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

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 and takes from the example only the launch mechanics.

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.