Implementation:InternLM Lmdeploy Gemm CpAsync
| Knowledge Sources | |
|---|---|
| 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);