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

> Master FP8 and FP4 GEMM kernels in DeepGEMM by transforming scaling factors using transform_sf_pair_into_required_layout. Learn the correct layouts for SM 9.0 and SM 10.0.

- Repository: [DeepSeek/DeepGEMM](https://github.com/deepseek-ai/DeepGEMM)
- Tags: how-to-guide
- Published: 2026-04-19

---

**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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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)

```python
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

```python

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

```python

# 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](https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/include/deep_gemm/impls/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](https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_gemm_1d1d.cuh) |
| [`csrc/apis/layout.hpp`](https://github.com/deepseek-ai/DeepGEMM/blob/main/csrc/apis/layout.hpp) | Layout transformation utilities (`transform_sf_into_required_layout`, `transform_sf_pair_into_required_layout`) | [layout.hpp](https://github.com/deepseek-ai/DeepGEMM/blob/main/csrc/apis/layout.hpp) |
| [`csrc/apis/gemm.hpp`](https://github.com/deepseek-ai/DeepGEMM/blob/main/csrc/apis/gemm.hpp) | Public Python API for FP8/FP4 GEMM and grouped variants | [gemm.hpp](https://github.com/deepseek-ai/DeepGEMM/blob/main/csrc/apis/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](https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/include/deep_gemm/ptx/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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.