Implementation:Sgl project Sglang Marlin Dequant

Knowledge Sources	Sgl_project_Sglang
Domains	CUDA_Kernels, Quantization, Marlin_GEMM
Last Updated	2026-02-10 00:00 GMT

Overview

GPU dequantization routines for converting quantized weights (INT4, INT8, FP4, FP8) to FP16/BF16 during Marlin GEMM kernel execution, using optimized PTX assembly instructions.

Description

The dequant.h header implements fast on-the-fly weight dequantization for the Marlin quantized GEMM kernel. The implementation uses a two-stage approach:

Stage 1 -- Bitwise operations: The lop3 template function provides a lookup-table based 3-input logical operation via PTX assembly (lop3.b32), used to extract and rearrange quantized values packed in int32 registers. The prmt template function performs byte permutation (prmt.b32) to reorder bytes between source registers. These operations extract nibbles (4-bit) or bytes (8-bit) from packed integer storage.

Stage 2 -- Floating-point computation: After bit extraction, floating-point operations complete the conversion to the target format. The file documents optimization strategies for fusing zero-point subtraction and scaling:

For INT4/INT8 with integer zero points: bias terms cancel out during subtraction, requiring no additional modification
For INT4/INT8 without zero points: scale and bias can be fused via FMA operations
For FP4/FP8: the scale factor and format multiplier can be pre-combined as scale_factor * multiplier
For INT8 to BF16: uses byte_perm which cannot be fused with scaling

The dequant template function is specialized for each combination of target type (half2 or nv_bfloat162) and weight type ID (kU4B8, kU8B128, kU4, kU8, etc.), with an optional skip_flop parameter to perform only bitwise extraction.

Usage

This header is included by marlin_template.h and used within the Marlin GEMM kernel's inner loop to dequantize weights from shared memory before tensor core MMA operations. It targets SM80+ (Ampere and newer) architectures.

Code Reference

Source Location

Repository: Sgl_project_Sglang
File: sgl-kernel/csrc/gemm/marlin/dequant.h
Lines: 1-504

Signature

namespace MARLIN_NAMESPACE_NAME {

// Lookup-table based 3-input logical operation (PTX lop3.b32)
template <int lut>
__device__ inline int lop3(int a, int b, int c);

// Byte permute operation (PTX prmt.b32)
template <int start_byte, int mask>
__device__ inline uint32_t prmt(uint32_t a);

// Main dequantization template
template <typename scalar_t2, sglang::ScalarTypeId w_type_id,
          bool skip_flop = false>
__device__ inline void dequant(int q, scalar_t2* frag_b);

// Specializations for INT4 (kU4B8) -> FP16
template <>
__device__ inline void dequant<half2, sglang::kU4B8.id(), true>(
    int q, half2* frag_b);
template <>
__device__ inline void dequant<half2, sglang::kU4B8.id(), false>(
    int q, half2* frag_b);

// Additional specializations for:
//   kU8B128 (INT8), kU4 (unsigned 4-bit), kU8 (unsigned 8-bit),
//   kFE4M3 (FP8 E4M3), kFE2M1 (FP4 E2M1)
//   targeting both half2 and nv_bfloat162

} // namespace MARLIN_NAMESPACE_NAME

Import

#include "dequant.h"

// Underlying dependency:
#include "marlin_dtypes.cuh"

I/O Contract

Inputs

Name	Type	Required	Description
q	int	Yes	Packed quantized weight values in a 32-bit integer register
w_type_id	sglang::ScalarTypeId	Yes (template)	Compile-time weight quantization type identifier (e.g., kU4B8, kU8B128)
skip_flop	bool	No (template)	If true, perform only bitwise extraction without the float conversion step

Outputs

Name	Type	Description
frag_b	scalar_t2*	Array of dequantized half2 or nv_bfloat162 values written to the B-fragment register

Usage Examples

// Inside the Marlin GEMM kernel inner loop
int q_val = shared_mem_B[offset];

// Dequantize INT4 weights to FP16 (with bias subtraction)
half2 frag_b[2];
dequant<half2, sglang::kU4B8.id(), false>(q_val, frag_b);

// Dequantize INT4 weights to FP16 (bitwise only, skip float op)
half2 frag_b_raw[2];
dequant<half2, sglang::kU4B8.id(), true>(q_val, frag_b_raw);

// Dequantize FP8 E4M3 weights to BF16
nv_bfloat162 frag_bf16[2];
dequant<nv_bfloat162, sglang::kFE4M3.id()>(q_val, frag_bf16);

Related Pages

Environment:Sgl_project_Sglang_CUDA_GPU_Runtime

Page Connections

Double-click a node to navigate. Hold to expand connections.

Principle

Implementation

Heuristic

Environment