DeepGEMM Scaling Factor Transformation Utilities: A Complete Guide to FP8 and FP4 Quantization

DeepGEMM's scaling factor transformation utilities are helper functions that compute, reshape, and pack scaling factors required for low-precision matrix multiplication kernels, converting high-precision FP32 tensors into hardware-optimized FP8 or FP4 representations with proper quantization metadata.

DeepGEMM is an open-source library for high-throughput GEMM operations optimized for MoE (Mixture of Experts) models. The scaling factor transformation utilities form a critical preprocessing layer that prepares quantization metadata before low-precision kernels execute. These utilities handle tensor shape calculations, UE8M0 format packing, and per-token or per-channel scaling factor generation across the codebase.

What Are Scaling Factor Transformation Utilities?

These utilities are specialized helper functions that bridge the precision gap between FP32 tensors and low-precision hardware kernels. They live in two primary locations within the DeepGEMM repository.

Core Utility Functions

Function	Purpose	Source Location
get_sf_shape	Computes the tensor shape (number of scales × tokens) required to store per-token or per-channel scaling factors, handling UE8M0 and column-major layout rules.	third-party/tilelang_ops/utils.py
get_sf_and_inv	Derives the floating-point scale and its reciprocal from a maximum absolute value (amax). Optionally rounds the scale to a power-of-two and packs it as an 8-bit integer for UE8M0.	third-party/tilelang_ops/utils.py
per_*_cast_to_fp8	Quantizes a 2-D tensor to FP8 (torch.float8_e4m3fn) while producing per-token, per-channel, per-block, or custom-dimension scaling factors.	deep_gemm/utils/math.py
per_*_cast_to_fp4	Similar to FP8 helpers but for 4-bit representation, handling optional UE8M0 packing.	deep_gemm/utils/math.py
pack_ue8m0_to_int / unpack_ue8m0_from_int	Converts UE8M0-packed 8-bit scale values to/from a 32-bit integer container.	deep_gemm/utils/math.py
ceil_to_ue8m0	Rounds a floating-point scale to the nearest representable UE8M0 value before packing.	deep_gemm/utils/math.py

When Are Scaling Factor Transformation Utilities Needed?

These utilities are required whenever converting high-precision tensors to low-precision formats for hardware-accelerated GEMM operations.

Scenario	Why the Utility is Required
FP8 quantization of activations/weights	Hardware kernels expect FP8 inputs with an 8-bit (or 4-bit) scale that rescales the integer representation back to FP32 during computation. Functions like per_token_cast_to_fp8 compute the per-token sf and optionally pack it for UE8M0.
FP4 (4-bit) quantization	FP4 packs two 4-bit codes into a single byte. per_token_cast_to_fp4 creates the codebook, computes the per-token scale based on the max absolute value of each token, and optionally packs the UE8M0 scale.
Column-major scaling layout	Certain Tile-Lang kernels (e.g., swiglu_apply_weight_to_fp8) require scaling factors stored in column-major order and aligned to 16-byte boundaries. get_sf_shape calculates the correctly aligned shape, and use_col_major_scales propagates this layout throughout the pipeline.
UE8M0 format conversion	When use_ue8m0 is true, scales must be stored as 8-bit unsigned exponents. ceil_to_ue8m0, pack_ue8m0_to_int, and unpack_ue8m0_from_int handle the conversion and packing/unpacking.
Mixed-precision GEMM kernels	The legacy and mega-MOE GEMM implementations call these utilities to prepare scaling factors before launching low-precision kernels, as seen in the test suite (tests/test_mega_moe.py, tests/test_attention.py).

Implementation Details and Usage Patterns

Shape Calculation for Scaling Factors

The get_sf_shape function in third-party/tilelang_ops/utils.py determines the correct tensor dimensions for storing scaling factors. It accounts for UE8M0 constraints and ensures 16-byte alignment required by Tile-Lang kernels.

from third_party.tilelang_ops.utils import get_sf_shape

num_tokens = 64
hidden_dim = 256
per_channel = 64

# Calculate shape for column-major UE8M0 scales

shape = get_sf_shape(
    num_tokens, 
    hidden_dim, 
    per_channel,
    use_ue8m0=True, 
    use_col_major_sf=True
)

# Returns (num_scales, num_sf_tokens) with proper alignment

Quantization with Per-Token Scaling

The deep_gemm/utils/math.py module provides high-level casting functions that handle both the quantization and scaling factor generation.

import torch
from deep_gemm.utils.math import per_token_cast_to_fp8, per_token_cast_to_fp4

# FP32 input tensor

x = torch.randn(64, 256, dtype=torch.float32)

# ---- FP8 quantization ----

x_fp8, sf_fp8 = per_token_cast_to_fp8(
    x, 
    use_ue8m0=False, 
    gran_k=128, 
    use_packed_ue8m0=False
)

# x_fp8: torch.float8_e4m3fn

# sf_fp8: float tensor of shape (64, 2) for per-token scaling

# ---- FP4 quantization ----

x_fp4, sf_fp4 = per_token_cast_to_fp4(
    x, 
    use_ue8m0=False, 
    gran_k=128, 
    use_packed_ue8m0=False
)

# x_fp4: packed int8 tensor (half the width)

# sf_fp4: float tensor for de-quantization

UE8M0 Scale Packing

When memory bandwidth is critical, scales can be packed into the UE8M0 format using ceil_to_ue8m0 and pack_ue8m0_to_int.

from deep_gemm.utils.math import ceil_to_ue8m0, pack_ue8m0_to_int, unpack_ue8m0_from_int

scale = 1.25e-3

# Round to nearest UE8M0 representable value

ue8m0_val = ceil_to_ue8m0(scale)

# Pack into 8-bit integer container

packed = pack_ue8m0_to_int(ue8m0_val)

# Unpack when needed for computation

unpacked = unpack_ue8m0_from_int(packed)

Key Source Files

File	Description
third-party/tilelang_ops/utils.py	Defines get_sf_shape and get_sf_and_inv for core shape calculations and scale derivation.
deep_gemm/utils/math.py	Implements FP8/FP4 casting functions, UE8M0 packing/unpacking, and mathematical helpers like ceil_to_ue8m0.
third-party/tilelang_ops/swiglu_apply_weight_to_fp8.py	Demonstrates how scaling tensors are passed to Tile-Lang kernels with column-major layout requirements.
tests/test_mega_moe.py	Contains concrete usage examples of scaling factor utilities in mega-MoE GEMM implementations.
tests/test_attention.py	Shows scaling factor preparation for attention kernel testing.

Summary

• Scaling factor transformation utilities in DeepGEMM compute, reshape, and pack quantization metadata required for low-precision matrix multiplication.
• Primary functions include get_sf_shape for dimension calculation, per_token_cast_to_fp8/per_token_cast_to_fp4 for quantization, and UE8M0 packers for memory efficiency.
• Required when converting FP32 tensors to FP8 or FP4 formats, especially for hardware kernels expecting specific scaling layouts (column-major, UE8M0-packed).
• Located in third-party/tilelang_ops/utils.py and deep_gemm/utils/math.py, with usage examples in tests/test_mega_moe.py.

Frequently Asked Questions

What is the difference between per-token and per-channel scaling factors?

Per-token scaling factors compute a single scale value for each token (row) in the input tensor, resulting in a scale tensor with shape (num_tokens, num_scales_per_token). Per-channel scaling factors compute scales across the hidden dimension (columns), producing different granularity. The per_token_cast_to_fp8 and per_channel_cast_to_fp8 functions in deep_gemm/utils/math.py handle these respective modes.

When should I use UE8M0 format for scaling factors?

Use UE8M0 (Unsigned Exponent 8-bit Mantissa 0) when memory bandwidth is constrained and you need to store scales as compact 8-bit unsigned integers rather than 32-bit floats. This format is particularly valuable in large-scale MoE deployments where scaling factor tensors can become memory-bound. Set use_ue8m0=True when calling get_sf_shape or the casting functions, and use ceil_to_ue8m0 to ensure values are representable in this format.

How do I calculate the correct shape for scaling factor tensors?

Use the get_sf_shape function from third-party/tilelang_ops/utils.py. Provide the number of tokens, hidden dimension, and channel granularity (per_chan), along with boolean flags for use_ue8m0 and use_col_major_sf. The function returns a tuple (num_scales, num_sf_tokens) that respects 16-byte alignment requirements and UE8M0 layout constraints needed by Tile-Lang kernels.

Why are scaling factors necessary for FP8 and FP4 matrix multiplication?

Scaling factors preserve numerical stability when converting high-dynamic-range FP32 values into low-precision 8-bit or 4-bit formats. During the GEMM operation, the hardware uses these scales to rescale the integer matrix product back to the correct magnitude, effectively simulating higher precision arithmetic. Without properly computed and formatted scaling factors from utilities like per_token_cast_to_fp8, the quantized values would suffer from excessive rounding error and numerical drift.

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