Implementation:InternLM Lmdeploy Gemm MatrixPtr

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

Overview

Provides matrix pointer resolution logic for flat, blocked, and indexed striding modes, translating runtime MatrixParam descriptors into device-side MatrixData with correct pointers and optional index arrays.

Description

This header defines the data structures and resolution functions that bridge the host-side matrix layout descriptions (MatrixParam) to device-side data access patterns (MatrixData):

• StridedPtr: A 16-byte aligned pair of (void* ptr, int stride)
• MatrixParam: Host-facing parameter bundle with pointer, stride, offset array, and index array
• MatrixData: Device-facing resolved pointer with optional index array for indirect access

The resolve<T, mode>() template function handles three striding modes:

• kFlat: Simple stride-based access, returns the pointer directly
• kBlocked: For blocked/batched GEMM, loads per-group StridedPtr from an array of pointers when stride==0, then applies offset-based displacement
• kIndexed: For indexed/ragged GEMM, resolves per-group index arrays and applies offset displacement

The to_param helper converts a MatrixLayout to a MatrixParam.

Usage

Called at the start of each GEMM kernel to resolve matrix operand pointers for the current group/batch element.

Code Reference

Source Location

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

Signature

struct StridedPtr { void* ptr; int stride; };
struct MatrixParam { void* ptr; int stride; int* offsets; int* idxs; };
struct MatrixData { StridedPtr ptr; const int* idxs; };

template<class T, Striding mode>
__device__ MatrixData resolve(const MatrixParam& param, int group_id);

Import

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

I/O Contract

Inputs

Name	Type	Required	Description
param	MatrixParam	Yes	Matrix pointer, stride, optional offsets/indices
group_id	int	Yes	Group/batch index for multi-group GEMM

Outputs

Name	Type	Description
MatrixData	struct	Resolved device pointer with stride and optional index array

Usage Examples

// Resolve a blocked matrix for group g
auto mat = resolve<half, Striding::kBlocked>(param, g);
// mat.ptr.ptr is the base pointer, mat.ptr.stride is the leading dimension

Related Pages

• Environment:InternLM_Lmdeploy_CUDA_GPU_Runtime