Implementation:InternLM Lmdeploy Gemm Format

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

Overview

Implements device-side data format conversion functors (Converter) for transforming between different quantized data types used by GEMM operands, including uint16-to-uint4 packing and uint16-to-uint8 conversion with byte permutation.

Description

The Converter template provides specialized conversion between storage types:

• Converter<T, T>: Identity conversion (pass-through)

• Converter<uint16_t, uint4_t>: Packs 16-bit values into 4-bit representation. Uses a two-step process: first truncates to 8-bit, then uses __byte_perm to pack pairs of 4-bit values into bytes. Operates in groups of 8 elements.

• Converter<uint16_t, uint8_t>: Truncates 16-bit values to 8-bit with a 3120 byte reordering pattern (elements 0,2,1,3) per group of 4, matching the MMA instruction's expected operand layout.

All conversions operate on Array<T, N> containers and use PRAGMA_UNROLL for compile-time loop unrolling.

Usage

Used by the layout conversion pipeline and weight preprocessing to transform quantized weights into the packed format expected by specific MMA instruction variants.

Code Reference

Source Location

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

Signature

template<class Tin, class Tout>
struct Converter {};

template<>
struct Converter<uint16_t, uint4_t> {
    static __device__ Array<uint4_t, 8> pack(const Array<uint8_t, 8>& vi);
    template<class U, int N>
    __device__ Array<uint4_t, N> operator()(const Array<U, N>& x);
};

template<>
struct Converter<uint16_t, uint8_t> {
    template<int N>
    __device__ Array<uint8_t, N> operator()(const Array<uint16_t, N>& x);
};

Import

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

I/O Contract

Inputs

Name	Type	Required	Description
x	Array<Tin, N>	Yes	Input data in source format

Outputs

Name	Type	Description
result	Array<Tout, N>	Converted data in target format

Usage Examples

Converter<uint16_t, uint4_t> conv;
Array<uint4_t, 8> packed = conv(input_u16);  // Pack 16-bit to 4-bit

Related Pages

• Environment:InternLM_Lmdeploy_CUDA_GPU_Runtime