Implementation:InternLM Lmdeploy Gemm Transform

Knowledge Sources	InternLM_Lmdeploy
Domains	GPU_Kernels, GEMM
Last Updated	2026-02-07 15:00 GMT

Overview

Implements pre-MMA data transformation functions that convert loaded operand data and optionally apply dequantization (scale + zero-point) before feeding data into the MMA instruction.

Description

This header provides transform policies applied between shared memory loading and MMA execution:

• Transform_Default: Simple identity transform that copies data from the loaded array into the MMA fragment format via reinterpret_cast with compile-time size validation.

• Transform_HMMA_16816: Dequantization transform for SM80 16x8x16 MMA. After loading quantized data, it applies per-element dequantization using group-wise scale/zero-point statistics. Supports multiple quantization formats including uint32_t packed scale+bias, uint8_t UE8M0 block scales (for both FP16 and BF16), and uint16_t scales. The dequant helper uses __hfma for fused multiply-add dequantization.

• Transform_HMMA_SIMT_B: Similar dequantization for Volta SM70 SIMT-based MMA with a simpler fragment layout.

Each transform handles the mapping between the loaded data layout (with potential packing) and the MMA instruction's expected fragment layout.

Usage

Selected as a template parameter of the GEMM operand configuration, applied in the mainloop between shared memory reads and MMA calls.

Code Reference

Source Location

• Repository: InternLM_Lmdeploy
• File: src/turbomind/kernels/gemm/transform.h

Signature

struct Transform_Default {
    template<class T, int Nf, int Mf, int K, int Nd, int Md, class S>
    __device__ static void apply(Array<T,Nf> (&frag)[K][Mf], int k,
                                  Array<T,Nd> (&data)[K][Md], S&, int div);
};

template<int StatStepS, int StatStepC>
struct Transform_HMMA_16816 {
    template<class F, int Nf, int Mf, int K, class D, int Nd, int Md, class S, int Ns, int Ms, int Ks>
    __device__ static void apply(Array<F,Nf> (&frag)[K][Mf], int k,
                                  Array<D,Nd> (&data)[K][Md],
                                  Array<S,Ns> (&stat)[Ks][Ms], int div);
};

Import

#include "src/turbomind/kernels/gemm/transform.h"

I/O Contract

Inputs

Name	Type	Required	Description
data	Array[K][Md]	Yes	Loaded quantized data from shared memory
stat	Array[Ks][Ms]	No	Quantization statistics (scale/zero-point), empty for Transform_Default
k	int	Yes	Current K iteration index
div	int	Yes	Group size divisor for quantization statistics

Outputs

Name	Type	Description
frag	Array[K][Mf]	Dequantized MMA-ready fragment in the instruction's expected format

Usage Examples

// Default transform (no dequantization)
Transform_Default::apply(frag, k, data, stat, 1);

// HMMA 16816 with dequantization
Transform_HMMA_16816<1, 2>::apply(frag, k, data, stat, group_size / CTA_K);

Related Pages

• Environment:InternLM_Lmdeploy_CUDA_GPU_Runtime