【免费下载链接】tiny-vllm
Build your own high performance LLM inference engine in C++ and CUDA - a smaller version of vLLM
tiny-vllm 是用 C++ 和 CUDA 构建的高性能 LLM 推理引擎,可以看作是 vLLM 的一个轻量学习版。本教程带你用 NVIDIA 官方的 Nsight Systems(nsys)和 Nsight Compute(ncu)对 tiny-vllm 做系统级 + 内核级两层性能剖析,快速定位 CUDA 内核瓶颈,掌握 LLM 推理引擎性能优化的完整方法论 🚀
1. 剖析对象:tiny-vllm 推理引擎长什么样
tiny-vllm 用不到几个源文件就实现了完整的 LLM 推理流程,非常适合拿来"解剖":
| 模块 | 路径 | 说明 |
|---|---|---|
| 主流程(prefill / decode) | src/main.cpp | 加载权重、批处理调度、cuBLAS 调用 |
| 全部 CUDA 内核 | src/kernels.cu | RMSNorm、RoPE、Softmax、PagedAttention 等 |
| 内核声明 | src/kernels.cuh | prefill / decode / pagedattn 三组内核 |
| 构建配置 | CMakeLists.txt | nvcc 编译器、cuBLAS 依赖 |
模型前向由embeddingGatherKernel、rmsNormKernel、ropeKernel、softmaxKernel、pagedAttentionKernel等自写 CUDA 内核,加上 cuBLAS 的cublasGemmEx矩阵乘法共同组成。性能剖析的关键问题就是:这一串内核里,谁占了大头?瓶颈在计算还是访存?
2. 剖析前准备:构建并运行 tiny-vllm
先获取代码并构建(仓库提供只读镜像):
git clone https://gitcode.com/gh_mirrors/ti/tiny-vllm cd tiny-vllm cmake -B build -G Ninja cmake --build build提示:CMakeLists.txt 中默认将
nvcc固定在/opt/cuda/bin/nvcc,如果你装在其他位置,需要改一下这两行,否则编译会失败。
运行需要把模型权重文件model.safetensors放在当前目录,并通过标准输入传入 token 序列。项目自带的 full_test.sh 就是一条完整的示例输入命令,直接执行它即可得到一次真实的 prefill + decode 全流程,这也是我们后续剖析的"测试负载"。
3. Nsight Systems:系统级时间线快速定位
nsys(Nsight Systems)回答的是宏观问题:整个程序的时间花在哪了——是 GPU 内核、内存拷贝(cudaMemcpy),还是 CPU 端等待?
项目已经把命令封装在 nsys.sh 中,核心是这两条:
nsys profile ./build/tiny-vllm --output="report1.nsys-rep" nsys stats report1.nsys-rep流程分三步:
- 采样:
nsys profile包裹一次完整推理运行,生成report1.nsys-rep报告文件; - 速览:
nsys stats直接在终端输出内核耗时排行表,不用打开 GUI; - 深看:把
.nsys-rep文件拖进 Nsight Systems 图形界面查看 CUDA 时间线。
看时间线时重点关注三块区域:
- prefill 段:输入 token 一次过完 16 层,
cublasGemmEx的 GEMM 内核通常是大头,计算密集; - decode 段:逐 token 生成,
cudaMemcpy(如 KV cache 搬运)和pagedAttentionKernel占比明显上升,访存特征更重; - 空隙:CPU 与 GPU 之间出现空白间隙,说明存在同步等待或内核启动开销。
一句话经验:先用 nsys 找出 Top 3 耗时内核,再用下一节的 ncu 对单个内核做微观剖析,这样不会陷进几百个内核的汪洋里。
4. Nsight Compute:单内核级别的瓶颈显微镜
ncu(Nsight Compute)能钻进单个内核内部,给出 SM 占用率、内存吞吐、warp 停顿等硬指标,是定位"为什么这个内核慢"的关键工具。
项目中的 ncu.sh 内容只有一行:
sudo /usr/local/cuda-12.8/bin/ncu --target-processes all ./build/tiny-vllm几个使用要点:
- 需要
sudo,因为 GPU 性能计数器默认只允许管理员读取(否则 ncu 会报权限错误); - 按实际 CUDA 版本替换路径,
which ncu可以找到正确位置; --target-processes all表示对程序派生的所有进程(包括 CUDA 运行时子进程)都采集数据;- ncu 会逐个内核重放采样,整个运行时间会比平时长很多,属正常现象。
在 Nsight Compute 界面中打开pagedAttentionKernel、rmsNormKernel、softmaxKernel等条目,重点看四个指标:
| 指标 | 含义 | 瓶颈信号 |
|---|---|---|
| Memory Throughput | 显存带宽利用率 | 高 → 访存受限,考虑减少重复读、合并访存 |
| SM Throughput | 计算单元利用率 | 低 + Memory 也低 → 可能是占用率不足或启动开销 |
| Occupancy | 同时在跑的 warp 比例 | 低 → 检查寄存器/共享内存用量,如rmsNormKernel的__shared__ float rms_vector[1024] |
| Warp State Statistics | warp 停顿原因 | 大量 Memory Latency 停顿 → 访存延迟未被掩盖 |
另一个值得用 ncu 对比观察的对象是 cuBLAS GEMM 内核。LLM 权重是 row-major 存储,而cublasGemmEx期望 column-major,tiny-vllm 通过转置技巧(CUBLAS_OP_T/CUBLAS_OP_N标志组合)避免了数据格式转换,原理图如下:
如果 ncu 里发现 GEMM 内核远未达到硬件峰值 FLOPs,除了检查矩阵尺寸和 dtype(本项目使用CUDA_R_16BF+CUBLAS_COMPUTE_32F),也可以考虑这种"零拷贝"的转置方案是否用对了。
5. 辅助工具:用 compute-sanitizer 排除内存错误
剖析前建议先做一次内存正确性检查,check.sh 已经写好:
compute-sanitizer ./build/tiny-vllmCUDA 内核越界、竞争写等问题会同时表现为"结果错误 + 性能莫名下降",先用 sanitizer 排除掉,性能数据才可信 ✅
6. 常见瓶颈模式速查清单
结合 nsys 时间线 + ncu 指标,对照这张表快速归类:
| 现象 | 可能瓶颈 | 优化方向 |
|---|---|---|
| GEMM 占比高且 FLOPs 利用率低 | 矩阵尺寸小 / 精度未对齐 | 合并请求增大 batch、统一 bf16 权重 |
decode 段大量cudaMemcpy | 小数据频繁搬运 | 复用缓冲、减少 host↔device 往返 |
| attention 内核 Memory Throughput 高 | KV cache 随机访问(分页) | 增大 block size、改善访问局部性 |
| 内核间时间线有空隙 | CPU 端启动开销 / 同步 | 减少内核数量、合并小内核 |
7. 小结
tiny-vllm 性能剖析的标准工作流可以总结为四步:
compute-sanitizer保证内核行为正确;nsys profile画出全局时间线,锁定 Top 耗时内核(nsys.sh);ncu对重点内核做单核剖析,读 Occupancy 与 Memory Throughput(ncu.sh);- 对照瓶颈模式表修改 src/kernels.cu 中的内核,再回到第 2 步回归验证。
因为整个推理引擎只有 src/main.cpp 和 src/kernels.cu 两个核心文件,你可以清楚地看到每一次优化前后 nsys 时间线的变化——这正是 tiny-vllm 作为"学习用推理引擎"的最大价值:把 LLM 推理的性能剖析从黑盒变成一道可推导的练习题 📊
【免费下载链接】tiny-vllm
Build your own high performance LLM inference engine in C++ and CUDA - a smaller version of vLLM
相关推荐
TinyNvidiaUpdateChecker:彻底告别GeForce Experience!轻量级NVIDIA驱动更新神器
TinyNvidiaUpdateChecker:彻底告别GeForce Experience!轻量级NVIDIA驱动更新神器 TinyNvidiaUpdateC
cuDF/libcudf 性能剖析实战指南:用 NVIDIA Nsight Systems 定位 GPU 瓶颈
cuDF/libcudf 性能剖析实战指南:用 NVIDIA Nsight Systems 定位 GPU 瓶颈 libcudf 是 cuDF 的 C++ GPU
数据分析数据工程机器学习Isaac Lab 性能剖析实战:用 Nsight Systems 定位 CPU/GPU 瓶颈与 NVTX Trace 配置详解
Isaac Lab 性能剖析实战:用 Nsight Systems 定位 CPU/GPU 瓶颈与 NVTX Trace 配置详解 Isaac Lab 官方支持通
人工智能强化学习机器人具身智能深度学习
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考