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

> Explore DeepGEMM's scaling factor transformation utilities for FP8 FP4 quantization. Learn how these helpers convert FP32 tensors into hardware-optimized low-precision formats with quantization metadata.

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

---

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

```python
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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/utils/math.py) module provides high-level casting functions that handle both the quantization and scaling factor generation.

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

```python
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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/tests/test_mega_moe.py) | Contains concrete usage examples of scaling factor utilities in mega-MoE GEMM implementations. |
| [`tests/test_attention.py`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/third-party/tilelang_ops/utils.py) and [`deep_gemm/utils/math.py`](https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/utils/math.py), with usage examples in [`tests/test_mega_moe.py`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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`](https://github.com/deepseek-ai/DeepGEMM/blob/main/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.