☰
tiny-vllm性能剖析教程:Nsight Systems与Nsight Compute定位CUDA内核瓶颈
2026/10/15 16:11:42 网站建设 项目流程

【免费下载链接】tiny-vllm

Build your own high performance LLM inference engine in C++ and CUDA - a smaller version of vLLM

项目地址:https://gitcode.com/gh_mirrors/ti/tiny-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.cuRMSNorm、RoPE、Softmax、PagedAttention 等
内核声明src/kernels.cuhprefill / decode / pagedattn 三组内核
构建配置CMakeLists.txtnvcc 编译器、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

流程分三步:

  1. 采样:nsys profile包裹一次完整推理运行,生成report1.nsys-rep报告文件;
  2. 速览:nsys stats直接在终端输出内核耗时排行表,不用打开 GUI;
  3. 深看:把.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 Statisticswarp 停顿原因大量 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-vllm

CUDA 内核越界、竞争写等问题会同时表现为"结果错误 + 性能莫名下降",先用 sanitizer 排除掉,性能数据才可信 ✅

6. 常见瓶颈模式速查清单

结合 nsys 时间线 + ncu 指标,对照这张表快速归类:

现象可能瓶颈优化方向
GEMM 占比高且 FLOPs 利用率低矩阵尺寸小 / 精度未对齐合并请求增大 batch、统一 bf16 权重
decode 段大量cudaMemcpy小数据频繁搬运复用缓冲、减少 host↔device 往返
attention 内核 Memory Throughput 高KV cache 随机访问(分页)增大 block size、改善访问局部性
内核间时间线有空隙CPU 端启动开销 / 同步减少内核数量、合并小内核

7. 小结

tiny-vllm 性能剖析的标准工作流可以总结为四步:

  1. compute-sanitizer保证内核行为正确;
  2. nsys profile画出全局时间线,锁定 Top 耗时内核(nsys.sh);
  3. ncu对重点内核做单核剖析,读 Occupancy 与 Memory Throughput(ncu.sh);
  4. 对照瓶颈模式表修改 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

项目地址:https://gitcode.com/gh_mirrors/ti/tiny-vllm
点击查看免费下载

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询