# 从零实现 C++ LLM 推理引擎：加载 Qwen2.5-7B 并跑通单层 Forward Pass

从零实现 C++ LLM 推理引擎：加载 Qwen2.5-7B 并跑通单层 Forward Pass

笔者最近在准备 LLM 推理框架方向的实习，发现从零写一个推理引擎是理解 Transformer 推理底层最直接的方式。以下用约 1000 行 C++/CUDA 搭建了一个极简推理 Runtime，涵盖 safetensors 权重加载、自定义 CUDA kernel 与 cuBLAS 矩阵运算，最终在 RTX 4090 上跑通 Qwen2.5-7B 单层 Forward，耗时 0.68ms。

为什么要自己写？

用 vLLM、SGLang 跑推理确实方便，但面试时一旦被追问到底层细节就容易答不上来。比如权重是怎么从 safetensors 文件加载到 GPU 的？一个 Transformer layer 到底需要哪些算子？cuBLAS 那套 row-major/column-major 的参数是怎么推导的？BF16 和 FP16 混用又会出什么问题？这些问题光看框架代码很难建立直觉，只有自己写一遍才能真正弄清楚。

---

架构总览

safetensors (mmap) → GpuTensor (RAII) → TransformerLayer::forward()
                                            ├── RMSNorm (custom kernel)
                                            ├── Q/K/V projection (cuBLAS)
                                            ├── RoPE (custom kernel)
                                            ├── GQA repeat_kv (custom kernel)
                                            ├── Attention scores (cublasHgemmStridedBatched)
                                            ├── Causal mask + softmax (custom kernel)
                                            ├── Output projection (cuBLAS)
                                            ├── Residual + RMSNorm
                                            └── SwiGLU MLP (cuBLAS + SiLU kernel)

项目结构：

cpp-inference-engine/
├── CMakeLists.txt              # CMake 构建 (FetchContent + CUDA Toolkit)
├── include/
│   ├── cuda_utils.cuh          # 错误检查宏 + cuBLAS RAII handle
│   ├── tensor.cuh              # FP16 GPU tensor (RAII, move-only)
│   ├── safetensors.h           # mmap 零拷贝权重加载
│   ├── kernels.cuh             # 9 个自定义 CUDA kernel
│   └── transformer.cuh         # TransformerConfig + TransformerLayer
├── src/
│   ├── tensor.cu / safetensors.cpp / kernels.cu
│   ├── transformer.cu          # 单层 forward (GQA + SwiGLU)
│   └── main.cpp                # CLI + benchmark

---

关键实现细节

1. Safetensors 解析：mmap + BF16 自动转换

Safetensors 格式很简洁：前 8 字节是 header 长度（little-endian uint64），随后是 JSON header，记录每个 tensor 的名字、dtype 和偏移量，再往后是连续排布的 tensor 数据。

uint64_t hdr_len = 0;
std::memcpy(&hdr_len, base, 8);            // 前 8 字节
auto j = nlohmann::json::parse(base + 8);  // JSON header
data_start_ = 8 + hdr_len;                 // 数据起始偏移

加载时直接用 mmap() 映射文件，跳过 read() 的内存拷贝。大模型权重动辄数 GB，这一步优化的收益显而易见。

需要特别注意的是 BF16 与 FP16 的区别。Qwen2.5 的 safetensors 存储的是 BF16（8 位指数 + 7 位尾数），而非 FP16（5 位指数 + 10 位尾数）。笔者一开始未注意 dtype 字段，读出的数据全为 NaN，排查了较长时间才定位到原因。最终在 host 端做 BF16→FP32→FP16 的两步转换解决了这个问题：

if (m.dtype == "BF16") {
    for (size_t i = 0; i < numel; ++i) {
        // BF16 → FP32：左移 16 位（BF16 本质上就是 FP32 的高 16 位）
        uint32_t fp32_bits = static_cast<uint32_t>(bf16_data[i]) << 16;
        // FP32 → FP16：提取 sign/exp/mantissa 并截断
        uint16_t sign = (fp32_bits >> 16) & 0x8000;
        int exp = ((fp32_bits >> 23) & 0xFF) - 127 + 15;
        uint16_t mant = (fp32_bits >> 13) & 0x03FF;
        if (exp <= 0)       h = sign;           // 下溢 → 0
        else if (exp >= 31) h = sign | 0x7C00;  // 上溢 → inf
        else                h = sign | (exp << 10) | mant;
    }
}

2. 自定义 CUDA Kernels

整个项目写了 9 个 kernel。以 RMSNorm 为例，它用 shared memory 做 parallel reduction，每个 block 负责一行：

__global__ void rms_norm_k(half* out, const half* x, const half* w,
                           int H, float eps) {
    extern __shared__ float sm[];
    float local = 0.f;
    for (int i = threadIdx.x; i < H; i += blockDim.x) {
        float v = __half2float(x[blockIdx.x * H + i]);
        local += v * v;
    }
    sm[threadIdx.x] = local;
    __syncthreads();
    // Reduction...
    float rms = rsqrtf(sm[0] / H + eps);
    for (int i = threadIdx.x; i < H; i += blockDim.x)
        out[i] = __float2half(val * rms * __half2float(w[i]));
}

RoPE 的实现稍微复杂一些：Q 和 K 的 head 数不同，这里用 unified indexing 将两者合并到同一个 kernel 中处理，省去分开 launch 的开销。

3. GQA（Grouped-Query Attention）

Qwen2.5-7B 采用 GQA：28 个 Q heads 共享 4 个 KV heads，比例为 7:1。实现上，先对 Q/K/V 投影结果做 transpose，变换到 [heads, seq_len, head_dim] 布局；接着用 repeat_kv kernel 将 KV 从 4 heads 扩展到 28 heads；最后调用 cublasHgemmStridedBatched 一次性计算所有 head 的 attention。

cublasHgemmStridedBatched(cublas_.get(),
    CUBLAS_OP_T, CUBLAS_OP_N,
    S, S, D, &alpha,
    Kt.data(), D, S * D,
    Qt.data(), D, S * D,
    &beta,
    scores.data(), S, S * S,
    n_heads);               // batch = 28 heads

4. cuBLAS Row-Major 陷阱

cuBLAS 默认 column-major，而 C++ 代码通常用 row-major，二者之间的转换是实际编码中最容易出错的地方。核心思路是利用一个数学等价关系：row-major 下的 C = A @ B^T 等价于 column-major 下的 C_col = B_col^T @ A_col，据此即可推导出正确的 cuBLAS 参数：

// Row-major: out[M,N] = input[M,K] @ W[N,K]^T
cublasHgemm(handle,
    CUBLAS_OP_T, CUBLAS_OP_N,
    N, M, K, &alpha,
    W, K,      // W[N,K] row → [K,N] col, lda=K
    input, K,  // input[M,K] row → [K,M] col, ldb=K
    &beta,
    out, N);   // out[M,N] row → [N,M] col, ldc=N

第一次推导时绕了不少弯路，但理解之后再看各开源框架中的 GEMM 调用，就不再觉得参数设置"反直觉"了。这个推导在面试中也经常被考到。

5. RAII 内存管理

GPU 资源一律用 RAII 托管，整个项目没有一处手动 cudaFree：

class GpuTensor {
    half* data_ = nullptr;
    bool owned_ = false;
public:
    GpuTensor(std::vector<int> shape) {
        CUDA_CHECK(cudaMalloc(&data_, numel() * sizeof(half)));
        owned_ = true;
    }
    ~GpuTensor() { if (owned_ && data_) cudaFree(data_); }
    GpuTensor(GpuTensor&& o) noexcept;        // move-only
    GpuTensor(const GpuTensor&) = delete;
};

---

运行结果

环境：RTX 4090 / CUDA 12.6 / Qwen2.5-7B-Instruct (BF16)

=== Tiny C++ Inference Engine ===
Config: hidden=3584, heads=28, kv_heads=4, head_dim=128
Loaded 9/9 weight tensors.

=== Results ===
Avg forward (1 layer): 0.68 ms
Output shape: [4, 3584]
Output[0, :8] = -0.1721 -0.1903 -0.3088 -0.3022 -0.1097 0.1157 -0.4243 0.0133

单层 forward 耗时 0.68ms，数值输出正常（无 NaN / Inf），可以确认 CUDA kernel 和 cuBLAS 的调用逻辑是正确的。

---

性能阶梯：从 Naive CUDA 到 Triton 到 cuBLAS

同期还做了 CUDA GEMM 两级优化与 Triton autotuned GEMM 的对比实验，用于建立从底到顶的性能参照：

实现	TFLOPS	占 cuBLAS %	代码量
Naive CUDA	~2	0.9%	~100 行 .cu
Tiled CUDA (shared memory)	~4	1.7%	~150 行 .cu
Triton DSL (autotuned)	163.5	96.7%	~40 行 Python
cuBLAS (vendor library)	169.1	100%	API 调用

Triton 仅用 40 行 Python 即可达到 cuBLAS 96.7% 的性能，主要得益于 @triton.autotune 对 BLOCK_M/N/K 和 pipeline stages 的自动搜索。

---

几点体会

safetensors 格式对 C++ 非常友好：JSON header 可直接用 nlohmann/json 解析，数据区通过 mmap 映射即可，几乎不需要额外的序列化库。BF16/FP16 混淆则是一个在生产环境中相当常见的问题，不查 dtype 就直接读取数据，NaN 的来源可能需要排查很长时间。

cuBLAS 的 row/column-major 转换在理解之前确实令人困惑，但推导清楚之后，再读各主流推理框架的 GEMM 调用代码就能看懂参数设置的逻辑了。对 GQA 而言，将 MHA → MQA → GQA 的演进脉络串起来理解，也就自然明白了 GQA 为何对 KV Cache 更友好。此外，用 RAII 管理 GPU 资源应该算是 C++ GPU 编程的基本功，否则很容易遇到 OOM 或资源泄漏的问题。

---

完整代码：GitHub - cpp-inference-engine

相关实验数据：AWQ INT4 量化 15.3→5.6GB（节省 63%）、vLLM TP=2 tensor parallel 验证、CUDA GEMM 两级优化 + Nsight Profiling