# How weight_dequant Handles Per-Block Scaling in DeepSeek-V3 FP8 GEMM

> Explore how weight_dequant manages per-block scaling for FP8 GEMM in DeepSeek-V3. Understand its custom Triton kernel for applying unique scaling factors to weight tiles.

- Repository: [DeepSeek/DeepSeek-V3](https://github.com/deepseek-ai/DeepSeek-V3)
- Tags: internals
- Published: 2026-02-26

---

**The `weight_dequant` function converts FP8-quantized weights back to full-precision FP32 by applying unique scaling factors to each `block_size × block_size` tile (default 128×128) through a custom Triton kernel that fetches per-block scales from a companion 2-D scale tensor.**

In the DeepSeek-V3 inference pipeline, the `weight_dequant` function serves as the critical bridge between compressed FP8 weights and full-precision computation. According to the DeepSeek-V3 source code, this utility decompresses quantized matrices using a per-block scaling strategy that preserves dynamic range across individual tiles. Understanding this mechanism is essential for developers working with the FP8 GEMM pathways or converting quantized checkpoints back to higher precision formats.

## Per-Block Scaling Architecture

The per-block scaling approach divides the weight matrix into a grid of `block_size × block_size` tiles, where each tile maintains its own scaling factor. This strategy allows the FP8 format to represent a wider dynamic range than a single global scale would permit.

- **Block dimensions**: Default 128×128 elements, configurable via the `block_size` parameter.
- **Scale tensor layout**: 2-D tensor of shape `(M // block_size, N // block_size)` stored in the companion buffer `s`.
- **Memory access pattern**: Each GPU block loads exactly one scale value corresponding to its spatial position in the output matrix.

## Kernel Implementation: weight_dequant_kernel

The actual dequantization logic resides in `weight_dequant_kernel` located in [`inference/kernel.py`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py) (lines 60-87). This Triton kernel executes on a 2-D grid where each program instance handles one spatial block of the output matrix.

### Grid Mapping and Scale Lookup

The kernel uses a 2-D indexing scheme where `pid_m` indexes the block row and `pid_n` indexes the block column:

```python
pid_m = tl.program_id(axis=0)
pid_n = tl.program_id(axis=1)
n = tl.cdiv(N, BLOCK_SIZE)
s = tl.load(s_ptr + pid_m * n + pid_n)  # Per-block scale fetch

```

The scale pointer arithmetic `s_ptr + pid_m * n + pid_n` computes the linear index into the 2-D scale tensor, where `n` represents the number of column blocks (`ceil(N / BLOCK_SIZE)`).

### Dequantization Computation

After loading the quantized FP8 values and converting to FP32, the kernel applies the block-specific scale:

```python
x = tl.load(x_ptr + offs, mask=mask).to(tl.float32)
y = x * s  # Dequantization: fp32_value = fp8_value * block_scale

tl.store(y_ptr + offs, y, mask=mask)

```

This multiplication happens entirely in register space on the GPU, ensuring high-throughput dequantization without intermediate host transfers.

## Python Wrapper and Launch Configuration

The host-side `weight_dequant` function (lines 89-110 in [`inference/kernel.py`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py)) validates input dimensions, allocates the output tensor, and configures the kernel launch grid:

```python
M, N = x.size()
grid = lambda meta: (triton.cdiv(M, meta['BLOCK_SIZE']),
                     triton.cdiv(N, meta['BLOCK_SIZE']))
weight_dequant_kernel[grid](x, s, y, M, N, BLOCK_SIZE=block_size)

```

The grid lambda creates a 2-D launch configuration where the number of blocks in each dimension equals the ceiling division of the matrix dimensions by the block size. This ensures complete coverage of the `M × N` weight matrix with properly sized tiles.

## Relationship to FP8 GEMM Operations

While `weight_dequant` explicitly reconstructs FP32 weights, the broader FP8 GEMM pathway in DeepSeek-V3 uses the same per-block scaling philosophy internally. The `fp8_gemm` function accepts separate scaling tensors (`a_s` for activations, `b_s` for weights) that follow identical per-block layouts.

**Key distinction**: The GEMM kernel fuses the scaling multiplication directly into the matrix multiplication, avoiding the explicit materialization of full-precision weights. The `weight_dequant` function becomes necessary when:
- Loading FP8 checkpoints for CPU fallback operations
- Converting weights to BF16 or FP32 for unsupported layer types  
- Debugging or analyzing weight distributions outside the quantized pathway

## Practical Usage Example

When working with DeepSeek-V3 FP8 checkpoints, you can reconstruct full-precision weights using the following pattern:

```python
import torch
from inference.kernel import weight_dequant

# FP8 weight matrix (M×N) with per-block scales

fp8_weight = torch.randn(1024, 1024, dtype=torch.float8_e4m3fn, device='cuda')

# Scale tensor: one value per 128×128 block (1024/128 = 8)

scale = torch.randn(8, 8, dtype=torch.float32, device='cuda')

# Dequantize to FP32

fp32_weight = weight_dequant(fp8_weight, scale, block_size=128)

# Use in standard PyTorch operations

x = torch.randn(1, 1024, device='cuda')
output = torch.nn.functional.linear(x, fp32_weight)

```

For production inference using the optimized `fp8_gemm` path, pass the quantized weights and scale tensors directly rather than calling `weight_dequant`, as this preserves memory bandwidth and keeps computation on-device.

## Summary

- **Per-block granularity**: The `weight_dequant` function processes FP8 weights in `block_size × block_size` tiles (default 128), with each tile using an independent scale factor stored in a 2-D companion tensor.
- **Kernel architecture**: The Triton implementation in [`inference/kernel.py`](https://github.com/deepseek-ai/DeepSeek-V3/blob/main/inference/kernel.py) uses a 2-D grid launch where each program ID maps to a specific spatial block, fetching scales via `s_ptr + pid_m * n + pid_n`.
- **Dequantization formula**: Full-precision reconstruction follows `FP32 = FP8 × block_scale`, executed in-register for maximum throughput.
- **GEMM relationship**: While `weight_dequant` materializes FP32 weights for general use, the native FP8 GEMM path fuses these scaling operations directly into the matrix multiplication kernel using identical per-block layouts.

## Frequently Asked Questions

### What is the default block size in weight_dequant?

The default block size is **128**, meaning each scaling factor applies to a 128×128 tile of the weight matrix. This parameter is configurable via the `block_size` argument in the `weight_dequant` function, though 128 represents the standard configuration used in DeepSeek-V3 checkpoints.

### How does the scale tensor shape relate to the weight matrix dimensions?

The scale tensor has shape `(M // block_size, N // block_size)` where `M` and `N` are the dimensions of the FP8 weight matrix. For example, a 1024×1024 weight matrix with block size 128 requires an 8×8 scale tensor, with each element representing the scaling factor for the corresponding spatial block.

### Why use per-block scaling instead of a single global scale?

Per-block scaling **preserves dynamic range across different regions of the weight matrix**. Neural network weights often exhibit varying magnitude distributions across different output channels or spatial positions. By assigning independent scales to each 128×128 block, the FP8 format can represent both small and large weight values with higher fidelity than a single global quantization scale would allow.

### When should I use weight_dequant versus the native fp8_gemm?

Use `weight_dequant` when you need **explicit access to full-precision weights**, such as for CPU inference, exporting to formats requiring FP32/BF16, or analyzing weight statistics. Use the native `fp8_gemm` path for production GPU inference, as it fuses the scaling operations into the matrix multiplication without materializing the full-precision tensor, preserving memory bandwidth and computational efficiency.