Implementation:InternLM Lmdeploy Gemm CpAsync

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

Overview

Provides inline PTX wrappers for the CUDA cp.async instruction, enabling asynchronous global-to-shared memory copies with configurable cache operations, eviction policies, and L2 cache hints.

Description

This header defines a template-based abstraction over the cp.async PTX instruction used on SM80+ GPUs. The CP_ASYNC struct is specialized for different combinations of cache operations (CacheOp::kGlobal, CacheOp::kAlways), copy sizes (4, 8, 16 bytes), and L2 prefetch hint sizes (64, 128, 256 bytes). Each specialization emits the appropriate inline assembly with predicated execution (conditional mask) and optional L2 cache policy hints.

Supporting enums include CacheOp (kDefault, kAlways, kGlobal), EvictPolicy (kEvictNormal, kEvictFirst, kEvictLast), and predefined cache policy bundles (cache_policy::Default, cache_policy::Stream, cache_policy::Reuse).

Usage

Used by the GEMM memory iterators (particularly GmemIteratorSm80) to asynchronously transfer tiles of matrix operands from global memory to shared memory without blocking the warp, hiding memory latency behind computation.

Code Reference

Source Location

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

Signature

enum class CacheOp { kDefault, kAlways, kGlobal };
enum class EvictPolicy { kEvictNormal, kEvictFirst, kEvictLast };

template<CacheOp, int size, int prefetch_size>
struct CP_ASYNC {
    __device__ static void apply(int smem_ptr, const void* src, bool mask);
    __device__ static void apply(int smem_ptr, const void* src, uint64_t cache_policy, bool mask);
};

Import

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

I/O Contract

Inputs

Name	Type	Required	Description
smem_ptr	int	Yes	Shared memory destination address (obtained via cast_smem_ptr_to_uint)
src	const void*	Yes	Global memory source pointer
mask	bool	Yes	Predicate controlling whether the copy executes
cache_policy	uint64_t	No	Optional L2 cache hint policy descriptor

Outputs

Name	Type	Description
(side effect)	shared memory	Data is asynchronously copied from global to shared memory

Usage Examples

// Issue a 16-byte async copy with global cache op and no prefetch hint
CP_ASYNC<CacheOp::kGlobal, 16, 0>::apply(smem_addr, global_ptr, true);

// Issue with L2 cache hint
CP_ASYNC<CacheOp::kAlways, 16, 128>::apply(smem_addr, global_ptr, l2_policy, mask);

Related Pages

• Environment:InternLM_Lmdeploy_CUDA_GPU_Runtime