海光Z100 DCU适配vLLM实录（三）：W4A16大M Prefill调优、HIP Attention长上下文与TaoToken配置骨架

1. 海光Z100上W4A16大M Prefill为什么必须换执行结构海光Z100 DCUgfx906上跑vLLMW4A16量化模型在Decode阶段表现尚可但Prefill一上来就会暴露一个结构性问题大M矩阵乘和M1解码根本不是同一类计算。W4A16指4比特权重、16比特激活权重长期以INT4压缩格式驻留显存运行时中间激活保持FP16或BF16。Decode时M1每步只读一遍权重INT4省显存带宽的收益非常直接Prefill时M变成几百到几千同一份INT4权重被不同M分块反复读取、反复解包省下来的显存流量全被重复位操作和格式转换吃掉。我实测下来2K Prefill时W4A16一度占据80%以上GPU内核时间8K仍有约67%。继续扫描Triton的BLOCK_M、BLOCK_N、BLOCK_K、num_warps只能拿到25%48%的局部改善Profiler明确指出瓶颈不在参数而在执行结构本身。最终采用的路径是大M Prefill开始计算某个线性投影时先把当前层INT4权重一次性反量化成临时FP16权重让这一层所有token共同复用再交给成熟的FP16稠密GEMM。临时展开只发生在当前层执行期间不会在模型加载时把整个27B模型永久转成FP16。这条路径在M越大时收益越明显。QKV投影M8192时原融合W4A16耗时125.677毫秒一次反量化加FP16稠密GEMM只需35.979毫秒加速3.49倍Gate/Up投影M8192加速3.53倍Down投影M8192加速3.58倍。放回Qwen3.6-27B完整模型后8K Prefill从151.5 tokens/s提升到297.8 tokens/s提升96.5%2K从185.0提升到475.2 tokens/s提升156.8%。这里有一个容易踩的坑临时FP16权重必须确认不会在HIP Graph捕获后变成每层永久驻留的副本。我们用M512、M2048、M8192做Changing-input Graph测试每轮改变Activation和Scale捕获、回放、释放后输出逐位一致工作区正常回收才把它设为gfx906上的默认Prefill方案。2. TaoToken统一Key通道把Z100推理服务接进可管理入口Z100上的vLLM服务跑通以后下一步是让外部调用有一个稳定的统一入口。TaoToken提供统一的API Key和API通道适合把国产DCU推理服务、模型对话、编码Agent等场景收敛到同一套鉴权与路由体系里。官网入口是 https://taotoken.net/?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content API地址是 https://taotoken.net/api 注意API地址不加UTM参数。接入前你需要准备三样东西一个可用的TaoToken API Key、Z100上已经启动的vLLM OpenAI兼容服务地址、以及一个能发HTTP请求的客户端。TaoToken的Key在控制台创建建议按用途分Key比如一个用于模型对话验证一个用于长期编码Agent避免混用后难以排查额度与调用来源。如果你只是先验证模型是否能通用模型对话页面最直接 https://taotoken.net/models?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。如果要长期跑编码或Agent任务建议看Coding Plan https://taotoken.net/coding-plan?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。Key管理在控制台 https://taotoken.net/console?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 具体Key创建在API Keys页 https://taotoken.net/api-keys?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。接入文档在 https://taotoken.net/doc?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。注意TaoToken是统一Key与API通道不是用来替代vLLM或编辑器的。Z100上的推理仍然由vLLM执行TaoToken负责的是调用入口、鉴权和路由管理。3. 可复制配置骨架config.toml与settings.json下面这份配置骨架面向Z100 DCU vLLM W4A16 HIP Attention长上下文场景。config.toml负责服务端启动参数settings.json负责客户端调用参数。你可以直接复制后按自己的模型路径和Key替换。3.1 config.tomlvLLM服务端启动配置# config.toml - 海光Z100 DCU vLLM 服务端配置骨架 [server] host 0.0.0.0 port 8000 api_key sk-your-taotoken-key # 替换为TaoToken API Key served_model_name qwen3.6-27b-awq # 对外暴露的模型名 [model] model_path /data/models/Qwen3.6-27B-AWQ-INT4 dtype float16 # gfx906上BF16走低效路径显式FP16 quantization awq tensor_parallel_size 2 # TP2单卡32GB显存建议TP2 max_model_len 131072 # 131K长上下文 gpu_memory_utilization 0.90 [prefill] # W4A16大M Prefill结构重写当前层INT4权重一次反量化后复用 enable_dequant_prefill true dequant_workspace_qkv_mib 80 # QKV投影临时FP16工作区约80MiB dequant_workspace_gate_up_mib 170 # Gate/Up投影约170MiB dequant_workspace_down_mib 85 # Down投影约85MiB [attention] # HIP Attention长上下文适配 attention_backend hip_gfx906 hip_attention_head_size 256 # Qwen3.6标准Attention Head Size hip_attention_gqa 6 # GQA6:1 hip_attention_kv_dtype fp8 # FP8 KV Cache hip_attention_softmax_segments 64 # Decode Softmax分段数64为平台点 hip_attention_fallback triton # 未覆盖语义回退Triton [graph] enforce_eager false hip_graph_mode FULL_DECODE_ONLY # Decode走GPU图Prefill不走 chunked_prefill true max_num_batched_tokens 8192 # 131K会被拆成约16个8192 Chunk [engine] # 131K Prefill后半段单Chunk可超130秒默认300秒RPC超时会被触发 rpc_timeout_seconds 1800这份配置里最关键的三处dtype显式设为float16因为gfx906缺少原生BF16矩阵计算能力BF16会走低效路径enable_dequant_prefill开启大M Prefill结构重写hip_attention_softmax_segments设为64这是TP2下8K Decode的实测平台点继续加到128没有收益TP1下128段反而退化。3.2 settings.json客户端调用配置{ api_base: https://taotoken.net/api, api_key: sk-your-taotoken-key, model: qwen3.6-27b-awq, default_params: { temperature: 0, max_tokens: 64, stream: true, stop: null }, long_context: { max_model_len: 131072, chunked_prefill: true, prefix_cache: false }, timeout: { connect_seconds: 30, read_seconds: 1800 } }客户端read超时要和引擎RPC超时对齐。131K Prefill的TTFT约1246.6秒如果客户端read超时还是默认的60秒或120秒请求会在服务端还在算的时候就被客户端断开。4. 验证请求从8K到131K的完整链路配置写好后先验证服务是否正常启动再验证长上下文链路。4.1 服务健康检查curl -s http://127.0.0.1:8000/health # 返回 {status:ok} 表示vLLM服务已就绪4.2 通过TaoToken通道发一次8K请求curl -s https://taotoken.net/api/v1/chat/completions \ -H Authorization: Bearer sk-your-taotoken-key \ -H Content-Type: application/json \ -d { model: qwen3.6-27b-awq, messages: [{role: user, content: 用一句话说明Prefill和Decode的区别}], temperature: 0, max_tokens: 64, stream: false }返回中如果看到choices[0].message.content有正常文本说明TaoToken Key、API通道、Z100上的vLLM服务三层已经打通。4.3 131K长上下文验证131K验证不要用短Prompt线性外推必须发真实约13万token的输入。测试口径统一为单请求、64个输出token、temperature0、关闭遇到结束符提前停止、测试前后检查Prefix Cache计数只要发生缓存命中该次Prefill结果作废。服务启动后第一次请求会触发JIT编译和缓存建立冷启动数据不作为正式成绩。import json, time, requests url https://taotoken.net/api/v1/chat/completions headers { Authorization: Bearer sk-your-taotoken-key, Content-Type: application/json } # 构造约131K token的输入实际使用时替换为你的长文本 long_prompt 请阅读以下长文档并总结要点\n (测试文本。 * 40000) payload { model: qwen3.6-27b-awq, messages: [{role: user, content: long_prompt}], temperature: 0, max_tokens: 64, stream: True } start time.time() first_token_time None token_count 0 with requests.post(url, headersheaders, jsonpayload, streamTrue, timeout1800) as r: for line in r.iter_lines(): if not line: continue if line.startswith(bdata: ) and line ! bdata: [DONE]: if first_token_time is None: first_token_time time.time() token_count 1 total time.time() - start ttft first_token_time - start if first_token_time else None decode_speed token_count / (total - ttft) if ttft and total ttft else 0 print(fTTFT: {ttft:.1f}s) print(fDecode: {decode_speed:.2f} tokens/s)在Z100 TP2、约130K输入、64个输出token、无Prefix Cache命中的条件下第二次干净运行的预期结果是Prefill约105.1 tokens/sTTFT约1246.6秒Decode约13.16 tokens/s。第一次冷运行会略慢Prefill约101.9 tokens/sTTFT约1285.5秒Decode约10.46 tokens/s。4.4 8K与16K对照同一套口径下8K和16K的预期结果如下输入长度TP1 PrefillTP1 DecodeTP2 PrefillTP2 Decode8192240.8 tokens/s27.79 tokens/s398.8 tokens/s35.56 tokens/s16320197.1 tokens/s22.92 tokens/s379.1 tokens/s32.20 tokens/sTP2不会在所有场景都得到2倍收益因为每层还有All-Reduce通信。当前TP2默认只让约256KB以下的小消息走Z100自定义All-Reduce更大的消息回到RCCL因为接近1MB以后自定义方案会落后于RCCL。5. 本篇常见错排查5.1 131K请求中途退出日志没有OOM第一反应容易判断成显存不足但实际检查日志会发现GPU显存并没有分配失败真正触发的是vLLM引擎进程之间的RPC调用超时。131K Prefill被Chunked Prefill拆成约16个8192-token Chunk前半段单个Chunk约95秒后半段随着KV增长到130秒以上相邻两个Chunk累计等待超过默认300秒后上层就把请求判断成RPC超时并终止引擎。解决办法是放宽Engine RPC执行时限而不是继续给KV Cache挤显存。5.2 GQA8、Head Size 128下随机出现NaNHIP Attention在Qwen3.6上连续运行很久都正常换成GQA8、Head Size 128后随机出现10242560个NaN相同输入的错误位置每次不同。这是典型的线程竞态最终定位到LDS中的FP8查找表构造完成后缺少一次__syncthreads()。一部分Wave已经开始读取查找表时另一部分Wave负责的数据还没写完。增加同步后原本最容易随机产生NaN的场景连续5次全部正常。GPU Kernel正确性不能只依赖某个模型连续跑了很多次没出错。5.3 Qwen3.8 Prefill远慢于Qwen3.6Qwen3.8默认Activation是BF16而gfx906缺少原生BF16矩阵计算能力很多BF16运算走低效路径。局部FP16 Bridge方案在高熵Prompt的Exact Token ID测试没有稳定通过最终采用显式--dtype float16让整个运行时Activation使用FP16INT4权重仍保持4比特存储。同一版本生产路径下Qwen3.8 FP16的1K Prefill约473.8 tokens/sBF16只有203.2 tokens/s。5.4 客户端提前断开长请求131K Prefill的TTFT超过1200秒如果客户端read超时还是默认60秒或120秒请求会在服务端还在计算时就被客户端断开。settings.json里的read_seconds要和引擎rpc_timeout_seconds对齐建议都设为1800秒。5.5 自定义All-Reduce在大消息下反而变慢gfx906小消息自定义All-Reduce在小尺寸下明显快于通用通信库但Payload接近1MB以后会落后于RCCL。当前默认只让约256KB以下的小消息走自定义路径更大的消息回到RCCL。如果你把所有消息都强制走自定义All-ReduceTP2长上下文Decode会明显退化。6. 接入与排障入口Z100上的vLLM服务跑通后统一Key和API通道建议按用途分流。排障和接入相关的问题先看API Keys页确认Key状态和额度再对照接入文档检查请求格式 https://taotoken.net/api-keys?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 和 https://taotoken.net/doc?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。如果只是验证模型输出是否正常用模型对话页面最快 https://taotoken.net/models?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。长期跑编码或Agent任务建议用Coding Plan管理调用 https://taotoken.net/coding-plan?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。控制台入口在 https://taotoken.net/console?utm_sourcetaotoken_aicg_blog_endutm_mediumcsdnutm_campaignrewriteutm_content 。这套配置骨架在Z100 TP2上跑131K时最需要盯住的三个参数是rpc_timeout_seconds、hip_attention_softmax_segments和dtype。前两个决定长上下文能不能完整跑完第三个决定Qwen3.8这类默认BF16的模型会不会在gfx906上走低效路径。把这三个调对剩下的就是按Profiler的热点分布逐层往下压。