How to Use FP8 and FP4 GEMM Kernels with Proper Scaling Factor Layouts in DeepGEMM

To use FP8 and FP4 GEMM kernels in DeepGEMM, you must transform scaling factors into the TMA-aligned, MN-major layout using transform_sf_pair_into_required_layout() before passing them to the kernel, with FP32 factors for SM 9.0 and packed UE8M0 INT32 for SM 10.0.

DeepGEMM provides high-performance General Matrix Multiply (GEMM) kernels that operate on FP8 and FP4 matrix data with per-row or per-column scaling factors. Proper handling of scaling factor layouts is critical because the kernels use Tensor Memory (TMEM) TMA operations that require specific alignment and memory ordering. This guide explains how to prepare scaling factors correctly for both NVIDIA Ampere (SM 9.0) and Blackwell (SM 10.0) architectures.

Architecture-Specific Kernel Implementations

DeepGEMM implements distinct kernel variants for different GPU architectures to optimize the FP8 and FP4 compute pipelines.

SM 9.0 (Ampere) FP8 Kernels

For SM 9.0 devices (e.g., A100), the kernels are defined in deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh and sm90_fp8_gemm_1d2d.cuh. These kernels expect FP32 scaling factors that are applied to the matrix multiply accumulator before the final store. The scaling logic is implemented in lines 55-63 of sm90_fp8_gemm_1d1d.cuh, where the kernel reads scaling factors from shared memory (smem_sfa / smem_sfb) and multiplies them with the accumulator.

SM 10.0 (Blackwell) FP8/FP4 Kernels

For SM 10.0 devices (e.g., H100), the kernel is defined in deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_gemm_1d1d.cuh. This kernel uses INT32 packed UE8M0 scaling factors (unsigned 8-bit mantissa with 0-bit exponent). The kernel fuses the scaling multiplication into the UMMA instruction (SM100_MMA_MXF4_SS in lines 18-38). The scaling factor handling occurs in lines 326-338, where the kernel unpacks the UE8M0 values and applies them to the FP8/FP4 operands.

Understanding Scaling Factor Layout Requirements

The kernels load scaling factors via Tensor Memory Accelerator (TMA) operations, which impose strict memory layout constraints.

TMA Alignment and MN-Major Order

Scaling factor tensors must be stored in TMA-aligned, MN-major layout that matches the GEMM block configuration. This means:

• Alignment: The buffer size must be a multiple of 128 bytes (get_tma_aligned_size in csrc/apis/layout.hpp ensures this).
• MN-major order: The kernel reads one whole M (or N) block per shared memory load, so the scaling factors must be contiguous across the MN dimension.

Data Type Differences

• SM 9.0: Uses raw FP32 scaling factors. The layout transformation (transform_sf_into_required_layout) aligns the tensor to TMA boundaries using get_mn_major_tma_aligned_tensor.
• SM 10.0: Packs FP32 scaling factors into UE8M0 format (8-bit unsigned). The transformation uses get_mn_major_tma_aligned_packed_ue8m0_tensor to pack and align the data into INT32 tensors.

Layout Transformation API

DeepGEMM provides utilities in csrc/apis/layout.hpp to handle the required layout transformations automatically.

Single Tensor Transformation

Use transform_sf_into_required_layout(sf, mn, k, recipe, ...) to transform a single scaling factor tensor. This function:

1. Selects the appropriate recipe (granularity tuple) based on architecture
2. Applies TMA alignment
3. Handles UE8M0 packing for SM 10.0

Pair Transformation for A and B Matrices

For GEMM operations, use transform_sf_pair_into_required_layout(sfa, sfb, m, n, k, ...) to transform both A-side and B-side scaling factors together. This function returns:

• The transformed tensors (sfa_t, sfb_t)
• The granularity of K used by the kernel (gran_k_a, gran_k_b)

The default recipe is automatically chosen for the current device:

• SM 9.0: (1, 1, 128) → FP32, 1-element MN, 128-element K
• SM 10.0: (1, 32, 128) or (1, 128, 128) → INT32, 1-element MN, 32/128-element K

Complete Implementation Examples

Basic FP8 GEMM (SM 9.0 or SM 10.0)

import torch
import deep_gemm

M, N, K = 1024, 1024, 128

# FP8 matrices stored as uint8 (interpreted as FP8 by kernel)

a = torch.randint(0, 255, (M, K), dtype=torch.uint8, device='cuda')
b = torch.randint(0, 255, (N, K), dtype=torch.uint8, device='cuda')

# Per-element scaling factors (FP32)

sfa = torch.randn(M, K, dtype=torch.float32, device='cuda')
sfb = torch.randn(N, K, dtype=torch.float32, device='cuda')

# Transform scaling-factor layouts (auto-selects recipe for device arch)

sfa_t, sfb_t, _, _ = deep_gemm.transform_sf_pair_into_required_layout(
    sfa, sfb, M, N, K)

# Output tensor (BF16 or FP32)

d = torch.empty((M, N), dtype=torch.bfloat16, device='cuda')

# Execute GEMM (kernel auto-selected based on architecture)

deep_gemm.fp8_fp4_gemm_nt(
    a=(a, sfa_t),
    b=(b, sfb_t),
    d=d,
    compiled_dims="nk")

M-Grouped GEMM with Contiguous Layout

# Grouped layout: G groups × N matrix, int32, contiguous

G = 4  # number of groups

grouped_layout = torch.arange(0, G, dtype=torch.int32, device='cuda').repeat_interleave(N)

# Transform scaling factors (same as non-grouped)

sfa_t, sfb_t, _, _ = deep_gemm.transform_sf_pair_into_required_layout(
    sfa, sfb, M, N, K)

# Execute grouped GEMM

deep_gemm.m_grouped_fp8_fp4_gemm_nt_contiguous(
    a=(a, sfa_t),
    b=(b, sfb_t),
    d=d,
    grouped_layout=grouped_layout,
    compiled_dims="nk")

K-Grouped GEMM (Variable K per Group)

# Variable K sizes per group

ks = [64, 128, 96]
ks_tensor = torch.tensor(ks, dtype=torch.int32, device='cuda')
recipe = (1, 1, 128)  # granularity for SM 10

# Transform K-grouped scaling factors

sfa_t = deep_gemm.transform_k_grouped_sf_into_required_layout(sfa, ks, ks_tensor, recipe)
sfb_t = deep_gemm.transform_k_grouped_sf_into_required_layout(sfb, ks, ks_tensor, recipe)

# Execute K-grouped GEMM

deep_gemm.k_grouped_fp8_gemm_nt_contiguous(
    a=(a, sfa_t),
    b=(b, sfb_t),
    d=d,
    ks=ks,
    ks_tensor=ks_tensor,
    recipe=recipe,
    compiled_dims="mn")

Key Source Files

File	Purpose	Link
deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh	SM 9.0 FP8 kernel (1-D-1-D layout) with FP32 scaling factors	sm90_fp8_gemm_1d1d.cuh
deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_gemm_1d1d.cuh	SM 10.0 FP8/FP4 kernel using packed UE8M0 INT32 scaling factors	sm100_fp8_fp4_gemm_1d1d.cuh
csrc/apis/layout.hpp	Layout transformation utilities (transform_sf_into_required_layout, transform_sf_pair_into_required_layout)	layout.hpp
csrc/apis/gemm.hpp	Public Python API for FP8/FP4 GEMM and grouped variants	gemm.hpp
deep_gemm/include/deep_gemm/ptx/tcgen05.cuh	Low-level PTX wrappers for block-scale MMA instructions (SM100_MMA_MXF4_SS)	tcgen05.cuh

Summary

• Architecture-specific formats: SM 9.0 uses FP32 scaling factors, while SM 10.0 uses packed UE8M0 INT32 format for higher density.
• Layout transformation is mandatory: Always call transform_sf_pair_into_required_layout() to convert scaling factors into TMA-aligned, MN-major layout before passing to GEMM kernels.
• Recipe selection: The transformation API automatically selects the correct granularity recipe (e.g., (1, 1, 128) for SM 9.0, (1, 32, 128) for SM 10.0) based on the detected device architecture.
• Grouped GEMM support: The same scaling factor transformation applies to M-grouped and K-grouped variants; only the layout tensor and API entry point change.
• Zero-copy fusion: Proper layout allows the kernel to fuse scaling multiplication into the UMMA instruction without extra memory traffic, maximizing FP8/FP4 throughput.

Frequently Asked Questions

What happens if I don't transform scaling factor layouts before calling the GEMM kernel?

If you pass untransformed scaling factors directly to fp8_fp4_gemm_nt(), the kernel will encounter illegal memory accesses or severe performance penalties because TMA (Tensor Memory Accelerator) operations require 128-byte aligned, MN-major contiguous memory layouts. The kernel reads scaling factors via smem_sfa and smem_sfb shared memory buffers that assume the transformed layout prepared by transform_sf_into_required_layout().

How do I choose between FP8 and FP4 on SM 10.0 GPUs?

The kernel selection happens automatically inside sm100_fp8_fp4_gemm_1d1d.cuh based on the input data type. When you pass uint8 tensors containing FP8 data, the kernel uses the FP8 pipeline; for FP4 operations, you provide the appropriate 4-bit packed format and the kernel invokes the SM100_MMA_MXF4_SS instruction. The scaling factor transformation remains identical for both precisions on SM 10.0, using the packed UE8M0 INT32 format.

Can I use the same scaling factors for grouped and non-grouped GEMMs?

Yes, the scaling factor transformation process is identical. You call transform_sf_pair_into_required_layout() to prepare sfa_t and sfb_t, then pass these transformed tensors to either the standard fp8_fp4_gemm_nt() or the grouped variants like m_grouped_fp8_fp4_gemm_nt_contiguous(). The only additional requirement for grouped GEMMs is providing the grouped_layout tensor (int32) that describes per-group offsets, as implemented in csrc/apis/gemm.hpp.

What is UE8M0 format and why does SM 10.0 use it for scaling factors?

UE8M0 stands for Unsigned 8-bit Mantissa, 0-bit Exponent—a custom 8-bit floating-point format used to compress scaling factors on SM 10.0 (Blackwell). According to csrc/apis/layout.hpp, the function get_mn_major_tma_aligned_packed_ue8m0_tensor packs FP32 scaling factors into this dense format, storing four UE8M0 values per 32-bit integer. This packing reduces memory bandwidth by 4× compared to FP32, which is essential for maintaining the high throughput of FP8/FP4 GEMM operations on Blackwell GPUs where the UMMA instruction fuses the scaled multiply-accumulate directly.

{
  "mcpServers": {
    "instagit": {
      "command": "npx",
      "args": ["-y", "instagit@latest"]
    }
  }
}