在国产算力落地过程中，“跑通模型”只是第一步，“跑出性价比”才是生产落地的核心。很多团队在海光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

初始使用默认参数启动：

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采集推理过程的内核耗时统计：

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节点启动。

# 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缓存，同步提升最大并发序列数：

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缓存分配器会产生大量显存碎片，高并发下频繁触发内存整理，推高长尾延迟。
操作：添加环境变量优化内存分配行为：

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为例，核心内核代码：

__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内归约替代全局同步，降低同步开销。

核心代码片段（简化版）：

// 针对 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：连续批处理调度策略调优

问题：默认调度策略偏向公平性，高并发下会频繁切换请求，降低批处理效率；纯吞吐优先场景可以调整调度策略，提升批量粒度。
操作：

vllm serve ... \
  --scheduling-policy fcfs \
  --max-num-batched-tokens 8192

原理：FCFS（先来先服务）调度+增大批处理Token数，让更多请求凑成一批执行，提升硬件利用率，摊薄调度开销。
收益：

  • 高并发下吞吐从39.1 → 42.5 token/s，提升8.7%
  • 代价是单请求平均延迟略有上升，适合离线批量、高吞吐优先场景。
优化8：前缀缓存（Prefix Caching）

问题：业务场景中大量请求带有相同的系统提示词，每次都重复计算Prefill阶段的前缀，浪费算力和带宽。
操作：开启前缀缓存功能：

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 生产级最终启动脚本

整合所有优化后的完整启动脚本，可直接参考使用：

#!/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次取中位数；测试时独占设备，避免其他进程干扰；不同
# 性能优化 # python # github # git
Logo
EazyDevelop社区

一站式 AI 云服务平台

更多推荐

cover

2026 AI 智能体平台深度横评：大厂产品怎么选，避坑要点总结

avatar EazyDevelop社区
cover

一个提示词，让 AiPy 帮我做了个能上集市的抠图智能体插件

avatar EazyDevelop社区
cover

从“人肉调度”到“自动派单”：家政公司服务流程数字化实战

avatar EazyDevelop社区