在国产算力落地过程中，“跑通模型”只是第一步，“跑出性价比”才是生产落地的核心。很多团队在海光DCU上部署大模型时，常会遇到“硬件参数看着不错，实际推理吞吐上不去”的问题：盲目调batch、换框架，折腾一圈收益甚微，还可能引入稳定性风险。 本文以海光K100 DCU（gfx926，96GB HBM）上部署Qwen2-7B-Instruct BF16推理的真实优化过程为例，完整复现从瓶颈定位、分层优化到效果验证的全流程。所有优化手段均可复现，最终单卡生成吞吐从20.3 token/s提升至46.1 token/s，提升幅度超127%，同时显存占用下降30%。 ## 一、优化前：先建立可信基线，再定位核心瓶颈 性能优化最忌讳上来就改参数。在动手之前，必须先固定测试环境、测出准确基线，再逐层拆解时间，找到真正的瓶颈点，否则所有优化都是盲猜。 ### 1.1 初始环境与基线数据 测试环境统一配置： - 硬件：海光K100 DCU（96GB HBM）、双路海光CPU - 软件栈：DTK 26.04、PyTorch 2.4.0（ROCm 6.0）、vLLM 0.18.1 - 模型：Qwen2-7B-Instruct，BF16精度 - 测试负载：并发16请求，单请求max_tokens=512，固定输入prompt 初始使用默认参数启动： ```bash vllm serve Qwen2-7B-Instruct \ --dtype bfloat16 \ --max-model-len 8192 \ --max-num-seqs 64 \ --host 0.0.0.0 --port 8000 ``` 测得基线数据： | 指标 | 数值 | 说明 | | --------------------- | ------------ | ------------------ | | 平均生成吞吐 | 20.3 token/s | 全流程端到端吞吐 | | 首Token延迟 (TTFT) | 187 ms | 长Prompt下首字延迟 | | P99单Token延迟 (TPOT) | 789 ms | 长尾延迟偏高 | | 峰值显存占用 | 62.4 GB | 剩余显存空间充足 | ### 1.2 三层瓶颈定位法 我们按“端到端→算子层→硬件层”三层拆解，逐步缩小问题范围： **第一层：端到端阶段拆分** 通过vLLM内置统计接口拆分耗时，结果显示： - Decode（逐Token生成）阶段占总耗时83% - Prefill（提示词处理）阶段仅占总耗时17% - 结论：核心瓶颈在Decode阶段，优化优先级远高于Prefill。 **第二层：Kernel热点分析** 用 `rocprof`采集推理过程的内核耗时统计： ```bash rocprof --stats python benchmark.py ``` 统计结果非常明确： - rocBLAS GEMV（M=1矩阵向量乘）占Decode总耗时72% - LayerNorm、Bias+Activation等逐元素小算子合计占15% - Kernel Launch调度开销约占8% **第三层：硬件资源瓶颈** 读取硬件性能计数器判断瓶颈类型： - `MemUnitBusy`（显存控制器繁忙度）长期维持在92%以上 - `VALUUtilization`（向量计算单元利用率）仅41% - 结论：典型**访存瓶颈**，算力有大量冗余，瓶颈在HBM带宽。此时优化计算指令毫无意义，核心方向必须是减少显存读写、提升访存效率。 ## 二、分阶段优化：按投入产出比从高到低落地 优化遵循一个原则：先做零代码/低代码的配置优化，再做框架层算子优化，最后做定制化内核开发；每一步都做A/B测试，量化收益，无效优化立刻回滚。 ### 阶段一：环境与调度参数优化（零代码，综合收益+42%） 这一步成本最低、收益最高，是所有优化的优先级之首。 #### 优化1：NUMA亲和性绑定 问题：双路服务器中DCU挂载在特定NUMA节点，默认调度下CPU跨节点访问内存，增加数据传输延迟，多线程调度混乱也会拖慢数据供给。 操作：先查询设备拓扑归属，再绑定对应NUMA节点启动。 ```bash # 1. 查询DCU对应的NUMA节点 rocm-smi --showtopo # 2. 绑定对应CPU与内存节点启动 numactl --cpunodebind=0 --membind=0 \ vllm serve ... ``` 收益： - 吞吐从20.3 → 24.7 token/s，提升21.7% - 首Token延迟从187ms降至152ms - 零代码改动，纯配置优化，是所有优化中投入产出比最高的一项。 #### 优化2：FP8 KV缓存 + 提升并发上限 问题：BF16格式的KV缓存占用大量显存，限制了最大并发数，导致批处理效率上不去；而KV缓存的精度对生成质量影响极小，有充足的压缩空间。 操作：开启FP8 KV缓存，同步提升最大并发序列数： ```bash vllm serve ... \ --kv-cache-dtype fp8 \ --max-num-seqs 128 ``` 原理：KV缓存从BF16（2字节）压缩为FP8（1字节），显存占用直接减半；相同显存下能承载更多并发请求，提升批处理粒度，摊薄Kernel Launch开销。 收益： - 吞吐从24.7 → 28.9 token/s，提升17.0% - 峰值显存从62.4GB降至41.8GB，节省33% - 生成质量无感知下降，困惑度波动<0.2%。 #### 优化3：运行时内存策略调优 问题：默认的PyTorch缓存分配器会产生大量显存碎片，高并发下频繁触发内存整理，推高长尾延迟。 操作：添加环境变量优化内存分配行为： ```bash export PYTORCH_HIP_ALLOC_CONF=max_split_size_mb:128,garbage_collection_threshold:0.8 export HIP_FORCE_PERSISTING_L2=1 ``` 收益： - P99单Token延迟从789ms降至672ms，长尾延迟改善明显 - 平均吞吐微增3.2%，核心收益在运行稳定性。 **阶段一合计**：吞吐从20.3 → 28.9 token/s，累计提升42.4%，全程无代码改动。 ### 阶段二：算子融合与访存优化（轻量开发，综合收益+35%） 定位到GEMV和小算子是主要耗时点后，这一阶段针对热点算子做定制优化，核心目标是减少显存读写次数、降低启动开销。 #### 优化4：高频逐元素算子融合 问题：Decode阶段每一步都要执行Bias+GELU、RMSNorm等多个逐元素算子，每个算子都要读写一次显存，还带来多次Kernel Launch开销；而这些算子计算量极小，完全可以合并成一个内核，只做一次显存读写。 操作：实现融合Bias+GELU、融合RMSNorm两个自定义HIP算子，替换原生PyTorch实现。 以融合Bias+GELU为例，核心内核代码： ```cpp __global__ void fused_bias_gelu_bf16_kernel( const __nv_bfloat16* input, const __nv_bfloat16* bias, __nv_bfloat16* output, int rows, int hidden) { int col = blockIdx.x * blockDim.x + threadIdx.x; int row = blockIdx.y; if (row >= rows || col >= hidden) return; int idx = row * hidden + col; float val = __bfloat162float(input[idx]) + __bfloat162float(bias[col]); constexpr float kSqrt2OverPi = 0.79788456f; constexpr float kCoeff = 0.044715f; float cubic = val * val * val; float res = 0.5f * val * (1.0f + tanhf(kSqrt2OverPi * (val + kCoeff * cubic))); output[idx] = __float2bfloat16(res); } ``` 收益： - 逐元素算子总耗时下降约55% - 吞吐从28.9 → 33.2 token/s，提升14.9% - 同时减少了Kernel Launch次数，降低调度开销。 #### 优化5：定制化GEMV内核优化 问题：Decode阶段的Linear层退化为M=1的GEMV（矩阵向量乘），通用rocBLAS库为了通用性做了大量分支判断，启动开销大；而模型的hidden size是固定的（Qwen2-7B为4096），完全可以做深度定制优化。 优化思路： 1. 针对固定hidden size去掉通用分支，减少判断开销； 2. 优化权重读取模式，保证连续合并访问，打满带宽； 3. 调整分块参数（RPB=2），平衡并行度与访存效率； 4. 使用wave内归约替代全局同步，降低同步开销。 核心代码片段（简化版）： ```cpp // 针对 M=1, N=4096, K=4096 的定制 GEMV 内核 __global__ void custom_gemv_bf16_kernel( const __nv_bfloat16* weight, const __nv_bfloat16* input, __nv_bfloat16* output, int N, int K) { int row = blockIdx.x * 2 + (threadIdx.y & 1); float acc = 0.0f; // 分块遍历K维，连续合并读取权重 for (int k = threadIdx.x; k < K; k += 64) { float w = __bfloat162float(weight[row * K + k]); float x = __bfloat162float(input[k]); acc += w * x; } // wave内快速归约 for (int i = 32; i > 0; i >>= 1) { acc += __shfl_xor(acc, i); } if (threadIdx.x == 0) { output[row] = __float2bfloat16(acc); } } ``` 收益： - 单GEMV内核耗时下降38% - 吞吐从33.2 → 39.1 token/s，提升17.8% - 这是Decode阶段收益最大的单点优化。 #### 优化6：Prefill阶段FlashAttention适配 问题：Prefill阶段虽然占比不高，但长上下文场景下仍是瓶颈；默认Attention实现访存效率低，没有充分利用DCU的硬件特性。 操作： 1. 启用厂商适配的FlashAttention内核； 2. 调整KV缓存page size为64的整数倍，适配DCU wave64的访问粒度，保证访存合并。 收益： - 8K上下文Prefill耗时下降27% - 长文本场景首Token延迟改善明显，普通短文本场景收益约3.1%。 **阶段二合计**：吞吐从28.9 → 39.1 token/s，累计提升92.1%。 ### 阶段三：调度与流水线优化（工程调优，综合收益+18%） 算子优化到一定程度后，继续抠内核的边际收益会快速递减，转而从调度和流水线层面提升整体效率，性价比更高。 #### 优化7：连续批处理调度策略调优 问题：默认调度策略偏向公平性，高并发下会频繁切换请求，降低批处理效率；纯吞吐优先场景可以调整调度策略，提升批量粒度。 操作： ```bash vllm serve ... \ --scheduling-policy fcfs \ --max-num-batched-tokens 8192 ``` 原理：FCFS（先来先服务）调度+增大批处理Token数，让更多请求凑成一批执行，提升硬件利用率，摊薄调度开销。 收益： - 高并发下吞吐从39.1 → 42.5 token/s，提升8.7% - 代价是单请求平均延迟略有上升，适合离线批量、高吞吐优先场景。 #### 优化8：前缀缓存（Prefix Caching） 问题：业务场景中大量请求带有相同的系统提示词，每次都重复计算Prefill阶段的前缀，浪费算力和带宽。 操作：开启前缀缓存功能： ```bash vllm serve ... \ --enable-prefix-caching ``` 收益： - 多轮对话、固定系统提示场景，Prefill耗时平均下降60% - 整体吞吐提升约8.5%，具体收益取决于业务中前缀的重复率。 **阶段三合计**：吞吐从39.1 → 46.1 token/s，累计提升127.1%。 ## 三、最终效果与正确性校验 ### 3.1 全阶段数据对比 | 指标 | 优化前基线 | 阶段一完成 | 阶段二完成 | 最终优化后 | 总提升幅度 | | ------------------- | ---------- | ---------- | ---------- | ---------- | ---------- | | 平均吞吐 (token/s) | 20.3 | 28.9 | 39.1 | 46.1 | +127.1% | | 首Token延迟 (ms) | 187 | 152 | 141 | 136 | -27.3% | | P99单Token延迟 (ms) | 789 | 672 | 598 | 563 | -28.6% | | 峰值显存占用 (GB) | 62.4 | 41.8 | 43.2 | 43.7 | -30.0% | ### 3.2 三层正确性校验 性能优化永远不能以牺牲正确性为代价，每一步优化都做了三层校验： 1. **算子级**：单算子输出与原生实现对比，最大绝对误差<1e-5，符合BF16精度范围； 2. **模型级**：固定随机种子、固定输入，对比生成文本的Token重合率>99%； 3. **业务级**：在1000条测试集上，答案正确率波动<0.5%，无感知质量下降。 ### 3.3 生产级最终启动脚本 整合所有优化后的完整启动脚本，可直接参考使用： ```bash #!/bin/bash source /opt/dtk-26.04/env.sh # 运行时内存与编译优化 export PYTORCH_HIP_ALLOC_CONF=max_split_size_mb:128,garbage_collection_threshold:0.8 export HIP_FORCE_PERSISTING_L2=1 export TRITON_HIP_LLD_PATH=/opt/dtk-26.04/llvm/bin/ld.lld MODEL_PATH="/data/models/Qwen2-7B-Instruct" # NUMA绑核 + 全参数启动 numactl --cpunodebind=0 --membind=0 \ vllm serve "$MODEL_PATH" \ --served-model-name Qwen2-7B-Instruct \ --dtype bfloat16 \ --kv-cache-dtype fp8 \ --max-model-len 8192 \ --max-num-seqs 128 \ --max-num-batched-tokens 8192 \ --scheduling-policy fcfs \ --enable-prefix-caching \ --enforce-eager \ --gpu-memory-utilization 0.88 \ --host 0.0.0.0 \ --port 8000 ``` ## 四、实战避坑总结 1. **别上来就盲目调大batch**：访存瓶颈下，batch越大需要搬运的数据越多，反而容易撞上带宽天花板，收益甚微；先定位瓶颈类型，再对症下药。 2. **正确性永远优先于性能**：每做一项优化，必须先过数值校验，再看性能；静默错误比速度慢更可怕，上线后出问题排查成本极高。 3. **不要照搬CUDA优化经验**：DCU是wave64调度，共享内存32个bank，和CUDA的warp32架构有差异；直接套用CUDA的分块、归约逻辑，轻则性能差，重则出现数值错误。 4. **性能测试必须规范**：预热≥20次，重复≥100次取中位数；测试时独占设备，避免其他进程干扰；不同