Implementation:InternLM Lmdeploy Gemm MainloopSm70

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

Overview

Software-pipelined GEMM mainloop for Volta (SM70) that alternates between register-based global memory fetches and shared memory MMA computation using double-buffered shared memory.

Description

The SM70 mainloop implements a classic double-buffered software pipeline for Volta GPUs. Since Volta lacks cp.async, data must first be loaded into registers from global memory, then stored to shared memory, with explicit __syncthreads() barriers between stages.

The pipeline structure uses SmemIter (pointer ping-pong for double buffering) and GroupIter (tracks when quantization scale operands U/V need to advance). Each K-tile iteration:

1. Fetches the next tile from global memory into register fragments
2. Synchronizes threads
3. Stores register fragments to shared memory
4. Loads MMA operands from shared memory and performs tiled MMA
5. Advances iterators

Supports quantization with grouped scales (U for A-operand, V for B-operand) that advance at configurable group sizes.

Usage

Used by GemmUniversal when targeting SM70 (Volta V100) GPUs.

Code Reference

Source Location

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

Signature

template<class Impl_>
struct MainloopSm70 {
    using Impl = Impl_;
    using FragC = typename Impl::FragC;

    __device__ void operator()(GmemIterA& gmem_A, GmemIterB& gmem_B,
                                GmemIterU& gmem_U, GmemIterV& gmem_V,
                                FragC& frag_C, int tile_iter,
                                SharedStorage& storage);
};

Import

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

I/O Contract

Inputs

Name	Type	Required	Description
gmem_A, gmem_B	GmemIterator	Yes	Global memory iterators for A and B operands
gmem_U, gmem_V	GmemIterator	Yes	Scale factor iterators
tile_iter	int	Yes	Number of K-tile iterations to execute
storage	SharedStorage	Yes	Shared memory for double buffering

Outputs

Name	Type	Description
frag_C	FragC&	Accumulated MMA result fragments

Usage Examples

// Inside GemmUniversal::operator()
MainloopSm70<Impl> mainloop{};
mainloop(gmem_A, gmem_B, gmem_U, gmem_V, frag_C, tile_iter, storage.mainloop);

Related Pages

• Environment:InternLM_Lmdeploy_CUDA_GPU_Runtime