delete old md
This commit is contained in:
root
@@ -36,7 +36,6 @@
|
||||
| schema YAML | muh/schema/*.yaml | 27 files | ✓ 完成 | 每个算法的参数空间定义 |
|
||||
| parse.py | muh/parse.py | ~150 | ✓ 基本可用 | .muh → JSON 解析（自实现 YAML parser）|
|
||||
| gen_yaml.py | muh/gen_yaml.py | ~80 | ✓ 完成 | .muh → computility-run.yaml |
|
||||
| gen_patch.py | muh/gen_patch.py | ~200 | ✓ 可提取 bi100_* | C++ header → vllm unified diff |
|
||||
| extract.py | muh/extract.py | ~100 | ✓ 完成 | CCCL tuning → schema 提取 |
|
||||
| muh_dispatch.py | muh_dispatch.py | ~400 | △ 概念完成 | CCCL-style 类型分派（未接入 vllm）|
|
||||
| muh_kernel_map.py | muh_kernel_map.py | ~350 | △ 手写常量 | 需要从 C++ headers 自动提取闭环 |
|
||||
@@ -44,98 +43,3 @@
|
||||
| test_smem_safety | muh/tests/test_smem_safety.py | 1 file | ✓ | 全算法 SMEM 安全检查 |
|
||||
|
||||
---
|
||||
|
||||
## 三、Decode 热路径 × 资产覆盖矩阵
|
||||
|
||||
```
|
||||
算法 CCCL muh schema bench test vllm注入点 竞赛权重
|
||||
───────────────── ───── ───── ────── ───── ───── ──────────────────────────────── ─────────
|
||||
reduce ✓ ✓ ✓ ✓ ✓ csrc/attention/paged_attention 83% (Output)
|
||||
scan ✓ ✓ ✓ ✓ ✓ csrc/attention/paged_attention 14% (Input)
|
||||
topk ✓ ✓ ✓ ✓ ✓ csrc/sampling/sampling_kernels per decode
|
||||
radix_sort ✓ ✓ ✓ ✓ ✓ csrc/sampling/sampling_kernels per decode
|
||||
transform ✓ ✓ ✓ ✓ ✓ csrc/activation/layernorm/rope 200×/token
|
||||
select_if ✓ ✓ ✓ △ ✓ csrc/sampling (top-p filter) per decode
|
||||
batch_memcpy ✓ ✓ ✓ △ △ csrc/cache_kernels 3% (Cache)
|
||||
for_each ✓ ✓ ✓ ✓ ✓ csrc (residual connections) per layer
|
||||
```
|
||||
|
||||
△ = benchmark/test 文件存在但名称不直接匹配（partition/if.cu 对应 select_if，copy/memcpy.cu 对应 batch_memcpy）
|
||||
|
||||
---
|
||||
|
||||
## 四、GitHub Issues 状态
|
||||
|
||||
### 已创建的 38 个 Issues（全 open，全有 labels）
|
||||
|
||||
**功能测试覆盖（#1-#16）**: 竞赛 50+ 功能测试用例的完整 PRD，每个含 PND 级 test cases 表
|
||||
|
||||
| 编号范围 | 前缀 | 数量 | 说明 |
|
||||
|----------|------|------|------|
|
||||
| #1-#14 | [FEA] | 14 | 功能测试: 非流式/流式/Tool/Reasoning/Cache/采样/结构化/多语言/多模态/校验/能力/截断/效果 |
|
||||
| #15-#16 | [EPIC] | 2 | 性能基准 + 开发环境 |
|
||||
| #17-#25 | [FEA]/[EPIC] | 9 | muh 语言设计: 语法/schema/codegen(yaml+patch+dockerfile)/bench/search/tuning提取 |
|
||||
| #26-#38 | [muh] | 13 | muh 算法标定: reduce/scan/radix_sort/select_if/scan_by_key/reduce_by_key/unique_by_key/transform/batch_memcpy/topk + gen_patch管道/benchmark runner/hardware校准 |
|
||||
|
||||
### Project/6 面板上的 Draft Issues（72 个，无 repo 关联）
|
||||
|
||||
来自后续对话生成，包含：
|
||||
- [muh] 语言规范 v1/v2
|
||||
- [muh] 20+ 个算法标定 items（adjacent_difference, batched_topk, find, histogram, merge, rle_encode 等）
|
||||
- [INFRA] CI同步/Build编译/Deploy部署/Verify回归
|
||||
- [BUG] scale_mem_bound / gen_patch 管道 / select_if 坍缩 / bytes_in_flight / reduce items
|
||||
- [CCCL-verify] 20 个 Thrust/CUB example 验证 items
|
||||
- [CCCL-test] 10 个 Catch2 测试矩阵 items
|
||||
- [muh-bench] 6 个 benchmark items
|
||||
- [muh-pipe] 端到端管道验证
|
||||
|
||||
**这 72 个 draft 需要转为真 issue。** 内容已经写好（body 含完整 test cases 表），只是缺少 repo 关联和 labels。
|
||||
|
||||
---
|
||||
|
||||
## 五、关键发现（SM count = 16）
|
||||
|
||||
Phanthy Cloud 实测确认 BI-V100 只有 **16 SMs**（不是规格书的 50c）。
|
||||
|
||||
影响范围：
|
||||
1. `hardware.cuh` — 已修正 sm_count=16
|
||||
2. `tuning_transform.cuh` — bytes_in_flight 基于 900/50=18 GB/s 已失效，应为 900/16=56 GB/s
|
||||
3. `tuning_reduce.cuh` — bi100_det_* 和 bi100_default 的 items 偏小（tile 仅用 23% SMEM）
|
||||
4. `tuning_scan.cuh` — lookback delay 基于 50 SM 的争用模型，16 SM 下争用更低、delay 可以更短
|
||||
5. 所有 benchmark 理论推导需要重跑
|
||||
|
||||
---
|
||||
|
||||
## 六、不需要 clone 更多 CCCL 的原因
|
||||
|
||||
完整 CCCL (github.com/NVIDIA/cccl) ≈ 40K 文件、1.2GB。我们有 8,900 文件 (74MB)。
|
||||
|
||||
已有的关键子集：
|
||||
- ✓ 全部 27 tuning headers（muh 从这里提取参数空间）
|
||||
- ✓ 全部 32 dispatch headers（tuning 参数化的对象）
|
||||
- ✓ 52 Thrust examples（正确性验证的 golden reference）
|
||||
- ✓ 18 CUB examples（device + block level API 验证）
|
||||
- ✓ 234 CUB Catch2 tests（回归测试矩阵）
|
||||
- ✓ 169 Thrust tests（Thrust 算法回归）
|
||||
- ✓ 78 CUB benchmarks（标定数据的来源）
|
||||
- ✓ 530 Thrust headers + 1,357 libcudacxx headers（编译依赖）
|
||||
|
||||
缺失的 ~31K 文件：
|
||||
- libcudacxx 深层 include（6K）— 编译时用 -I 指向安装路径
|
||||
- cudax 实验模块（800）— 竞赛不用
|
||||
- cmake/CI 基础设施（5K）— 平台用 Dockerfile 构建
|
||||
- Python 绑定 / 文档 / 其他（19K）— 不相关
|
||||
|
||||
---
|
||||
|
||||
## 七、信创魔盒核心差异（竞赛定位）
|
||||
|
||||
> "信创魔盒是基于系统级的架构，内置算法因子，用 EngineX 引擎把模型内部的算法因子重新置换——不是单纯的连接器。"
|
||||
|
||||
muh 在这个架构中的角色：
|
||||
- CCCL 的 `policy_selector` 是 NVIDIA 为自家 GPU 写的"算法因子"
|
||||
- muh 的 `policy_selector` 是为天垓100 写的等效"算法因子"
|
||||
- EngineX 把 CCCL 的 NVIDIA 算法因子替换成 muh 的天垓100 算法因子
|
||||
- 不是适配层（60% 精度），是置换层（目标 ≥100% 精度在天垓100 硬件约束下的最优解）
|
||||
|
||||
竞赛成绩 = 算法因子置换的精度 × 硬件实测标定的覆盖度。
|
||||
|
||||
@@ -1,120 +0,0 @@
|
||||
# CCCL ↔ muh 完整 Gap 分析
|
||||
|
||||
> 生成时间: 2026-08-06 | HEAD: 2a7ca10 | 26 算法全量扫描
|
||||
|
||||
## 核心数据
|
||||
|
||||
| 指标 | 值 | 说明 |
|
||||
|------|------|------|
|
||||
| CCCL 算法总数 | 26 | cub/device/dispatch/tuning/ 下所有 tuning_*.cuh |
|
||||
| muh tuning headers | 26 | 1:1 文件对应 ✓ |
|
||||
| CCCL 代码行 | 18,094 | 所有 tuning_*.cuh 总和 |
|
||||
| muh 代码行 | 3,568 | 19.7% 覆盖率 |
|
||||
| CCCL benchmark 注释 | 299 | `ipt_N.tpb_M ... speedup` 格式的数据点 |
|
||||
| SM100 模板特化 | 157 | NVIDIA 为 SM100 跑出的最优配置数 |
|
||||
| BI-V100 命名 struct | 37 | muh 中 `bi100_*` struct 数量 |
|
||||
| 有 bi100 struct 的算法 | 3/26 | reduce(14个), scan(22个), for(1个) |
|
||||
| 有 SMEM 保护的算法 | 16/26 | scale_mem_bound 或 while loop |
|
||||
|
||||
## 关键发现
|
||||
|
||||
### 1. 只有 reduce 和 scan 达到了"READY"状态
|
||||
|
||||
reduce 和 scan 是唯一两个同时具备 bi100 命名 struct + SMEM 保护 + 完整 policy_selector 的算法。但即便如此，这些 struct 的值全部是从 SM100 推导的**理论值**，没有一个在 BI-V100 上实测过。
|
||||
|
||||
### 2. 其余 24 个算法停留在"inline only"
|
||||
|
||||
"inline only" 意味着 muh header 里有 policy_selector，但它的值是硬编码在 if/else 分支里的，不是通过命名 struct 暴露的。gen_patch.py 提取不到这些值（它只认 `struct bi100_*` 模式）。
|
||||
|
||||
### 3. CCCL 有 299 个 benchmark 数据点，muh 有 0 个
|
||||
|
||||
CCCL 的 benchmark 注释格式完美定义了目标：
|
||||
```
|
||||
ipt_22.tpb_384.ns_1904.dcid_6.l2w_830.trp_1.ld_0 1.148442 0.997167 1.139902 1.462651
|
||||
```
|
||||
四个数字 = 四个 problem size 下的加速比。muh 需要在 BI-V100 上产出同样格式的 299 个数据点来填充所有空位。
|
||||
|
||||
### 4. 竞赛瓶颈不在代码量而在实测数据
|
||||
|
||||
- 代码架构已经搭好（26 个 header + policy_selector + gen_patch 管道）
|
||||
- 缺的是 BI-V100 实测数据来替换理论值
|
||||
- 没有实测数据，所有 bi100_* struct 的值都是猜的
|
||||
|
||||
## 26 算法状态矩阵
|
||||
|
||||
| 算法 | CCCL 行 | muh 行 | CCCL BM | SM100 特化 | bi100 struct | SMEM✓ | 状态 |
|
||||
|------|---------|--------|---------|-----------|-------------|-------|------|
|
||||
| reduce | 478 | 297 | 7 | 6 | 14 | ✓ | ✓ READY |
|
||||
| scan | 1,525 | 591 | 18 | 12 | 22 | ✓ | ✓ READY |
|
||||
| for | 78 | 51 | 0 | 0 | 1 | ✗ | ⚠ no SMEM |
|
||||
| topk | 121 | 113 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
| transform | 549 | 185 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
| batch_memcpy | 227 | 95 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
| select_if | 2,729 | 459 | 84 | 52 | 0 | ✓ | △ inline |
|
||||
| radix_sort | 2,381 | 222 | 70 | 0 | 0 | ✓ | △ inline |
|
||||
| scan_by_key | 2,008 | 145 | 30 | 17 | 0 | ✓ | △ inline |
|
||||
| reduce_by_key | 1,735 | 171 | 32 | 22 | 0 | ✓ | △ inline |
|
||||
| unique_by_key | 1,539 | 166 | 29 | 21 | 0 | ✓ | △ inline |
|
||||
| three_way_partition | 788 | 99 | 13 | 9 | 0 | ✓ | △ inline |
|
||||
| rle_non_trivial_runs | 691 | 68 | 8 | 8 | 0 | ✗ | △ inline |
|
||||
| segmented_sort | 640 | 189 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| rle_encode | 626 | 63 | 4 | 7 | 0 | ✗ | △ inline |
|
||||
| histogram | 363 | 76 | 4 | 3 | 0 | ✗ | △ inline |
|
||||
| segmented_radix_sort | 311 | 48 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| batch_memcpy | 227 | 95 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
| merge_sort | 193 | 83 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| segmented_reduce | 189 | 51 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
| batched_topk | 186 | 66 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| merge | 180 | 89 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| segmented_scan | 158 | 45 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| adjacent_difference | 118 | 77 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| find_bound_sorted_values | 106 | 47 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
| find | 90 | 39 | 0 | 0 | 0 | ✓ | △ inline |
|
||||
| transform_tile | 85 | 33 | 0 | 0 | 0 | ✗ | △ inline |
|
||||
|
||||
## gen_patch 管道状态
|
||||
|
||||
当前 gen_patch.py 跑出来的结果：
|
||||
|
||||
```
|
||||
READ reduce: bi100_plus_float32_o4 → {items:24, threads:512, vec:2}
|
||||
READ scan: bi100_sm90_float32 → {threads:128, items:24}
|
||||
READ topk: __inline_topk__ → {threads:512, bits_per_pass:11}
|
||||
READ transform: __inline_transform__ → {bytes_in_flight:64}
|
||||
READ for: bi100_default → {threads:256, items:4}
|
||||
SKIP 其余 21 个算法: no bi100_* structs
|
||||
```
|
||||
|
||||
**0 个 patch 生成**——因为 VLLM_INJECTION_POINTS 映射表中的 key 与当前 struct 字段名不匹配。这是管道断裂点。
|
||||
|
||||
## CCCL benchmark 源码作为 muh 的输入规范
|
||||
|
||||
CCCL bench/reduce/base.cuh 定义了 benchmark 框架：
|
||||
- 参数空间：`%RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1` / `%RANGE% TUNE_THREADS_PER_BLOCK tpb 128:1024:32`
|
||||
- 输出格式：`ipt_N.tpb_M.ipv_K speedup0 speedup1 speedup2 speedup3`
|
||||
- 四个 problem size：`Elements{io}` = 2^16, 2^20, 2^24, 2^28
|
||||
|
||||
muh 的 bench_bi100.py 已经有 topk 的实测数据（最佳配置：ipt=4, tpb=512, ld=0），
|
||||
但 reduce/scan/transform 还没跑。
|
||||
|
||||
## CCCL 已有的可直接利用的资产
|
||||
|
||||
| 资产类型 | 数量 | 路径 | 用途 |
|
||||
|----------|------|------|------|
|
||||
| CUB benchmarks | 80 .cu | cccl_upstream/cub/benchmarks/bench/ | 参数空间搜索框架 |
|
||||
| CUB tests | 243 .cu | cccl_upstream/cub/test/ | 正确性验证 |
|
||||
| CUB examples | 18 .cu | cccl_upstream/cub/examples/ | API 验证 |
|
||||
| Thrust examples | 52 .cu | cccl_upstream/thrust/examples/ | 算法验证 |
|
||||
| muh schemas | 27 .yaml | muh/schema/ | 参数空间定义 |
|
||||
|
||||
总计 420 个 .cu 文件可直接编译运行在 BI-V100 上产出数据。
|
||||
|
||||
## 下一步行动
|
||||
|
||||
优先级按竞赛权重排序：
|
||||
|
||||
1. **reduce 实测** (Output TPS × 16.796 = 83%): 用 bench/reduce/sum.cu 框架，在 BI-V100 上扫描 ipt∈[7,24] × tpb∈{128..1024:32} × ipv∈{1,2,4}
|
||||
2. **scan 实测** (decode softmax): 用 bench/scan/exclusive/sum.cu 框架，额外标定 LookbackDelay
|
||||
3. **topk 补全** (sampling): 已有部分数据，需要补 batch=4 和 bits_per_pass 对比
|
||||
4. **gen_patch 闭环**: 修复 VLLM_INJECTION_POINTS 映射，让 gen_patch 真正产出可用 patch
|
||||
5. **50+ 功能测试**: 在 patch 后的 vllm 上跑竞赛功能验证
|
||||
@@ -1,134 +0,0 @@
|
||||
================================================================================
|
||||
CCCL vs muh 精确比对审计报告
|
||||
================================================================================
|
||||
|
||||
### 1. scale_mem_bound 函数 parity check
|
||||
------------------------------------------------------------
|
||||
float32 (CCCL SM100 reduce) CCCL=( 16i, 512t,tile= 32768B) muh=( 16i, 512t,tile= 32768B) ✓
|
||||
float64 (CCCL SM100 reduce) CCCL=( 8i, 640t,tile= 40960B) muh=( 8i, 640t,tile= 40960B) ✓
|
||||
accum8 (CCCL SM100 reduce) CCCL=( 7i, 512t,tile= 28672B) muh=( 7i, 512t,tile= 28672B) ✓
|
||||
scan 4B (CCCL SM100 scan) CCCL=( 22i, 384t,tile= 33792B) muh=( 22i, 384t,tile= 33792B) ✓
|
||||
scan 8B (CCCL SM100 scan) CCCL=( 11i, 416t,tile= 36608B) muh=( 11i, 416t,tile= 36608B) ✓
|
||||
det float32 SM90 CCCL=( 13i, 224t,tile= 11648B) muh=( 13i, 224t,tile= 11648B) ✓
|
||||
det float64 SM86 CCCL=( 5i, 128t,tile= 5120B) muh=( 5i, 128t,tile= 5120B) ✓
|
||||
1-byte type CCCL=( 32i, 256t,tile= 8192B) muh=( 32i, 256t,tile= 8192B) ✓
|
||||
2-byte type CCCL=( 32i, 256t,tile= 16384B) muh=( 32i, 256t,tile= 16384B) ✓
|
||||
16-byte type (int128) CCCL=( 4i, 256t,tile= 16384B) muh=( 4i, 256t,tile= 16384B) ✓
|
||||
SMEM cap test (should trigger) CCCL=( 8i, 768t,tile= 49152B) muh=( 8i, 768t,tile= 49152B) ✓
|
||||
→ scale_mem_bound: FULL PARITY ✓
|
||||
|
||||
### 2. reduce tuning: CCCL SM100值 → BI-V100 scale_mem_bound适配后
|
||||
------------------------------------------------------------
|
||||
CCCL benchmarked on SM100 → muh should use scale_mem_bound for BI-V100
|
||||
Key: reduce loads to REGISTERS not SMEM → SMEM cap rarely triggers
|
||||
|
||||
float32_plus_o4 @4B: scaled=(16i, 512t) tile= 32768B (66.7%)
|
||||
float32_plus_o4 @8B: scaled=( 8i, 512t) tile= 32768B (66.7%)
|
||||
float64_plus_o4 @4B: scaled=(16i, 640t) tile= 40960B (83.3%)
|
||||
float64_plus_o4 @8B: scaled=( 8i, 640t) tile= 40960B (83.3%)
|
||||
accum8_plus_o4 @4B: scaled=(15i, 512t) tile= 30720B (62.5%)
|
||||
accum8_plus_o4 @8B: scaled=( 7i, 512t) tile= 28672B (58.3%)
|
||||
accum8_plus_o8 @4B: scaled=(15i, 512t) tile= 30720B (62.5%)
|
||||
accum8_plus_o8 @8B: scaled=( 7i, 512t) tile= 28672B (58.3%)
|
||||
det_float32_sm90 @4B: scaled=(13i, 224t) tile= 11648B (23.7%)
|
||||
det_float32_sm90 @8B: scaled=( 6i, 224t) tile= 10752B (21.9%)
|
||||
det_float32_sm86 @4B: scaled=( 6i, 224t) tile= 5376B (10.9%)
|
||||
det_float32_sm86 @8B: scaled=( 3i, 224t) tile= 5376B (10.9%)
|
||||
det_float64_sm86 @4B: scaled=(11i, 128t) tile= 5632B (11.5%)
|
||||
det_float64_sm86 @8B: scaled=( 5i, 128t) tile= 5120B (10.4%)
|
||||
default_fallback @4B: scaled=(16i, 256t) tile= 16384B (33.3%)
|
||||
default_fallback @8B: scaled=( 8i, 256t) tile= 16384B (33.3%)
|
||||
|
||||
### 3. muh bi100 reduce当前值 vs CCCL参考
|
||||
------------------------------------------------------------
|
||||
muh改用了更大的items (24 vs SM100的16)来补偿16 SMs
|
||||
这是对的——reduce加载到寄存器,SMEM不是瓶颈
|
||||
|
||||
★ float32 plus (paged_attention score reduction — 83% weight):
|
||||
CCCL SM100: items=16, threads=512, vec=2
|
||||
muh BI-V100: items=24, threads=512, vec=2
|
||||
理由: 16 SMs vs 148 SMs, 每个CTA需要处理更多数据
|
||||
tile对比: SM100=512*16*4=32768B | BI-V100=512*24*4=49152B (exactly 48KB)
|
||||
→ items=24 用满了SMEM → 合理但有风险,如果BlockReduce实际占SMEM则溢出
|
||||
→ 但注释说reduce不用BlockLoad(loads to registers) → 安全
|
||||
|
||||
### 4. scan tuning: CCCL SM100 → BI-V100 SMEM约束
|
||||
------------------------------------------------------------
|
||||
Scan DOES use BlockLoad staging in SMEM → tile_bytes ≤ 49152 is HARD
|
||||
|
||||
lookback_1B_o4 @1B: tpb= 512 ipt=18 tile= 9216B ✓
|
||||
lookback_1B_o4 @2B: tpb= 512 ipt=18 tile= 18432B ✓
|
||||
lookback_1B_o4 @4B: tpb= 512 ipt=18 tile= 36864B ✓
|
||||
lookback_1B_o4 @8B: tpb= 512 ipt=18 tile= 73728B ✗ OVERFLOW → max_items=12
|
||||
lookback_2B_o4 @1B: tpb= 512 ipt=13 tile= 6656B ✓
|
||||
lookback_2B_o4 @2B: tpb= 512 ipt=13 tile= 13312B ✓
|
||||
lookback_2B_o4 @4B: tpb= 512 ipt=13 tile= 26624B ✓
|
||||
lookback_2B_o4 @8B: tpb= 512 ipt=13 tile= 53248B ✗ OVERFLOW → max_items=12
|
||||
lookback_4B_o4 @1B: tpb= 384 ipt=22 tile= 8448B ✓
|
||||
lookback_4B_o4 @2B: tpb= 384 ipt=22 tile= 16896B ✓
|
||||
lookback_4B_o4 @4B: tpb= 384 ipt=22 tile= 33792B ✓
|
||||
lookback_4B_o4 @8B: tpb= 384 ipt=22 tile= 67584B ✗ OVERFLOW → max_items=16
|
||||
lookback_8B_o4 @1B: tpb= 416 ipt=23 tile= 9568B ✓
|
||||
lookback_8B_o4 @2B: tpb= 416 ipt=23 tile= 19136B ✓
|
||||
lookback_8B_o4 @4B: tpb= 416 ipt=23 tile= 38272B ✓
|
||||
lookback_8B_o4 @8B: tpb= 416 ipt=23 tile= 76544B ✗ OVERFLOW → max_items=14
|
||||
lookback_1B_o8 @1B: tpb= 384 ipt=14 tile= 5376B ✓
|
||||
lookback_1B_o8 @2B: tpb= 384 ipt=14 tile= 10752B ✓
|
||||
lookback_1B_o8 @4B: tpb= 384 ipt=14 tile= 21504B ✓
|
||||
lookback_1B_o8 @8B: tpb= 384 ipt=14 tile= 43008B ✓
|
||||
lookback_4B_o8 @1B: tpb= 416 ipt=19 tile= 7904B ✓
|
||||
lookback_4B_o8 @2B: tpb= 416 ipt=19 tile= 15808B ✓
|
||||
lookback_4B_o8 @4B: tpb= 416 ipt=19 tile= 31616B ✓
|
||||
lookback_4B_o8 @8B: tpb= 416 ipt=19 tile= 63232B ✗ OVERFLOW → max_items=14
|
||||
lookback_8B_o8 @1B: tpb= 320 ipt=22 tile= 7040B ✓
|
||||
lookback_8B_o8 @2B: tpb= 320 ipt=22 tile= 14080B ✓
|
||||
lookback_8B_o8 @4B: tpb= 320 ipt=22 tile= 28160B ✓
|
||||
lookback_8B_o8 @8B: tpb= 320 ipt=22 tile= 56320B ✗ OVERFLOW → max_items=19
|
||||
|
||||
关键发现:
|
||||
- scan lookback_4B_o4: items=22, threads=384 → tile@4B=33792 ✓ tile@8B=67584 ✗
|
||||
- scan lookback_8B_o4: items=23, threads=416 → tile@8B=76544 ✗
|
||||
- 这些值在SM100上是安全的(228KB SMEM),但在BI-V100(48KB)上必须降级
|
||||
- muh已经做了降级(用scale_mem_bound),但需要验证降级后的值是否正确
|
||||
|
||||
### 5. CCCL benchmark format解析
|
||||
------------------------------------------------------------
|
||||
NVIDIA的benchmark注释格式:
|
||||
ipt_<items>.tpb_<threads>.ns_<delay>.dcid_<algo>.l2w_<latency>.trp_<transpose>.ld_<load>
|
||||
后跟4个浮点数: 在[2^16, 2^20, 2^24, 2^28]四个problem size下的speedup
|
||||
|
||||
dcid映射:
|
||||
0 = no_delay
|
||||
1 = fixed_delay
|
||||
2 = exp_backoff
|
||||
3 = exp_backoff_jitter
|
||||
4 = exp_backoff_jitter_window
|
||||
5 = exp_backon_jitter_window
|
||||
6 = exp_backon_jitter
|
||||
7 = exp_backon
|
||||
|
||||
### 6. 竞赛关键路径优先级
|
||||
------------------------------------------------------------
|
||||
Token吞吐加权值 = Output_TPS × 16.796 + Input_TPS × 2.799 + Cache_TPS × 0.56
|
||||
→ Output_TPS权重83%, Input_TPS权重14%, Cache_TPS权重3%
|
||||
|
||||
decode热路径 (Output TPS):
|
||||
1. paged_attention score reduction → reduce (DONE: muh tuned)
|
||||
2. softmax denominator prefix-sum → scan (DONE: muh tuned)
|
||||
3. top-k/top-p sampling → topk/radix_sort (DONE: muh tuned)
|
||||
4. RMSNorm/SiLU/RoPE element-wise → transform (DONE: muh tuned)
|
||||
|
||||
prefill热路径 (Input TPS):
|
||||
5. flash_attention → scan + reduce
|
||||
6. MoE expert routing → select_if + reduce_by_key
|
||||
|
||||
cache热路径 (Cache TPS):
|
||||
7. KV cache block copy → batch_memcpy (DONE: muh tuned)
|
||||
|
||||
### 7. 待验证的关键问题
|
||||
------------------------------------------------------------
|
||||
1. reduce items=24: 虽然loads to registers, 但实际BlockReduce<WARP_REDUCTIONS>的SMEM用量需要确认
|
||||
2. scan delay参数: 0.5x/0.6x缩放是启发式, 需要BI-V100实测L2 write latency
|
||||
3. LOAD_LDG vs LOAD_DEFAULT: topk bench显示BI-V100上LOAD_DEFAULT更快, reduce/scan可能同理
|
||||
4. SM count=16 → wave efficiency: 所有tuning都需要重新算occupancy
|
||||
5. transform bytes_in_flight: 从18GB/s改为56GB/s后items需要相应增大
|
||||
@@ -1,104 +0,0 @@
|
||||
# CCCL → vllm Kernel Pattern Mapping
|
||||
## BI-V100 Competition Reference
|
||||
|
||||
### Pattern 1: Multi-field Reduction (paged_attention)
|
||||
|
||||
**CCCL source**: `thrust/examples/bounding_box.cu`, `summary_statistics.cu`
|
||||
**vllm kernel**: `paged_attn.py` → ixformer paged_attention_v1/v2
|
||||
|
||||
```
|
||||
CCCL: transform_reduce(begin, end, unary_op, init, binary_op)
|
||||
vllm: for each KV block: score = Q·K, max_score = reduce_max, exp_sum = reduce_sum
|
||||
```
|
||||
|
||||
**Tuning surface**:
|
||||
- `_PARTITION_SIZE`: controls how many KV tokens per CTA in V2 mode
|
||||
- V1/V2 dispatch threshold: `total_tiles vs 2 × sm_count`
|
||||
- BI-V100: 16 SMs → V2 beneficial when seq_len > 1024 (2 waves of 16 CTAs × 512 partition)
|
||||
|
||||
**CCCL parameter**: `ReducePassPolicy{threads=512, items=24, vec=2, WARP_REDUCTIONS, LDG}`
|
||||
|
||||
### Pattern 2: Prefix Scan + Transform (softmax)
|
||||
|
||||
**CCCL source**: `thrust/examples/simple_moving_average.cu`, `cub/benchmarks/bench/scan/exclusive/sum.cu`
|
||||
**vllm kernel**: `prefix_prefill.py` context_attention_fwd_kernel
|
||||
|
||||
```
|
||||
CCCL: inclusive_scan(begin, end, output, plus<float>)
|
||||
vllm: for each BLOCK_N chunk: qk = Q·K, m_new = max(m_old, max(qk)),
|
||||
l_new = l_old * exp(m_old - m_new) + sum(exp(qk - m_new))
|
||||
```
|
||||
|
||||
**Tuning surface**:
|
||||
- `BLOCK_M`: Q tile rows (32 or 64 for BI-V100)
|
||||
- `BLOCK_N`: K/V sweep width (32 or 64)
|
||||
- `NUM_WARPS`: 4 (16 SMs don't benefit from 8 warps per CTA)
|
||||
- `num_stages`: 1 (no cp.async) or 2 (software pipeline)
|
||||
|
||||
**CCCL parameter**: `ScanLookbackPolicy{threads=384, items=22, WARP_TRANSPOSE, DEFAULT, WARP_SCANS, {backon_jitter_window, 952, 415}}`
|
||||
|
||||
### Pattern 3: Transform (activation functions)
|
||||
|
||||
**CCCL source**: `cub/benchmarks/bench/transform/babelstream.cu`
|
||||
**vllm kernel**: Triton SiLU, GeLU, RMSNorm kernels (via `_custom_ops.py`)
|
||||
|
||||
```
|
||||
CCCL: transform(begin, end, output, silu_op) // x * sigmoid(x)
|
||||
vllm: @triton.jit def silu_kernel(x): tl.sigmoid(x) * x
|
||||
```
|
||||
|
||||
**Tuning surface**:
|
||||
- `bytes_in_flight`: 64KB on BI-V100 (56 GB/s per-SM × 1100ns latency)
|
||||
- Triton `num_stages=2` maps to BIF=64KB (2× prefetch window)
|
||||
- `SMEM = 49152` (fixed by _custom_ops.py)
|
||||
|
||||
**CCCL parameter**: `TransformPrefetchPolicy{threads=256, bif=64KB, prefetch_stride=128}`
|
||||
|
||||
### Pattern 4: TopK (sampling)
|
||||
|
||||
**CCCL source**: `cub/benchmarks/bench/topk/keys.cu`
|
||||
**vllm kernel**: sampling_kernels (precompiled .so)
|
||||
|
||||
```
|
||||
CCCL: DeviceTopk::TopK(keys, k, output)
|
||||
vllm: ixformer topk_sampling → radix_sort + select partial
|
||||
```
|
||||
|
||||
**Tuning surface** (via .so, limited):
|
||||
- `bits_per_pass`: 11 for float32 (32 bits / 3 passes)
|
||||
- Thread count: 512 (baked into .so)
|
||||
|
||||
### Pattern 5: Triton Flash Attention (all patterns combined)
|
||||
|
||||
**CCCL source**: All of the above + `cub/agent/agent_scan.cuh` union SMEM model
|
||||
**vllm kernel**: `triton_flash_attention.py`
|
||||
|
||||
```
|
||||
Q_resident × K_streaming × softmax_online → Output
|
||||
= transform_reduce (Q·K) + scan (softmax) + transform (V matmul)
|
||||
```
|
||||
|
||||
**Tuning surface**: 17 existing + 19 new autotune configs from gen_config.py
|
||||
**Key configs for BI-V100**:
|
||||
```python
|
||||
# Best for long context (seq_len > 4096):
|
||||
Config(BLOCK_M=64, BLOCK_N=64, num_warps=4, num_stages=2) # 40KB SMEM, 1 CTA/SM
|
||||
|
||||
# Best for short context (seq_len < 1024):
|
||||
Config(BLOCK_M=32, BLOCK_N=32, num_warps=2, num_stages=2) # 32KB SMEM, 2 CTAs/SM
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
### CCCL Asset Utilization Summary
|
||||
|
||||
| CCCL Asset | Files | Used for BI-V100 | Competition Impact |
|
||||
|-----------|-------|-------------------|-------------------|
|
||||
| Tuning headers (26) | 18094 lines | 3568 lines (20%) | P0: reduce/scan/transform |
|
||||
| CUB benchmarks (80) | reduce/scan/topk/transform | benchmark framework | P0: parameter search |
|
||||
| Thrust examples (52) | summary_stats/bounding_box/norm | pattern mapping | P1: architecture understanding |
|
||||
| CUB tests (243) | correctness verification | 0% (need BI-V100) | P2: correctness |
|
||||
| libcudacxx (1463) | type traits, atomics | implicit (via CUB) | Infra |
|
||||
|
||||
**Total usable CCCL assets**: 5205 files in cccl_upstream
|
||||
**Competition-critical subset**: ~30 files (5 tuning headers + 10 benchmarks + 15 examples)
|
||||
@@ -1,89 +0,0 @@
|
||||
# CCCL ↔ muh Tuning Header Gap Report
|
||||
|
||||
> **Generated**: 2026-08-06 (auto-analyzed from source code)
|
||||
> **Source of truth**: `cccl_upstream/cub/cub/device/dispatch/tuning/tuning_*.cuh`
|
||||
> **muh headers**: `muh/include/muh/tuning/tuning_*.cuh`
|
||||
|
||||
## Executive summary
|
||||
|
||||
- **26 algorithms** have both CCCL original and muh BI-V100 tuning headers.
|
||||
- muh covers **19% of CCCL lines** (3568 / 18094).
|
||||
- CCCL contains **294 benchmark annotations** across all algorithms. muh has **1 benchmarked algorithm** (scan, partial).
|
||||
- The **#1 gap** is not code coverage — it's the absence of BI-V100 benchmark data in `ipt_N.tpb_M speedup` format.
|
||||
|
||||
## Per-algorithm coverage
|
||||
|
||||
| Algorithm | CCCL lines | muh lines | Coverage | CCCL bench pts | muh bi100 structs | muh benchmarked? |
|
||||
|-----------|-----------|----------|----------|---------------|-------------------|-----------------|
|
||||
| reduce | 478 | 297 | 62% | 6 | 14 | ✗ |
|
||||
| scan | 1525 | 591 | 38% | 16 | 22 | ✓ (partial) |
|
||||
| topk | 121 | 113 | 93% | 0 | 0 | ✗ |
|
||||
| radix_sort | 2381 | 222 | 9% | 70 | 0 | ✗ |
|
||||
| select_if | 2729 | 459 | 16% | 82 | 0 | ✗ |
|
||||
| scan_by_key | 2008 | 145 | 7% | 30 | 0 | ✗ |
|
||||
| reduce_by_key | 1735 | 171 | 9% | 32 | 0 | ✗ |
|
||||
| unique_by_key | 1539 | 166 | 10% | 29 | 0 | ✗ |
|
||||
| three_way_partition | 788 | 99 | 12% | 13 | 0 | ✗ |
|
||||
| rle_non_trivial_runs | 691 | 68 | 9% | 8 | 0 | ✗ |
|
||||
| segmented_sort | 640 | 189 | 29% | 0 | 0 | ✗ |
|
||||
| rle_encode | 626 | 63 | 10% | 4 | 0 | ✗ |
|
||||
| transform | 549 | 185 | 33% | 0 | 0 | ✗ |
|
||||
| histogram | 363 | 76 | 20% | 4 | 0 | ✗ |
|
||||
| segmented_radix_sort | 311 | 48 | 15% | 0 | 0 | ✗ |
|
||||
| batch_memcpy | 227 | 95 | 41% | 0 | 0 | ✗ |
|
||||
| batched_topk | 186 | 66 | 35% | 0 | 0 | ✗ |
|
||||
| merge_sort | 193 | 83 | 43% | 0 | 0 | ✗ |
|
||||
| merge | 180 | 89 | 49% | 0 | 0 | ✗ |
|
||||
| segmented_reduce | 189 | 51 | 26% | 0 | 0 | ✗ |
|
||||
| segmented_scan | 158 | 45 | 28% | 0 | 0 | ✗ |
|
||||
| adjacent_difference | 118 | 77 | 65% | 0 | 0 | ✗ |
|
||||
| find | 90 | 39 | 43% | 0 | 0 | ✗ |
|
||||
| find_bound_sorted_values | 106 | 47 | 44% | 0 | 0 | ✗ |
|
||||
| transform_tile | 85 | 33 | 38% | 0 | 0 | ✗ |
|
||||
| for | 78 | 51 | 65% | 0 | 1 | ✗ |
|
||||
| **TOTAL** | **18094** | **3568** | **19%** | **294** | **37** | **1/26** |
|
||||
|
||||
## Reduce: CCCL SM100 → muh BI-V100 divergence analysis
|
||||
|
||||
### SM100 benchmark annotations in CCCL
|
||||
```
|
||||
ipt_15.tpb_512.ipv_2 1.020 1.000 1.018 1.058 (geo=1.024) — accum8, offset4
|
||||
ipt_15.tpb_512.ipv_1 1.019 1.000 1.017 1.057 (geo=1.023) — accum8, offset8
|
||||
ipt_16.tpb_512.ipv_2 1.061 1.000 1.065 1.167 (geo=1.072) — float32, offset4
|
||||
ipt_16.tpb_640.ipv_1 1.018 1.000 1.016 1.057 (geo=1.022) — float64, offset4
|
||||
ipt_13.tpb_224 1.107 1.010 1.097 1.317 (geo=1.127) — deterministic float32 (sm90)
|
||||
ipt_6.tpb_224 1.034 1.000 1.032 1.091 (geo=1.039) — deterministic float32 (sm86)
|
||||
```
|
||||
|
||||
### Key divergences
|
||||
|
||||
| Parameter | CCCL SM100 | muh BI-V100 | Rationale | Risk |
|
||||
|-----------|-----------|------------|-----------|------|
|
||||
| float32+plus items | 16 | 24 | Compensate for 16 vs 148 SMs | Unvalidated: may hurt L1 hit rate |
|
||||
| float64+plus threads | 640 | 384 | Clean 12-warp config | May underutilize vs 20-warp original |
|
||||
| float64+plus vec | 1 | 2 | 16B vectorized loads | Alignment risk with non-contiguous data |
|
||||
| det float32 items | 13 | 32 | More work per CTA on 16 SMs | 2.5× register pressure increase |
|
||||
| accum1/2/16 | absent | added | Extrapolated from scaling | Not in CCCL SM100, completely theoretical |
|
||||
|
||||
## Scan: lookback delay calibration gap
|
||||
|
||||
CCCL SM100 lookback delay parameters (from benchmark annotations):
|
||||
- `delay_ns` range: 228 – 1904 ns
|
||||
- `dcid` (delay constructor ID) range: 1 – 7
|
||||
- `l2_write_latency` range: 520 – 965 ns
|
||||
|
||||
These are calibrated on SM100's 50MB L2 cache. BI-V100 has 6MB L2 → delay parameters need re-calibration. Current muh values use heuristic scaling (SM100 × 0.5 for ns, × 0.6 for l2w) without hardware validation.
|
||||
|
||||
## Priority action items (by Output TPS impact)
|
||||
|
||||
| # | Algorithm | CCCL bench pts needed | vllm hot path | Weight |
|
||||
|---|-----------|----------------------|---------------|--------|
|
||||
| 1 | reduce | 6 | paged_attention score reduction | 83% |
|
||||
| 2 | scan | 16 (8 remaining) | softmax denominator | 83% |
|
||||
| 3 | topk | 0 (format from radix_sort) | vocab=152064 sampling | 83% |
|
||||
| 4 | radix_sort | 70 | logit sorting for top-k/top-p | 83% |
|
||||
| 5 | select_if | 82 | top-p token filtering | 83% |
|
||||
| 6 | transform | 0 (no CCCL benches) | RMSNorm/SiLU/RoPE | 10-15% |
|
||||
| 7 | scan_by_key | 30 | per-sequence softmax | ~5% |
|
||||
| 8 | reduce_by_key | 32 | per-sequence aggregation | ~3% |
|
||||
| 9 | batch_memcpy | 0 | KV cache block copy | 3% |
|
||||
193
CODEPATH_MAP.md
193
CODEPATH_MAP.md
@@ -1,193 +0,0 @@
|
||||
# 代码路径时序图 — 从HTTP请求到GPU kernel的完整链路
|
||||
|
||||
## 一、请求入口到引擎调用
|
||||
|
||||
```
|
||||
HTTP POST /v1/chat/completions
|
||||
│
|
||||
├─ api_server.py → FastAPI route handler
|
||||
│ └─ serving_chat.py:create_chat_completion() [line ~140]
|
||||
│ ├─ protocol.py:ChatCompletionRequest.model_validate()
|
||||
│ │ └─ max_completion_tokens → max_tokens 映射 [line 418]
|
||||
│ │ └─ extra="allow" (Sub168用extra="forbid"导致400)
|
||||
│ │
|
||||
│ ├─ chat_utils.py → 消息格式化 + 多模态处理
|
||||
│ │ └─ content=None容错 (Sub168这里崩)
|
||||
│ │
|
||||
│ ├─ serving_chat.py [line 175-213] → enable_thinking逻辑
|
||||
│ │ ├─ tool_choice=auto + tools存在 → enable_thinking=False
|
||||
│ │ ├─ thinking.type=disabled → enable_thinking=False
|
||||
│ │ └─ 默认 → enable_thinking=True
|
||||
│ │
|
||||
│ ├─ serving_chat.py [line 250-252] → n值检查
|
||||
│ │ └─ n>2 → 400 (n=2允许传入引擎)
|
||||
│ │
|
||||
│ └─ engine_client.generate() [line 355]
|
||||
│ └─ try/except ValueError + catch-all Exception
|
||||
│
|
||||
├─ computility-run.yaml → vLLM启动参数
|
||||
│ ├─ --max-num-seqs 2 (防止n=2崩溃)
|
||||
│ ├─ --max-model-len 80000
|
||||
│ ├─ --enforce-eager (禁用CUDA Graph)
|
||||
│ ├─ --enable-prefix-caching
|
||||
│ └─ --tool-call-parser qwen3_coder
|
||||
│
|
||||
└─ 如果引擎crash → 后续所有请求Connection Refused
|
||||
(Sub508的根因: t2_n_2触发, 30个FAIL级联)
|
||||
```
|
||||
|
||||
## 二、模型前向传播 — 逐层链路
|
||||
|
||||
```
|
||||
Qwen3_5ForCausalLM.forward() [qwen3_5.py line 1214]
|
||||
│
|
||||
└─ Qwen3_5Model.forward() [line 1094]
|
||||
│
|
||||
├─ embed_tokens(input_ids)
|
||||
│
|
||||
└─ for layer in self.layers: # 36层 (Qwen3.6-27B典型配置)
|
||||
│
|
||||
├─ GemmaRMSNorm(hidden_states, residual)
|
||||
│ └─ ☆ 可用ixformer: fused_add_rms_norm(input, residual, weight, eps)
|
||||
│
|
||||
├─ [linear_attention层] GatedDeltaNet.forward() [line 407]
|
||||
│ │
|
||||
│ ├─ CoreX dispatch尝试 [line 416-425]
|
||||
│ │ └─ _use_corex_gdn=False (base image无corex_gdn模块)
|
||||
│ │
|
||||
│ └─ _pytorch_forward() [line 435] ← 当前执行路径
|
||||
│ │
|
||||
│ ├─ 投影: in_proj_qkv, in_proj_z, in_proj_b, in_proj_a
|
||||
│ │ └─ ☆ 每个是F.linear → 可用ixformer.matmul
|
||||
│ │
|
||||
│ ├─ [prefill] 逐序列循环 [line 463-555]
|
||||
│ │ │
|
||||
│ │ ├─ F.conv1d (causal conv)
|
||||
│ │ │ └─ ☆ 可用ixformer.conv2d (需reshape)
|
||||
│ │ │
|
||||
│ │ ├─ F.silu → ☆ 可用ixformer.silu_and_mul
|
||||
│ │ │
|
||||
│ │ ├─ g计算: -A_log.exp() * softplus(a+dt_bias)
|
||||
│ │ │ └─ 当前: clamp(-8,4)后exp, softplus.clamp(max=10)
|
||||
│ │ │
|
||||
│ │ └─ _torch_chunk_gated_delta_rule() [line 152-247]
|
||||
│ │ │
|
||||
│ │ ├─ g.clamp(-5,2).cumsum(-1).clamp(-20,20) ← NaN修复点
|
||||
│ │ ├─ decay_mask = exp(g差) ← 所有exp在clamp后
|
||||
│ │ ├─ attn矩阵: k_beta @ key.T * decay_mask
|
||||
│ │ │ └─ ☆ 三角求解循环 → 无法用ixformer加速
|
||||
│ │ │ (这是纯序列依赖: attn[i] += attn[i,:i] @ attn[:i,:i])
|
||||
│ │ ├─ state更新循环: for i in chunks [line 219-232]
|
||||
│ │ │ ├─ q @ k.T * decay ← ☆ ixformer.matmul可加速
|
||||
│ │ │ ├─ q * exp(g) @ state ← ☆ ixformer.matmul可加速
|
||||
│ │ │ └─ state更新: state * exp(g) + k.T @ v_new
|
||||
│ │ │ └─ ☆ ixformer.matmul可加速
|
||||
│ │ └─ 最终: core_out → transpose → to(dtype)
|
||||
│ │
|
||||
│ ├─ [decode] 单token路径 [line 558-638]
|
||||
│ │ ├─ _torch_causal_conv1d_update
|
||||
│ │ │ └─ 逐通道点积 → ☆ ixformer.gemv可加速
|
||||
│ │ ├─ g_t = g.clamp(-20,2).exp_() ← NaN修复点
|
||||
│ │ ├─ temporal_state.mul_(g_t) ← 状态衰减
|
||||
│ │ ├─ torch.bmm(k, state) ← ☆ ixformer.matmul可加速
|
||||
│ │ └─ state.baddbmm_(k, delta) ← ☆ ixformer.matmul可加速
|
||||
│ │
|
||||
│ └─ GemmaRMSNorm + out_proj
|
||||
│ └─ ☆ ixformer.rms_norm + ixformer.matmul
|
||||
│
|
||||
├─ [full_attention层] Qwen3_5FullAttention.forward() [line 737]
|
||||
│ └─ 标准vLLM Attention → XFormers后端
|
||||
│ └─ ☆ 已使用ixformer.flash_attn_func (base image配置)
|
||||
│
|
||||
├─ GemmaRMSNorm(hidden_states, residual)
|
||||
│ └─ ☆ ixformer.fused_add_rms_norm
|
||||
│
|
||||
└─ [MLP/MoE] Qwen3_5MLP 或 Qwen3_5MoeSparseBlock
|
||||
│
|
||||
├─ [MLP] gate_up_proj → silu_and_mul → down_proj
|
||||
│ └─ ☆ 全部可用ixformer: matmul + silu_and_mul + matmul
|
||||
│
|
||||
└─ [MoE] Qwen3_5MoeSparseBlock.forward() [line 974]
|
||||
├─ gate(hidden) → router_logits
|
||||
├─ softmax → topk → renormalize (纯PyTorch, 无硬件加速)
|
||||
├─ _pure_pytorch_experts() [line 897]
|
||||
│ ├─ [decode T=1] 批量GEMM: 3次kernel launch
|
||||
│ │ └─ F.linear(x, w13_sel.reshape(-1,H)) ← ☆ ixformer.matmul
|
||||
│ │ └─ F.silu(gate) * up ← ☆ ixformer.silu_and_mul (需reshape)
|
||||
│ │ └─ torch.bmm(w2_sel, act) ← ☆ ixformer.matmul
|
||||
│ └─ [prefill] 逐expert循环 ← 性能瓶颈
|
||||
│ └─ 每个expert: F.linear × 2 + silu
|
||||
│ └─ ☆ 可用ixformer.matmul但循环开销不变
|
||||
└─ shared_expert: gate_up → silu_and_mul → down → sigmoid gate
|
||||
└─ ☆ 全部可用ixformer
|
||||
```
|
||||
|
||||
## 三、ixformer可用原语 vs 当前使用情况
|
||||
|
||||
| ixformer原语 | 签名 | 当前是否使用 | 可替换的PyTorch调用 |
|
||||
|-------------|------|------------|-------------------|
|
||||
| `matmul` | `matmul(input, other, out, transa, transb, alpha, beta)` | ❌ 未使用 | F.linear, torch.mm, torch.bmm, @ |
|
||||
| `softmax` | `softmax(input, dim)` | ❌ 未使用 | torch.softmax (MoE路由) |
|
||||
| `rms_norm` | `rms_norm(input, weight, output, eps)` | ❌ 未使用 | GemmaRMSNorm内部 |
|
||||
| `fused_add_rms_norm` | `fused_add_rms_norm(input, residual, weight, eps, scale)` | ❌ 未使用 | residual + layernorm 两步 |
|
||||
| `silu_and_mul` | `silu_and_mul(input, output)` | ❌ 未使用 | SiluAndMul层, F.silu(g)*up |
|
||||
| `conv2d` | `conv2d(input, weight, bias, stride, padding, dilation, groups)` | ❌ 未使用 | F.conv1d (causal conv) |
|
||||
| `flash_attn_func` | `flash_attn_func(q, k, v, dropout_p, softmax_scale, causal)` | ✅ XFormers后端使用 | full_attention层 |
|
||||
| `gemv` | `gemv(x, A)` | ❌ 未使用 | decode路径小矩阵乘 |
|
||||
| `scaled_dot_product_attention` | `sdpa(query, key, value, attn_mask, dropout_p, is_causal)` | ❌ 未使用 | 可替代chunk内QK^T计算 |
|
||||
|
||||
**关键发现：9个可用原语中只有1个（flash_attn_func）被使用，而且不是我们的代码使用的——是base image的XFormers后端自动调用的。我们的代码对ixformer的利用率是0%。**
|
||||
|
||||
## 四、Sub168 vs Sub508 性能差距的代码解释
|
||||
|
||||
```
|
||||
Sub168 (8.49s for d01):
|
||||
base image native qwen3_5.py
|
||||
├─ corex_gdn: 使用libcorex_gdn.so的fused GDN kernel ← 不存在于我们的base image
|
||||
├─ corex_moe: 使用libcorex_moe.so的fused MoE kernel ← 不存在于我们的base image
|
||||
└─ 所有底层ops由ixformer后端加速 (matmul/rms_norm/softmax等)
|
||||
|
||||
Sub508 (95.85s for d01):
|
||||
我们的自定义 qwen3_5.py
|
||||
├─ GatedDeltaNet: 纯PyTorch (cumsum→exp→NaN→nan_to_num→全零)
|
||||
├─ MoE: 纯PyTorch循环 (每expert单独F.linear)
|
||||
└─ 底层ops全部用PyTorch默认kernel (未调用ixformer)
|
||||
```
|
||||
|
||||
## 五、优化路径 — 用ixformer原语替换PyTorch
|
||||
|
||||
### 立即可做 (不改算法, 只换kernel):
|
||||
1. **matmul**: 所有F.linear/torch.bmm/@ → ixformer.matmul
|
||||
2. **silu_and_mul**: MLP和MoE的silu*gate → ixformer.silu_and_mul
|
||||
3. **rms_norm**: GemmaRMSNorm内部 → ixformer.rms_norm
|
||||
4. **fused_add_rms_norm**: residual+norm两步 → 一步fused
|
||||
5. **softmax**: MoE路由softmax → ixformer.softmax
|
||||
|
||||
## 六、功能测试FAIL根因分析（6个非crash FAIL）
|
||||
|
||||
```
|
||||
FAIL类型A: NaN导致模型输出质量问题 (修NaN后自愈)
|
||||
├─ d03_tool_call: tools=0 — 模型不能输出<tool_call> XML
|
||||
├─ d07_reasoning_plus_content: content[0] — 模型不输出</think>
|
||||
├─ d10_thinking_disable_ctk: 乱码 — 模型logits被NaN扭曲
|
||||
├─ t1a_thinking_true: reasoning[0] — output.text为空→parser返回空
|
||||
└─ t1c_thinking_default: reasoning[0] — 同上
|
||||
|
||||
FAIL类型B: 请求处理层问题
|
||||
└─ d05_multimodal: HTTP 400 — 多模态请求验证失败
|
||||
|
||||
FAIL类型C: 引擎crash级联 (修max-num-seqs=2后自愈)
|
||||
└─ t2_n_2 → t3/t4/t5/t6/t7/t8/t9/t10/t12/t13/t14/t15/t16 全部HTTP 500 (25个)
|
||||
|
||||
当前代码状态:
|
||||
NaN修复: ✅ cumsum前clamp[-5,2] + 后clamp[-20,20] + A_log clamp[-8,4]
|
||||
引擎防崩: ✅ max-num-seqs=2 + catch-all Exception
|
||||
ixformer加速: ✅ matmul/bmm/softmax接入12处热路径
|
||||
reasoning parser: ✅ qwen3已注册，部署正确
|
||||
tool parser: ✅ qwen3_coder已注册，adjust_request禁thinking
|
||||
|
||||
预期: NaN修复后模型质量恢复 → 类型A的5个FAIL自愈
|
||||
max-num-seqs=2 → 类型C的25个FAIL自愈
|
||||
剩余: d05_multimodal需要单独debug
|
||||
预估: 45/51 PASS (88%)
|
||||
```
|
||||
@@ -1,165 +0,0 @@
|
||||
# comp 168 Docker 诊断 → .so 开发清单
|
||||
|
||||
> 基于 `2d5232c5d6bc` (comp 168 docker log, 3786 行)
|
||||
> 当前 HEAD: `b25fc53e` (414 commits)
|
||||
|
||||
## 一、comp 168 日志三大致命问题
|
||||
|
||||
| # | 错误 | 出现次数 | 根因 | 状态 |
|
||||
|---|------|----------|------|------|
|
||||
| 1 | `GDN NaN frac=0.9998` | 16次(layer 0-4) | 我们的 GDN prefill 实现产生 NaN → replace with zeros → 模型质量归零 | **P0 未修** |
|
||||
| 2 | `vllm_moe_topk_softmax not found` | 39次 | `ixformer.functions` 没有 Python binding → fallback to Python for 循环 | **P0 需 .so** |
|
||||
| 3 | `CUDA OOM 32 MiB` | 17次 | `max_model_len=100000` 超过 KV cache 容量 → engine 死亡 | ✅ 已修为 80000 |
|
||||
|
||||
## 二、真机探测确认的事实
|
||||
|
||||
从你贴的真机 probe 输出：
|
||||
|
||||
```
|
||||
ixformer.functions 有:
|
||||
✓ silu_and_mul, rms_norm, fused_add_rms_norm, rotary_embedding
|
||||
✓ flash_attn_*, vllm_single_query_cached_kv_attention_v2
|
||||
✓ vllm_cache_ops_reshape_and_cache, vllm_swap_blocks, vllm_copy_cache
|
||||
✗ vllm_moe_topk_softmax (不存在!)
|
||||
✗ moe_compute_token_index_api (不存在!)
|
||||
✗ moe_w16a16_group_gemm (不存在!)
|
||||
|
||||
libixformer.so 中:
|
||||
✓ 上述函数全部存在 (C++ 符号, xllm 的 ixformer.h 声明了它们)
|
||||
但 Python binding (_C.so) 没有暴露
|
||||
```
|
||||
|
||||
**结论**: MoE 7 步 pipeline 中的 topk_softmax / gen_idx / expand / group_gemm / combine 全部需要通过 `ix_moe_bridge.so` 桥接。
|
||||
|
||||
## 三、需要开发/修复的 .so 清单
|
||||
|
||||
### SO-1: `ix_moe_bridge.so` (MoE 7步 pipeline) — ✅ 代码已有，需真机编译验证
|
||||
|
||||
**源码**: `ex_engine/csrc/ix_moe_bridge.cpp` (258行)
|
||||
**编译**: `ex_engine/precompile_ix_bridge.py` → `torch.utils.cpp_extension.load(-lixformer)`
|
||||
**状态**: 代码写好了，Dockerfile 有 build step，但从未在真机验证过编译成功
|
||||
|
||||
真机验证命令:
|
||||
```bash
|
||||
cd /workspace/ex_engine
|
||||
python3 precompile_ix_bridge.py
|
||||
ls -la build/ix_moe_bridge*.so
|
||||
python3 -c "import torch; from torch.utils.cpp_extension import load; m=load('test', sources=['csrc/ix_moe_bridge.cpp'], extra_ldflags=['-L/usr/local/corex/lib64/python3/dist-packages/ixformer', '-lixformer']); print(dir(m))"
|
||||
```
|
||||
|
||||
### SO-2: GDN prefill 修复 — **P0 最高优先级**
|
||||
|
||||
**现状**: 我们的 `_torch_chunk_gated_delta_rule` 在 fp16 下产生 99.98% NaN
|
||||
**参考**: `upstream_ref/xllm/core/layers/npu_torch/qwen3_gated_delta_net_base.cpp` (576行)
|
||||
|
||||
关键差异:
|
||||
- xllm 用 `fp32` accumulation: `decay_mask = ... .exp().float()`
|
||||
- xllm 用 `torch::matmul` 而不是自定义 chunk kernel
|
||||
- xllm 的 recurrent state 管理有精确的 `clamp(-20, 20)` 限制
|
||||
|
||||
**解决方案**: 不写新 .so，而是从 xllm 搬运 GDN 的 PyTorch 实现（C++ torch ops, 全 fp32 accumulation），替换我们的 chunk kernel。
|
||||
|
||||
### SO-3: `_custom_ops.py` patch — ✅ 已有 fallback 逻辑
|
||||
|
||||
base image 的 `_custom_ops.py` 调用 `ixf_F.vllm_moe_topk_softmax` 时会报错。
|
||||
但 comp 168 的 base 镜像绕过了 `_custom_ops`，直接走 `corex_moe.py` 的 7 步 pipeline。
|
||||
|
||||
**如果 base 有 corex_moe.py**: 不需要 patch
|
||||
**如果 base 没有 corex_moe.py**: 我们的版本 + ix_moe_bridge.so 补位
|
||||
|
||||
## 四、upstream 已有、不需要重写的代码
|
||||
|
||||
| upstream 文件 | 行数 | 我们的对应文件 | 搬运状态 |
|
||||
|--------------|------|---------------|---------|
|
||||
| `xllm/core/kernels/ilu/ixformer.h` | 147 | `ex_engine/csrc/ilu/ixformer.h` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/fused_moe.cpp` | 99 | `ex_engine/csrc/ilu_kernel_fused_moe.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/layers/ilu/fused_moe.cpp` | 797 | `ex_engine/csrc/ilu_layer_fused_moe.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/attention.cpp` | 162 | `ex_engine/csrc/ilu_kernel_attention.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/layers/ilu/attention.cpp` | 189 | `ex_engine/csrc/ilu_layer_attention.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/norm.cpp` | 50 | `ex_engine/csrc/ilu_kernel_norm.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/activation.cpp` | 32 | `ex_engine/csrc/ilu_kernel_activation.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/rope.cpp` | 31 | `ex_engine/csrc/ilu_kernel_rope.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/group_gemm.cpp` | 39 | `ex_engine/csrc/ilu_kernel_group_gemm.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/ilu/matmul.cpp` | 73 | `ex_engine/csrc/ilu_kernel_matmul.cpp` | ✅ 已搬 |
|
||||
| `xllm/core/layers/npu_torch/qwen3_gated_delta_net_base.cpp` | 576 | `ex_engine/csrc/qwen3_gated_delta_net_base.cpp` | ✅ 已搬 |
|
||||
| `ds_vllm/csrc/moe/topk_softmax_kernels.cu` | 874 | `ex_engine/csrc/moe_v055/topk_softmax_kernels.cu` | ✅ 已搬 |
|
||||
| `xllm/core/kernels/cuda/moe/moe_topk_softmax_kernels.cuh` | ~400 | `ex_engine/csrc/moe/moe_topk_softmax_kernels.cuh` | ✅ 已搬 |
|
||||
|
||||
## 五、真机验证 checklist
|
||||
|
||||
在真机上按顺序执行:
|
||||
|
||||
```bash
|
||||
# 1. 验证 ix_moe_bridge.so 编译
|
||||
cd /workspace/ex_engine && python3 precompile_ix_bridge.py
|
||||
ls build/ix_moe_bridge*.so # 必须存在
|
||||
|
||||
# 2. 验证符号解析
|
||||
python3 -c "
|
||||
import torch
|
||||
import importlib.util
|
||||
spec = importlib.util.spec_from_file_location('ix', 'build/ix_moe_bridge.cpython-310-x86_64-linux-gnu.so')
|
||||
m = importlib.util.module_from_spec(spec)
|
||||
spec.loader.exec_module(m)
|
||||
print([x for x in dir(m) if not x.startswith('_')])
|
||||
# 应输出: ['topk_softmax', 'moe_gen_idx', 'moe_expand_input', 'moe_group_gemm',
|
||||
# 'silu_and_mul', 'moe_combine_result', 'paged_attention', 'rms_norm',
|
||||
# 'fused_add_rms_norm', 'linear', 'reshape_and_cache', 'rotary_embedding']
|
||||
"
|
||||
|
||||
# 3. 验证 topk_softmax 功能
|
||||
python3 -c "
|
||||
import torch
|
||||
# ... load ix_moe_bridge ...
|
||||
gating = torch.randn(4, 64, device='cuda', dtype=torch.float32)
|
||||
tw = torch.empty(4, 8, device='cuda', dtype=torch.float32)
|
||||
ti = torch.empty(4, 8, device='cuda', dtype=torch.int32)
|
||||
tei = torch.empty(4, 8, device='cuda', dtype=torch.int32)
|
||||
m.topk_softmax(tw, ti, tei, gating)
|
||||
print('topk_weights:', tw)
|
||||
print('topk_ids:', ti)
|
||||
"
|
||||
|
||||
# 4. 验证 GDN 不再 NaN
|
||||
# (需要先修复 GDN prefill 代码)
|
||||
|
||||
# 5. 启动服务验证
|
||||
python3 -m vllm.entrypoints.openai.api_server --model /model ...
|
||||
```
|
||||
|
||||
## 六、最关键发现：07-23 的 base image 自带完整 corex_* chain
|
||||
|
||||
**07-23 日志证据** (dockerrizhi.txt):
|
||||
```
|
||||
corex_gdn.py:56 → Loaded fused CoreX GDN decode operator from /usr/local/corex/lib64/libcorex_gdn.so ✅
|
||||
corex_gdn.py:228 → Using fused CoreX GDN prefill operator ✅
|
||||
corex_moe.py:339 → Using CoreX fused MoE prefill operator: tokens=4096, kernel=expert-grouped-wmma ✅
|
||||
corex_fa2.py:333 → Using CoreX FA2 packed prefill: B=2 Hq=4 Hkv=1 D=256 ✅
|
||||
corex_fa2.py:507 → Using CoreX paged FA2 chunked prefill ✅
|
||||
```
|
||||
|
||||
**08-07 日志**: 零条 corex_* 加载记录。取而代之的是 `qwen3_5.py:445 NaN in prefill` + `_custom_ops.py:58 topk_softmax not found`。
|
||||
|
||||
**根因**: 08-07 提交部署了我们自己的 `qwen3_5.py`，覆盖了 base image 自带的版本，打断了 `corex_gdn.py` / `corex_moe.py` / `corex_fa2.py` 的调用链。
|
||||
|
||||
**当前状态**: `patch_ops.sh v2` 已经有条件跳过逻辑（`_QW_SIZE > 1000 → KEEPING IT`），但需要确保下次提交时不再触发 qwen3_5.py 覆盖。
|
||||
|
||||
**结论**: 如果 base image 有工作的 corex_* chain，我们只需要:
|
||||
1. 不覆盖 qwen3_5.py
|
||||
2. 只部署 serving 层（protocol/serving_chat/api_server/tool_parser）
|
||||
3. `max_model_len=80000`（已修）
|
||||
4. `ix_moe_bridge.so` 作为备用（如果 base 的 _custom_ops 有路径碰到 topk_softmax）
|
||||
|
||||
## 七、代码量评估
|
||||
|
||||
| 组件 | 文件数 | 总行数 | 状态 |
|
||||
|------|--------|--------|------|
|
||||
| ex_engine/csrc (C++) | 39 | ~8000 | 全部已有，需真机编译 |
|
||||
| ex_engine/python (Python) | 7 | ~1200 | 全部已有，dispatch chain 完整 |
|
||||
| qwen3_6_scripts (serving) | 20+ | ~6000 | 全部已有，patch_ops.sh 管部署 |
|
||||
| upstream_ref (xllm reference) | 500+ | ~100K | 参考用，关键文件已搬到 ex_engine |
|
||||
|
||||
**结论**: 代码量是够的。问题不是代码不够，而是:
|
||||
1. GDN NaN 没修（需要用 xllm 的 fp32 accumulation 逻辑替换）
|
||||
2. ix_moe_bridge.so 从未在真机编译成功
|
||||
3. 没有 "不允许 fallback" 的硬要求落实到代码里
|
||||
@@ -1,121 +0,0 @@
|
||||
# 竞赛对比分析 & 修复计划
|
||||
|
||||
## 一、核心数据对比
|
||||
|
||||
| 模块 | 对手 Sub168 | 我们 Sub508 | 差距 |
|
||||
|------|-----------|-----------|------|
|
||||
| **functional** | 48/52 PASS (92.3%) | 21/51 PASS (41.2%) | **-51%** |
|
||||
| **case_truncation** | score=1.0 (8192 tokens输出完整) | score=0.0 (引擎崩溃) | **致命** |
|
||||
| **replay_tencent** | score=60194 (94/881成功,tps avg 11.86) | score=0.0 (881/881 connection refused) | **致命** |
|
||||
| **opencompass** | 0.0 (server也崩了) | 0.0 (同上) | 平 |
|
||||
| **总分** | **60194.6** | **0.0** | -- |
|
||||
|
||||
## 二、Sub508 崩溃根因链
|
||||
|
||||
```
|
||||
t2_n_2 (n=2请求) → get_scheduler_config() 异常 → 引擎进程死亡
|
||||
→ 后续所有请求 Connection Refused → 30个FAIL级联
|
||||
→ case_truncation/replay/opencompass 全部0分
|
||||
```
|
||||
|
||||
**关键事实：t2_n_2 崩溃发生在 06:42:45，之后所有模块都是在引擎已死的情况下跑的。**
|
||||
|
||||
## 三、对手 Sub168 的弱点（我们已经修复的）
|
||||
|
||||
1. **`max_completion_tokens` 被拒** — 对手 `extra="forbid"` 导致 replay 中所有带此字段的请求返回 400。我们已添加该字段到 protocol.py，replay 中不会被拒。
|
||||
2. **`tool_calls` content=None 被拒** — 对手的 replay preflight 失败（"Each message must have at least one of 'content' or 'reasoning_content'"）。我们已修复 chat_utils.py 中 content=None 的处理。
|
||||
3. **d06_cache_hit FAIL** — 对手没有 prefix caching，我们 PASS。
|
||||
4. **t3_max_tokens_1/64/max 3个FAIL** — 对手也有3个max_tokens测试失败。
|
||||
|
||||
**对手 replay 中 787/881 失败(89.3%)，只有 94 个成功。我们的目标是超越这个。**
|
||||
|
||||
## 四、我们需要修复的问题（按优先级排序）
|
||||
|
||||
### P0 — 引擎稳定性（决定能否拿分的前提）
|
||||
|
||||
| 问题 | 根因 | 修复位置 |
|
||||
|------|------|----------|
|
||||
| **t2_n_2 → 引擎崩溃级联** | `get_scheduler_config()` 异常 + n>1 未处理 | `qwen3_6_scripts/serving_chat.py` + `protocol.py` |
|
||||
| **引擎OOM死亡** | 单个长请求耗尽GPU内存后整个进程死 | 需要在 worker/model_runner.py 加 OOM catch |
|
||||
|
||||
已有 commit 修复（994c657 clamp n>1, c241764 try-catch scheduler），但 **Sub508 用的是修复前的代码**。Sub509 日志确认 d01 能跑（95.85s），但 d03 仍然 FAIL。
|
||||
|
||||
### P1 — d03_tool_call FAIL（功能测试核心分）
|
||||
|
||||
**Sub508**: `tools=0 finish=stop reasoning[0]` (49.04s)
|
||||
**Sub509**: `tools=0 finish=stop reasoning[0]` (49.04s)
|
||||
**对手**: `tool=get_weather args="{'city': 'Beijing'}" finish=tool_calls` (2.12s)
|
||||
|
||||
**根因分析**：
|
||||
- 对手 d03 只用了 2.12s，模型直接输出 tool_call XML，tool parser 正确解析
|
||||
- 我们用了 49.04s，模型在 thinking 中耗尽了时间，没有产生 `<tool_call>` 标签
|
||||
- commit e0344b1 说"禁用 tool_call 请求的 thinking"，但 Sub509 的 d03 仍显示 `reasoning[0]`
|
||||
- **真正的问题**：当 `tool_choice=auto` 且有 tools 时，需要在 chat_template 中设置 `enable_thinking=False`，否则 Qwen3 会先 think 再输出，大量token浪费在思考上
|
||||
|
||||
**修复方案**：在 `serving_chat.py` 的 `create_chat_completion` 中，当检测到 `request.tools` 且 `tool_choice != "none"` 时，在 `chat_template_kwargs` 中注入 `enable_thinking=False`。
|
||||
|
||||
### P1 — d05_multimodal HTTP 400
|
||||
|
||||
对手 PASS (content[374])，我们 HTTP 400。
|
||||
可能是多模态请求格式/图片解码问题。需要检查 chat_utils.py 的图片处理路径。
|
||||
|
||||
### P1 — d07_reasoning_plus_content
|
||||
|
||||
对手 PASS (reasoning[3489] content[962])，我们 FAIL (reasoning[131] content[0])。
|
||||
模型 think 后不产生 content。这是模型行为问题，但可以通过调低 thinking budget 或调整 temperature 来缓解。
|
||||
|
||||
### P2 — t1a_thinking_true / t1c_thinking_default
|
||||
|
||||
对手 PASS (reasoning[541] / [411])，我们 FAIL (reasoning[0])。
|
||||
**根因**：模型在短回答场景下不触发 thinking。可能需要在 chat_template 中确保 `enable_thinking=True` 是默认值。检查 Qwen3.6 的 chat_template 是否正确注入了 `<think>` 标签。
|
||||
|
||||
### P2 — d10_thinking_disable_ctk 乱码输出
|
||||
|
||||
对手输出 `'4'`（正确），我们输出乱码 `"presت< **sama一..."`。
|
||||
模型在 thinking disabled 模式下输出质量极差。这是模型+chat_template 的交互问题。
|
||||
|
||||
### P3 — 速度差距
|
||||
|
||||
| 测试 | 对手 | 我们 | 倍数 |
|
||||
|------|------|------|------|
|
||||
| d01 | 8.49s | 95.85s | **11x慢** |
|
||||
| d04 | 17.78s | 128.74s | **7x慢** |
|
||||
| d03 | 2.12s | 49.04s | **23x慢** |
|
||||
|
||||
速度问题核心：BI-V100 硬件本身比 NVIDIA GPU 慢，但 10x 的差距说明还有架构问题。对手的 output_tps 平均 11.86，decode 阶段 tps 在 2.4-22.7 之间。
|
||||
|
||||
## 五、修复代码的具体文件
|
||||
|
||||
需要修改的文件（全部在 `qwen3_6_scripts/` 中，会被 patch_ops.sh 部署）：
|
||||
|
||||
1. **`serving_chat.py`** — tool_call 时注入 `enable_thinking=False`
|
||||
2. **`protocol.py`** — 确认 `extra="forbid"` 已经去掉（已做），确认 `thinking` 字段被正确传递
|
||||
3. **`chat_utils.py`** — 多模态请求处理、content=None 容错
|
||||
4. **`model_runner.py`** — OOM recovery
|
||||
5. **`qwen3_5.py`** — 检查模型是否正确处理 `enable_thinking` 参数
|
||||
6. **`computility-run.yaml`** — 考虑调整 `--max-num-seqs` / `--gpu-memory-utilization`
|
||||
|
||||
## 六、对手的 replay 得分结构
|
||||
|
||||
对手 881 个请求中：
|
||||
- 94 个成功 (10.7%)
|
||||
- 77 个因 `max_completion_tokens` extra_forbidden 而 400
|
||||
- 704 个 connection refused（server也崩了！）
|
||||
- output_tps_avg = 11.86, output_tps_p50 = 12.97
|
||||
|
||||
**关键发现：对手的 server 也在 replay 后期崩溃了（704 个 connection refused）。但他在崩溃前完成了 94 个请求。**
|
||||
|
||||
我们的优势：
|
||||
- 我们已修复 `max_completion_tokens` → 对手的 77 个 400 我们不会有
|
||||
- 我们已修复 `tool_calls content=None` → 对手的 tool preflight fail 我们不会有
|
||||
- 我们有 prefix caching → 对手没有
|
||||
|
||||
**如果我们能保持引擎稳定不崩溃，仅靠不拒绝 max_completion_tokens 的请求，就能多处理 77+ 个请求，超过对手。**
|
||||
|
||||
## 七、下一步行动
|
||||
|
||||
1. 修复 `serving_chat.py`：tool_call 时禁用 thinking
|
||||
2. 确认 n>1 clamp 和 scheduler try-catch 在 patch 文件中生效
|
||||
3. 测试 OOM 恢复逻辑
|
||||
4. 调整 computility-run.yaml 参数确保稳定性
|
||||
5. 提交部署，跑测试
|
||||
@@ -1,94 +0,0 @@
|
||||
# EngineX vllm Injection Point Map
|
||||
|
||||
> **Source**: `enginex-vllm-bi100-qwen36-main.zip` (101MB, 1444 files)
|
||||
> **Generated**: 2026-08-02 from full source analysis
|
||||
|
||||
---
|
||||
|
||||
## 关键发现
|
||||
|
||||
### 1. 不是 C++ CUDA 文件注入 — 是 Python 层
|
||||
|
||||
EngineX vllm 的 CUDA kernels 全部预编译在 `ixformer.functions` (ixf_F) 中，打包在基础镜像里。
|
||||
`_custom_ops.py` 是 Python 薄封装层，调用 `ixf_F.vllm_single_query_cached_kv_attention()` 等。
|
||||
|
||||
**没有 .cu 文件可以直接 patch。** muh 的 gen_patch.py 需要改为 patch Python 文件，不是 C++ 文件。
|
||||
|
||||
### 2. paged_attention_v2 未实现
|
||||
|
||||
```python
|
||||
def paged_attention_v2(...) -> None:
|
||||
raise NotImplementedError()
|
||||
```
|
||||
|
||||
且 `use_v1 = True` 硬编码覆盖了启发式逻辑。所有 decode 都走 v1。
|
||||
|
||||
### 3. 实际可调参数 (THE TUNING SURFACE)
|
||||
|
||||
| 参数 | 文件 | 当前值 | 作用 | 优先级 |
|
||||
|------|------|--------|------|--------|
|
||||
| `_PARTITION_SIZE` | `vllm/attention/ops/paged_attn.py:13` | 512 | PagedAttention partition (v2 用) | 低 (v2 disabled) |
|
||||
| `use_v1` | `paged_attn.py:128` | `True` (hardcoded) | 强制 v1 | **P0** — 解锁 v2 可能提升长序列 |
|
||||
| `BLOCK` | `prefix_prefill.py:712` | 128 (cc≥80) / 64 | Triton prefill tile size | **P0** — 直接影响 Input TPS |
|
||||
| `NUM_WARPS` | `prefix_prefill.py:713` | 8 | Triton warp count | **P0** |
|
||||
| `BLOCK_SIZE_M/N/K` | `fused_moe.py:342-344` | 64/64/32 | MoE kernel tile | **P0** — Qwen3.6 是 MoE |
|
||||
| `get_max_shared_memory` | `_custom_ops.py:892` | `32 * 1024` | SMEM 上限声明 | **P0** — 可能错误限制性能 |
|
||||
| Triton flash attention configs | `triton_flash_attention.py:214-303` | 8 个 triton.Config | Triton autotune 搜索空间 | P1 |
|
||||
|
||||
### 4. SMEM 32KB vs 48KB 冲突
|
||||
|
||||
`_custom_ops.py:892` 返回 `32 * 1024` (32KB)。
|
||||
但 `hardware.cuh` 和 muh 假设 49152 (48KB)。
|
||||
如果 BI-V100 实际 SMEM 是 32KB，则 muh 所有 tuning 的 SMEM 约束都需要从 48KB 降到 32KB。
|
||||
|
||||
### 5. ixf_F kernel 列表 (不可改，只能调参)
|
||||
|
||||
| Python 封装 | ixf_F 调用 | 说明 |
|
||||
|-------------|-----------|------|
|
||||
| `paged_attention_v1` | `ixf_F.vllm_single_query_cached_kv_attention` | decode 核心 |
|
||||
| `silu_and_mul` | `ixf_F.silu_and_mul` | SwiGLU 激活 |
|
||||
| `rms_norm` | `ixf_F.rms_norm` | LayerNorm |
|
||||
| `fused_add_rms_norm` | `ixf_F.fused_add_rms_norm` | 融合残差+norm |
|
||||
| `rotary_embedding` | `ixf_F.vllm_rotary_embedding_neox` | RoPE 位置编码 |
|
||||
| `reshape_and_cache` | `ixf_F.vllm_cache_ops_reshape_and_cache` | KV cache 写入 |
|
||||
| `copy_blocks` | `ixf_F.copy_blocks` | prefix cache block 复制 |
|
||||
| `moe_align_block_size` | `ixf_F.vllm_moe_align_block_size` | MoE token 排列 |
|
||||
| `invoke_fused_moe_kernel` | `ixf_F.vllm_invoke_fused_moe_kernel` | MoE GEMM |
|
||||
| `topk_softmax` | `ixf_F.vllm_moe_topk_softmax` | MoE routing |
|
||||
| `cutlass_scaled_mm` | `ixf_F.w8a8` | INT8 矩阵乘 |
|
||||
|
||||
### 6. Triton kernels (可直接修改)
|
||||
|
||||
这些是 Python Triton JIT 编译的 kernel，可以直接改源码：
|
||||
|
||||
- `prefix_prefill.py` — 3 个 `_fwd_kernel` 变体 (context attention)
|
||||
- `triton_flash_attention.py` — Triton flash attention (8 个 autotune configs)
|
||||
- `fused_moe.py` — MoE GEMM kernel (Triton, 自定义 config)
|
||||
|
||||
---
|
||||
|
||||
## muh 策略修正
|
||||
|
||||
### 旧策略 (假设 C++ injection)
|
||||
```
|
||||
CCCL tuning_*.cuh → muh bi100_* → gen_patch.py → C++ #define 注入 → 编译 .so
|
||||
```
|
||||
|
||||
### 新策略 (实际 Python injection)
|
||||
```
|
||||
层1: Python 参数调优
|
||||
paged_attn.py: _PARTITION_SIZE, use_v1
|
||||
prefix_prefill.py: BLOCK, NUM_WARPS
|
||||
fused_moe.py: BLOCK_SIZE_M/N/K
|
||||
_custom_ops.py: get_max_shared_memory (32KB→实测值)
|
||||
|
||||
层2: Triton kernel 优化
|
||||
prefix_prefill.py: 3 个 _fwd_kernel — tile size, loop structure
|
||||
triton_flash_attention.py: autotune config 添加 BI-V100 特化
|
||||
fused_moe.py: MoE GEMM kernel tune
|
||||
|
||||
层3: CCCL/muh 知识迁移
|
||||
用 CCCL 的 tuning 方法论指导 Triton kernel 参数选择
|
||||
不是直接注入 C++ 值，而是把 CCCL 的 policy_selector 逻辑
|
||||
翻译成 Triton constexpr 参数
|
||||
```
|
||||
@@ -1,150 +0,0 @@
|
||||
# Engine Code Path Timeline: Sub168 vs Our Sub508/509
|
||||
|
||||
**Purpose**: Anyone reading this repo can understand the exact runtime difference in 2 minutes instead of re-deriving from raw logs.
|
||||
|
||||
## 1. Boot Sequence Comparison
|
||||
|
||||
```
|
||||
TIME SUB168 (07-23, score=60194) OUR SUB508 (08-07, score=0)
|
||||
──────────────────────────────────────────────────────────────────────────────────
|
||||
+0s api_server.py:530 → vLLM 0.6.3 api_server.py:530 → vLLM 0.6.3
|
||||
max_model_len=256000 max_model_len=256000 (same)
|
||||
max_num_seqs=2, gpu_mem=0.95 max_num_seqs=2, gpu_mem=0.95 (same)
|
||||
chunked_prefill=True chunked_prefill=True (same)
|
||||
|
||||
+10s model_runner.py:1074 load start model_runner.py:1119 load start
|
||||
↑ DIFFERENT line number ↑ DIFFERENT line number
|
||||
↑ (base image native model_runner) ↑ (our patched model_runner)
|
||||
|
||||
+18s weights = 17.3529 GB weights = 16.2303 GB
|
||||
↑ 1.1GB MORE (corex state buffers) ↑ 1.1GB LESS (no corex buffers)
|
||||
|
||||
+180s corex_gdn.py:56 → load libcorex_gdn.so qwen3_5.py:445 → NaN in prefill layer 0
|
||||
corex_gdn.py:228 → GDN prefill OK ↑ PyTorch GDN produces NaN (99.98%)
|
||||
corex_moe.py:339 → MoE prefill OK qwen3_5.py:913 → FusedMoE FAILED
|
||||
corex_fa2.py:333 → FA2 prefill OK ↑ ixformer.functions missing topk_softmax
|
||||
↑ ALL THREE CoreX accelerators loaded ↑ ZERO accelerators, all fallback
|
||||
|
||||
+182s GPU blocks: 19259 GPU blocks: ~19000 (similar)
|
||||
Ready to serve Ready to serve (but 10x slower)
|
||||
```
|
||||
|
||||
## 2. Call Chain During Inference
|
||||
|
||||
### Sub168 (with CoreX) — d01_basic_nostream: 8.49s
|
||||
```
|
||||
serving_chat.py → create_chat_completion()
|
||||
→ engine.generate()
|
||||
→ model_runner.py:1074 execute_model()
|
||||
→ qwen3_5.py:1421 Qwen3_5ForCausalLM.forward()
|
||||
→ qwen3_5.py:1165 Qwen3_5Model.forward() (decoder layers loop)
|
||||
→ qwen3_5.py:1086 Qwen3_5DecoderLayer.forward()
|
||||
├─ GatedDeltaNet layers (4 of 36):
|
||||
│ ├─ PREFILL: corex_gdn.py:228 → libcorex_gdn.so (fused CUDA kernel)
|
||||
│ └─ DECODE: corex_gdn.py:138 → libcorex_gdn.so (fused CUDA kernel)
|
||||
├─ MoE layers (all 36):
|
||||
│ ├─ PREFILL: corex_moe.py:339 → libcorex_moe.so (expert-grouped-wmma)
|
||||
│ └─ DECODE: corex_moe.py:249 → libcorex_moe.so (fused MoE decode)
|
||||
└─ Attention (32 of 36 layers):
|
||||
├─ PREFILL: corex_fa2.py:333 → libcorex_fa2.so (packed FA2)
|
||||
└─ DECODE: corex_fa2.py:225 → libcorex_fa2.so (paged decode)
|
||||
```
|
||||
|
||||
### Our Sub508 (no CoreX) — d01_basic_nostream: 95.87s (11.3x slower)
|
||||
```
|
||||
serving_chat.py → create_chat_completion()
|
||||
→ engine.generate()
|
||||
→ model_runner.py:1119 execute_model()
|
||||
→ qwen3_5.py:1369 Qwen3_5ForCausalLM.forward() (52 lines shorter!)
|
||||
→ qwen3_5.py:???? Qwen3_5Model.forward()
|
||||
→ qwen3_5.py:???? Qwen3_5DecoderLayer.forward()
|
||||
├─ GatedDeltaNet layers (4 of 36):
|
||||
│ ├─ PREFILL: pure PyTorch conv1d → matmul → softmax (NaN!)
|
||||
│ └─ DECODE: pure PyTorch _torch_causal_conv1d_update
|
||||
├─ MoE layers (all 36):
|
||||
│ ├─ PREFILL: PyTorch loop over unique_eids (SLOW)
|
||||
│ └─ DECODE: PyTorch batched GEMM fallback
|
||||
└─ Attention (32 of 36 layers):
|
||||
├─ PREFILL: xformers _run_sdpa_fallback (patched, matmul+softmax)
|
||||
└─ DECODE: xformers _run_sdpa_fallback
|
||||
```
|
||||
|
||||
## 3. The Crash Chain (Sub508/509 → Score 0)
|
||||
|
||||
```
|
||||
FUNCTIONAL TEST SEQUENCE:
|
||||
d01_basic_nostream ✓ PASS (95.87s — slow but works)
|
||||
d02_stream_usage ✓ PASS (1.84s)
|
||||
d03_tool_call ✗ FAIL (49.04s — model thinks instead of emitting tool XML)
|
||||
d04_reasoning ✓ PASS (128.74s)
|
||||
... more tests pass ...
|
||||
t2_n_2 ✗ FAIL → HTTP 500 → ENGINE PROCESS DIES
|
||||
↓
|
||||
t3_max_tokens_none ✗ FAIL → HTTP 500 (engine dead, Connection Refused)
|
||||
t3_max_tokens_1 ✗ FAIL → HTTP 500
|
||||
t3_max_tokens_64 ✗ FAIL → HTTP 500
|
||||
... 25 more tests ...
|
||||
t16c_empty_messages ✗ FAIL → HTTP 500
|
||||
───────────────────────────────────────
|
||||
functional score: 21/51 = 0.412 (passed before crash)
|
||||
|
||||
case_truncation → Connection Refused → score=0.0
|
||||
replay_tencent → 881/881 Connection Refused → score=0.0
|
||||
opencompass → Connection Refused → score=0.0
|
||||
───────────────────────────────────────
|
||||
TOTAL: 0.0 (engine was dead for 90% of evaluation)
|
||||
```
|
||||
|
||||
## 4. CoreX Dispatch Gap — The 52-Line Difference
|
||||
|
||||
Sub168's qwen3_5.py has ~1421 lines. Ours has 1369.
|
||||
The missing ~52 lines are CoreX dispatch wrappers:
|
||||
|
||||
```python
|
||||
# WHAT SUB168 HAS (reconstructed from log evidence):
|
||||
|
||||
# In GatedDeltaNet.__init__:
|
||||
try:
|
||||
from vllm.model_executor.models.corex_gdn import CoreXGDN
|
||||
self._corex_gdn = CoreXGDN(...) # loads libcorex_gdn.so
|
||||
except ImportError:
|
||||
self._corex_gdn = None
|
||||
|
||||
# In GatedDeltaNet.forward() prefill path:
|
||||
if self._corex_gdn is not None:
|
||||
result = self._corex_gdn.prefill(...) # → corex_gdn.py:228
|
||||
else:
|
||||
result = self._pytorch_prefill(...) # our current pure PyTorch
|
||||
|
||||
# In Qwen3_5MoE.forward():
|
||||
try:
|
||||
from vllm.model_executor.models.corex_moe import corex_moe_forward
|
||||
result = corex_moe_forward(...) # → corex_moe.py:339
|
||||
except:
|
||||
result = self._pytorch_moe_forward(...) # our current loop
|
||||
```
|
||||
|
||||
## 5. Environment Variables (already set in YAML)
|
||||
|
||||
```yaml
|
||||
VLLM_COREX_GDN_LIBRARY: /usr/local/corex/lib64/libcorex_gdn.so
|
||||
VLLM_COREX_MOE_LIBRARY: /usr/local/corex/lib64/libcorex_moe.so
|
||||
VLLM_COREX_FA2_LIBRARY: /usr/local/corex/lib64/libcorex_fa2.so
|
||||
```
|
||||
|
||||
These .so files exist in the base image. The Python wrappers
|
||||
(`corex_gdn.py`, `corex_moe.py`, `corex_fa2.py`) also exist in
|
||||
the base image at:
|
||||
`/usr/local/corex/lib/python3/dist-packages/vllm/model_executor/models/`
|
||||
|
||||
**Our qwen3_5.py simply never imports them.**
|
||||
|
||||
## 6. What Needs To Happen
|
||||
|
||||
Add try/except CoreX dispatch in 3 places in qwen3_5.py:
|
||||
1. `GatedDeltaNet.forward()` — prefill + decode paths
|
||||
2. `Qwen3_5MoE.forward()` — prefill + decode MoE dispatch
|
||||
3. Attention — already handled by xformers patches (corex_fa2 is separate)
|
||||
|
||||
CCCL pattern: `dispatch_with_env` — try native kernel first, fallback on error.
|
||||
Our Python equivalent: `try: corex_forward() except: pytorch_forward()`
|
||||
@@ -1,124 +0,0 @@
|
||||
# project_6 真实状态报告
|
||||
|
||||
生成时间: 2026-08-05, commit 96f6465
|
||||
|
||||
## 一句话总结
|
||||
|
||||
**enginex 没有 .cu 源码，gen_patch 的 C++ injection 管道全部失效。** 实际可用的优化路径只有 Python/Triton 层面的参数调优。muh 的 27 个 C++ tuning headers 是正确的架构设计，但在竞赛引擎上无处注入。
|
||||
|
||||
---
|
||||
|
||||
## 1. 竞赛引擎的致命事实
|
||||
|
||||
```
|
||||
gen_patch.py 第 47 行:
|
||||
WARNING: ALL csrc/*.cu targets are DEAD — files do not exist.
|
||||
enginex-vllm-bi100-qwen36 ships: Python + precompiled .so + Triton.
|
||||
No .cu source files. gen_patch patches have zero effect.
|
||||
```
|
||||
|
||||
enginex 交付物 = Python 文件 + 预编译 .so + Triton kernels。
|
||||
不提供 C 源码 → 无法修改 CUDA kernel → C++ tuning header 无法注入到 vllm 的编译产物里。
|
||||
|
||||
**真正的优化路径:**
|
||||
- Triton kernels (prefix_prefill.py, paged_attn.py): 可以改 BLOCK、NUM_WARPS 等 JIT 参数
|
||||
- Python 配置层 (computility-run.yaml): max_model_len、gpu_memory_utilization 等
|
||||
- 模型适配 (qwen3_5.py): MoE routing、attention 实现
|
||||
|
||||
## 2. 已有的 benchmark 数据 (真实的)
|
||||
|
||||
| 算法域 | 已跑配置数 | 来源 |
|
||||
|--------|-----------|------|
|
||||
| flash_attn | 22 configs | bi100_configs.json, SMEM 约束扫描 |
|
||||
| prefill (Triton) | 9 configs | bi100_configs.json, BLOCK×NUM_WARPS |
|
||||
| MoE | 5 configs | bi100_configs.json, BLOCK_SIZE_M |
|
||||
| reduce/scan/topk CUB | 0 | bench_bi100.py 已写但需要 BI-V100 硬件才能跑 |
|
||||
|
||||
## 3. muh C++ headers vs CCCL 覆盖率
|
||||
|
||||
| 算法 | muh 行数 | CCCL 行数 | 覆盖率 | 竞赛优先级 |
|
||||
|------|---------|---------|--------|-----------|
|
||||
| reduce | 297 | 478 | 62% | **P0** — Output TPS 83% 权重 |
|
||||
| scan | 352 | 1525 | 23% | **P0** — softmax 累积 |
|
||||
| topk | 113 | 121 | 93% | **P0** — sampling 路径 |
|
||||
| transform | 185 | 549 | 33% | P1 — RMSNorm/SiLU |
|
||||
| select_if | 459 | 2729 | 16% | P1 — token filtering |
|
||||
| radix_sort | 222 | 2381 | 9% | P1 — full sort path |
|
||||
| scan_by_key | 145 | 2008 | 7% | P1 — per-seq softmax |
|
||||
| reduce_by_key | 171 | 1735 | 9% | P1 — score aggregation |
|
||||
| unique_by_key | 166 | 1539 | 10% | P1 — KV cache dedup |
|
||||
| 其余 18 个 | 33-189 | 78-788 | 10-65% | P2 |
|
||||
|
||||
总计: muh 3618 行 vs CCCL 17000+ 行 = 平均 21% 覆盖率
|
||||
|
||||
## 4. CCCL 资产完整性
|
||||
|
||||
cccl_upstream/ 34MB, 3432 files — 是精选提取, 不是 full clone。
|
||||
|
||||
**已有 (竞赛必需的全有):**
|
||||
- 27/27 tuning headers ✓
|
||||
- 32/32 dispatch implementations ✓
|
||||
- 25/25 agent kernels ✓
|
||||
- 60/60 Thrust examples ✓
|
||||
- 243 CUB tests ✓
|
||||
- 78 CUB benchmark .cu files ✓
|
||||
- 230 Thrust tests ✓
|
||||
- 48 Thrust benchmark algorithms ✓
|
||||
|
||||
**不需要 full clone。** 缺的 ~21000 文件是 CI/CD、cudax、Python bindings、docs。
|
||||
|
||||
## 5. 真正的行动路径
|
||||
|
||||
### 短期 (功能测试通过)
|
||||
竞赛门控: 50+ 功能测试全通过 + 效果偏差 ≤ ±4%
|
||||
|
||||
关键文件:
|
||||
- `computility-run.yaml` — 控制 vllm 启动参数
|
||||
- `qwen3_6_scripts/qwen3_5.py` (588行) — MoE 模型适配
|
||||
- `prefix_prefill.py` — Triton prefill kernel, 可调 BLOCK/NUM_WARPS
|
||||
- `paged_attn.py` — Triton decode kernel
|
||||
|
||||
### 中期 (性能优化)
|
||||
目标: Token 吞吐加权值 ≥ 8000
|
||||
|
||||
```
|
||||
加权值 = Output_TPS × 16.796 + Input_TPS × 2.799 + Cache_TPS × 0.56
|
||||
```
|
||||
|
||||
**Output TPS (83%):** decode kernel → paged_attn.py Triton 参数优化
|
||||
**Input TPS (14%):** prefill kernel → prefix_prefill.py Triton 参数优化
|
||||
**Cache TPS (3%):** prefix caching 配置
|
||||
|
||||
### 长期 (如果能编译 C++)
|
||||
如果能获取 EngineX 的 C 编译环境:
|
||||
- muh C++ headers 可以直接注入
|
||||
- bench_bi100.py 的 CUB parameter sweep 可以在 BI-V100 上跑
|
||||
- 这条路 ROI 最高但依赖竞赛方提供编译链
|
||||
|
||||
## 6. 代码架构
|
||||
|
||||
```
|
||||
project_6/
|
||||
├── computility-run.yaml ← 竞赛提交配置 (直接影响评测)
|
||||
├── baseline.muh ← muh 格式的 vllm 配置
|
||||
├── Dockerfile ← 竞赛镜像构建
|
||||
├── cccl_upstream/ ← CCCL 精选 (34MB, 3432 files)
|
||||
│ ├── cub/ ← CUB: dispatch/tuning/agent/test/bench
|
||||
│ ├── thrust/ ← Thrust: examples/testing/benchmarks
|
||||
│ └── libcudacxx/ ← CUDA 标准库
|
||||
├── muh/ ← kernel tuning 框架 (544KB)
|
||||
│ ├── include/muh/tuning/ ← 27 个 BI-V100 tuning headers
|
||||
│ ├── bench_bi100.py ← CUB parameter sweep runner
|
||||
│ ├── gen_patch.py ← vllm patch 生成 (C++ 注入点已死)
|
||||
│ ├── gen_yaml.py ← computility-run.yaml 生成
|
||||
│ └── parse.py ← .muh 配置解析器
|
||||
├── muh_kernel_map.py ← CCCL 算法 → vllm kernel 映射
|
||||
├── muh_dispatch.py ← 运行时 policy 分派
|
||||
├── vllm/ ← vllm 引擎源码 (11MB Python)
|
||||
├── vllm_adapter/ ← Qwen3.5 模型适配 + 部署脚本
|
||||
├── qwen3_6_scripts/ ← Qwen3.6 patch 集合 (576KB, 25+ patches)
|
||||
├── prefix_prefill.py ← Triton prefill kernel (可调优)
|
||||
├── paged_attn.py ← Triton decode kernel (可调优)
|
||||
├── attention.py ← Attention 实现
|
||||
└── enginex-vllm-bi100-qwen36-main.zip ← 竞赛基础引擎 (97MB)
|
||||
```
|
||||
@@ -1,101 +0,0 @@
|
||||
# project_6 真实状态 v2
|
||||
|
||||
更新时间: 2026-08-06, 基于完整代码阅读
|
||||
|
||||
## 核心事实
|
||||
|
||||
**enginex 没有 .cu 源码。gen_patch 的 C++ injection 全部失效。** 但这不是终点。
|
||||
|
||||
实际可优化的三条路径：
|
||||
|
||||
### 路径 1: Triton kernel 参数调优 (直接有效)
|
||||
|
||||
文件: `prefix_prefill.py` (895行), `paged_attn.py` (794行)
|
||||
状态: 22 个 flash_attn 配置 + 9 个 prefill 配置已计算 SMEM，未上机实测
|
||||
关键参数:
|
||||
- prefill: BLOCK_M, BLOCK_N, NUM_WARPS (已有 SMEM 约束扫描)
|
||||
- decode: _PARTITION_SIZE=512 (硬编码), V1/V2 切换阈值
|
||||
- 竞赛权重: Output TPS×16.796(83%) + Input TPS×2.799(14%)
|
||||
|
||||
gen_patch.py 第 87-103 行已经指向了这些真正的 injection points:
|
||||
```python
|
||||
('prefill', 'BLOCK_M'): [('prefix_prefill.py', 'BLOCK')],
|
||||
('flash_attn', 'BLOCK_M'): [('vllm/attention/ops/triton_flash_attention.py', 'BLOCK_M')],
|
||||
('moe', 'BLOCK_SIZE_M'): [('vllm/model_executor/layers/fused_moe/fused_moe.py', 'BLOCK_SIZE_M')],
|
||||
```
|
||||
|
||||
### 路径 2: 模型适配 (功能门控)
|
||||
|
||||
文件: `vllm_adapter/qwen3_5.py` (588行), `qwen3_6_scripts/` (25+ patches)
|
||||
状态: MoE 256 experts top-8 注册完成，treat ALL layers as full attention
|
||||
待验证: TP=4 加载, reasoning 分离, tool_call parsing
|
||||
竞赛门控: 50+ 功能测试全通过 + 效果偏差 ≤±4%
|
||||
|
||||
### 路径 3: vllm Python 层配置优化 (低风险高收益)
|
||||
|
||||
文件: `computility-run.yaml`, `baseline.muh`
|
||||
关键发现 from paged_attn.py:
|
||||
- 第 99 行: `use_v1 = True` 硬编码禁用了 V2 — 对 100K token 序列这是性能杀手
|
||||
- `_PARTITION_SIZE = 512` 硬编码 — 应该根据 SM count=16 动态调整
|
||||
- `max_num_seqs: 1` — 限制了批处理并行度
|
||||
- `--enable-prefix-caching` — 已开启，但 cache copy kernel 未优化
|
||||
|
||||
## CCCL 资产的真实价值
|
||||
|
||||
CCCL 的价值不在于 C++ 注入（已证实失效），而在于：
|
||||
|
||||
1. **参数空间知识**: 27 个 tuning_*.cuh 告诉我们 NVIDIA 在 3 代 GPU 上搜索了哪些参数维度
|
||||
- reduce: ipt×tpb×ipv = 1044 个组合
|
||||
- scan: ipt×tpb×ns×dcid×l2w×trp×ld = ~26B 个（剪枝后可管理）
|
||||
- 这些维度完全适用于 Triton kernel 的等价参数
|
||||
|
||||
2. **benchmark 数据**: 199 条标注告诉我们在不同 problem size 下的加速比分布
|
||||
- 小数据量(<16M): 大多数优化无效（speedup≈1.0）
|
||||
- 大数据量(>256M): 加速比显著（最高 1.58x）
|
||||
- 这意味着 decode（小 batch）和 prefill（大 batch）需要不同策略
|
||||
|
||||
3. **约束模型**: scale_mem_bound, SMEM 公式, occupancy 计算
|
||||
- BI-V100: 16 SM, 48KB SMEM, 900GB/s BW
|
||||
- per-SM BW = 56 GB/s ≈ B200 水平
|
||||
- bytes_in_flight = 64KB (bench_bi100.py 已验证)
|
||||
|
||||
4. **算法映射**: muh_kernel_map.py 的 VLLM_KERNEL_MAP 精确映射了每个 vllm kernel 对应的 CCCL 算法
|
||||
- paged_attention → reduce (summary_statistics.cu Welford pattern)
|
||||
- softmax → scan
|
||||
- sampling → topk + radix_sort
|
||||
- normalization → transform + reduce
|
||||
|
||||
## bench_bi100.py 的实际作用
|
||||
|
||||
bench_bi100.py (713行) 是真正的工具 — 它用 PyTorch CUDA 操作模拟 CCCL benchmark:
|
||||
- 不需要编译 C++，不需要 nvbench
|
||||
- 直接在 BI-V100 上跑 torch.sum/torch.cumsum/torch.topk
|
||||
- 输出 CCCL 格式: `ipt_N.tpb_M.ipv_K speedup0 speedup1 speedup2 speedup3`
|
||||
- 搜索空间定义完整: reduce 1044 组合, scan 剪枝后可管理, topk/transform 都有
|
||||
|
||||
**但它需要 BI-V100 硬件才能跑。** 在 Phanthy Cloud 上部署就能开始标定。
|
||||
|
||||
## 代码覆盖率 (muh vs CCCL)
|
||||
|
||||
| 算法 | muh 行 | CCCL 行 | 比率 | 竞赛价值 |
|
||||
|------|--------|---------|------|---------|
|
||||
| reduce | 297 | 478 | 62% | 最高 — Output TPS 83% |
|
||||
| topk | 113 | 121 | 93% | 高 — 每次 decode |
|
||||
| scan | 370 | 1525 | 24% | 高 — softmax |
|
||||
| transform | 185 | 549 | 34% | 中 — RMSNorm/SiLU |
|
||||
| select_if | 459 | 2729 | 17% | 中 — token filter |
|
||||
| radix_sort | 222 | 2381 | 9% | 中 — full sort |
|
||||
| scan_by_key | 145 | 2008 | 7% | 中 — per-seq scan |
|
||||
| reduce_by_key | 171 | 1735 | 10% | 中 — score aggregation |
|
||||
| unique_by_key | 166 | 1539 | 11% | 低 — KV dedup |
|
||||
| 其余 18 个 | 33-189 | 78-788 | varies | 低 |
|
||||
|
||||
muh 总计 3618 行 / CCCL 17000+ 行 = 21% 平均覆盖率。
|
||||
reduce 和 topk 覆盖率最高（62%、93%），正好是竞赛权重最大的两个算法。
|
||||
|
||||
## 下一步具体行动
|
||||
|
||||
1. **在 Phanthy Cloud 上跑 bench_bi100.py** — 产出 BI-V100 真实 benchmark 数据
|
||||
2. **把 benchmark 结果回填到 Triton kernel 参数** — prefix_prefill.py 的 BLOCK/NUM_WARPS
|
||||
3. **修复 paged_attn.py 的 V2 禁用** — 对长序列性能至关重要
|
||||
4. **功能测试回归** — 确保 qwen3_5.py 适配通过 50+ 用例
|
||||
@@ -1,68 +0,0 @@
|
||||
# MoE 函数符号真相 (2026-08-17 确认)
|
||||
|
||||
## 结论
|
||||
|
||||
那5个 MoE 函数**确实不在任何镜像预装的 .so 里**。另一位开发者说的是对的。
|
||||
|
||||
但它们也**不需要**在预装 .so 里——它们是自编译的。
|
||||
|
||||
## 5个函数的正确命名空间
|
||||
|
||||
```
|
||||
ixformer::infer::topk_softmax
|
||||
ixformer::infer::moe_compute_token_index_api
|
||||
ixformer::infer::moe_expand_input
|
||||
ixformer::infer::moe_w16a16_group_gemm
|
||||
ixformer::infer::moe_output_reduce_sum
|
||||
```
|
||||
|
||||
**注意**: 是 `ixformer::infer`，不是 `ixformer::kernels::infer`。
|
||||
|
||||
## 声明 vs 实现的关系
|
||||
|
||||
| 位置 | 角色 |
|
||||
|------|------|
|
||||
| `ixformer_sdk/csrc/include/ixformer/kernels/kernels.h` | **头文件声明** (namespace `ixformer::kernels::infer`) — C++ 模板声明，给 SDK 用的 |
|
||||
| `ex_engine/csrc/moe_ops_impl.cu` | **CUDA 实现** (namespace `ixformer::infer`) — 自己写的 kernel，不依赖任何 .so |
|
||||
| `ex_engine/csrc/ix_full_bridge_v2.cpp` | **pybind11 桥** — forward-declare 然后调用 moe_ops_impl.cu 里的实现 |
|
||||
| `ex_engine/build_moe_bridge.sh` | **构建脚本** — 把 v2.cpp + moe_ops_impl.cu 一起编译成 ix_full_bridge_v2.so |
|
||||
|
||||
## 符号表搜索结果 (4个 .so 全部搜过)
|
||||
|
||||
| .so 文件 | MoE 函数 | 结论 |
|
||||
|----------|----------|------|
|
||||
| `libixformer.so` (3937 symbols) | 无 topk_softmax/moe_compute_token_index 等 | 只有 `reduce_sum` (通用的) |
|
||||
| `_ixformer_torch.so` (49 symbols) | 完全没有 MoE | 只有 norm/rope/cache/attn |
|
||||
| `_C.so` (6 symbols) | 几乎空壳 | 只有 PyInit |
|
||||
| `libcuinfer.so` (270 symbols) | 只有 cuinferTopK (不是 MoE 的) | GEMM/BLAS 级别 |
|
||||
|
||||
## 构建链
|
||||
|
||||
```
|
||||
patch_ops.sh
|
||||
└→ build_moe_bridge.sh
|
||||
└→ ninja/CppExtension 编译:
|
||||
ix_full_bridge_v2.cpp + moe_ops_impl.cu
|
||||
→ ix_full_bridge_v2.so (包含5个MoE函数的实现)
|
||||
```
|
||||
|
||||
## `ixformer::kernels::infer` vs `ixformer::infer` 的区别
|
||||
|
||||
- `ixformer::kernels::infer` — SDK 头文件 (kernels.h) 中的声明，使用 raw pointer + cudaStream_t
|
||||
- 例: `void moe_topk_softmax(const T *gating_output, T *topk_weights, int *topk_indices, ...)`
|
||||
- `ixformer::infer` — 我们自己实现的 PyTorch wrapper，使用 torch::Tensor
|
||||
- 例: `void topk_softmax(torch::Tensor& topk_weights, torch::Tensor& topk_indices, ...)`
|
||||
|
||||
`moe_ops_impl.cu` 是直接写 CUDA kernel（不调用 kernels.h 模板），然后暴露 Tensor API。
|
||||
|
||||
## Python 调用链
|
||||
|
||||
```python
|
||||
# 通过 ixformer SDK (需要真机上的 _C.so 包含 infer 子模块):
|
||||
import ixformer._C as ops
|
||||
ops.infer.moe_topk_softmax(...) # 如果 _C.so 有实现
|
||||
|
||||
# 通过 ex_engine bridge (我们自编译的):
|
||||
import ix_full_bridge_v2 as bridge
|
||||
bridge.topk_softmax(...) # 来自 moe_ops_impl.cu
|
||||
```
|
||||
@@ -1,208 +0,0 @@
|
||||
# MUH Project Checkpoint
|
||||
|
||||
> **最后更新**: 2026-07-30
|
||||
> **GitHub Project**: github.com/users/dylanyunlon/projects/6
|
||||
> **代码仓库**: github.com/dylanyunlon/project_6
|
||||
> **竞赛截止**: 2026-09-30
|
||||
|
||||
---
|
||||
|
||||
## 一、项目是什么
|
||||
|
||||
参加信创模盒 ModelHub XC 的"模型适配引擎竞赛-第一届"。目标是优化 vllm 引擎，让 Qwen3.6-35B-A3B 在天数智芯天垓100（4×BI-V100 GPU）上跑出最高的 Token 吞吐加权值。
|
||||
|
||||
**计分公式**:
|
||||
```
|
||||
Token吞吐加权值 = Output TPS × 16.796 + Input TPS × 2.799 + Cache TPS × 0.56
|
||||
```
|
||||
|
||||
Output TPS 权重占 83%——decode 阶段优化收益最大。
|
||||
|
||||
**奖项**:
|
||||
- 基础奖 200,000 积分（1:1 兑现金）: 通过全部功能/效果测试 + 性能达标(≥8000)
|
||||
- 高级奖 +100,000: 加权值提升 ≥ 30%
|
||||
- 特级奖 +50,000: 加权值提升 ≥ 50%
|
||||
|
||||
## 二、竞赛测评流程
|
||||
|
||||
参赛者提交的是 **Git 仓库地址**（在 dev.modelhub.org.cn 上）。平台自动执行：
|
||||
|
||||
1. **构建镜像**: 读取仓库根目录的 `Dockerfile`，基于基础镜像 `harbor.4pd.io/modelhubxc/enginex-iluvatar/bi100-3.2.3-x86-ubuntu20.04-py3.10-poc-llm-infer:v1.2.3` 构建
|
||||
2. **启动服务**: 读取 `computility-run.yaml` 的 `command`，在 4×天垓100 容器里启动 vllm api server（模型权重平台预挂载在 `/model`）
|
||||
3. **功能测试（门控）**: 50+ 个 OpenAI 兼容 API 测试用例，全部通过才进入下一步
|
||||
4. **效果测试（门控）**: 标准 benchmark 偏差 ≤ ±4%
|
||||
5. **性能测试（排名）**: 计算加权值
|
||||
|
||||
**你能改的**: Dockerfile + vllm 源码 + computility-run.yaml 启动参数。模型本身不能改。
|
||||
|
||||
## 三、muh 是什么
|
||||
|
||||
muh 是我们设计的 **tuning DSL（领域特定语言）**，用于：
|
||||
|
||||
1. 把 CCCL 的 tuning pattern（block_threads / items_per_thread / load_algorithm / cache_modifier 等）抽象成硬件无关的参数空间
|
||||
2. 针对天垓100 的硬件特性搜索最优参数组合
|
||||
3. Codegen 输出实际的 vllm kernel 修改 + computility-run.yaml + Dockerfile
|
||||
|
||||
**为什么需要它**: CCCL 有 27 个 tuning_*.cuh 文件（17000+ 行），每个算法都有针对不同 NVIDIA SM 架构的特化参数。天垓100 不是 NVIDIA GPU，不能直接用这些参数，但 tuning 的维度（block size、warp 策略、shared memory 用量、prefetch 策略）是通用的。muh 让迁移过程变成"改配置 + 跑 benchmark"而不是"手改 kernel + 祈祷"。
|
||||
|
||||
**muh 的状态**: v0.3 — 6个算法的C++ tuning headers已就绪(reduce/scan/topk/transform/batch_memcpy/for)，compile_test 33项通过，gen_patch.py从C++ headers提取bi100值生成vllm patches。参数值从CCCL SM100复制，等BI-V100实测替换。
|
||||
|
||||
## 四、已完成的工作
|
||||
|
||||
### 4.1 Project 6 已有 16 个真实 GitHub Issue（不是 Draft）
|
||||
|
||||
都在 `dylanyunlon/project_6` 仓库里，已关联到 GitHub Project 6，有 label 和 Priority：
|
||||
|
||||
| # | 标题 | Labels | Priority |
|
||||
|---|------|--------|----------|
|
||||
| 1 | [FEA] 非流式基础对话 | 基本功能,vllm,天垓100,Qwen3.6 | P0 |
|
||||
| 2 | [FEA] 流式对话 SSE | 基本功能,vllm | P0 |
|
||||
| 3 | [FEA] Tool Calling | 基本功能,vllm,Qwen3.6 | P0 |
|
||||
| 4 | [FEA] Reasoning/Thinking 分离 | 基本功能,thinking,Qwen3.6 | P0 |
|
||||
| 5 | [FEA] Prefix Cache | 基本功能,性能测试,vllm | P0 |
|
||||
| 6 | [FEA] 采样参数边界 | 采样参数,vllm | P1 |
|
||||
| 7 | [FEA] max_tokens 边界 | max_tokens,vllm | P1 |
|
||||
| 8 | [FEA] 结构化输出 | 结构化输出,vllm | P0 |
|
||||
| 9 | [FEA] 多语言 Emoji | 多语言,Qwen3.6 | P1 |
|
||||
| 10 | [FEA] 多模态 base64 PNG | 多模态,基本功能,Qwen3.6 | P0 |
|
||||
| 11 | [FEA] 参数校验 | 参数校验,vllm | P1 |
|
||||
| 12 | [FEA] 基础能力 | 基础能力,vllm,Qwen3.6 | P0 |
|
||||
| 13 | [FEA] 输出截断 | 截断测试,vllm | P1 |
|
||||
| 14 | [FEA] 效果测试 | 效果测试,Qwen3.6,天垓100 | P0 |
|
||||
| 15 | [EPIC] 性能基准 | 性能测试,天垓100,vllm | P0 |
|
||||
| 16 | [EPIC] 开发环境与代码提交 | infra,天垓100 | P1 |
|
||||
|
||||
这 16 个覆盖了竞赛功能测试的所有 50+ 用例。每个 issue 的 body 里都有 PND 级别的测试用例表（前置条件 + 原子步骤 + 二值判定标准）。
|
||||
|
||||
### 4.2 仓库里已有 NVIDIA CCCL 代码
|
||||
|
||||
`project_6/cccl_upstream/` 目录下包含完整的 CCCL：
|
||||
- `cub/` — GPU 原语（reduce, scan, sort, topk, block/warp/device 三层）
|
||||
- `thrust/` — 高层算法 + 60 个示例
|
||||
- `libcudacxx/` — CUDA C++ 标准库
|
||||
- `cudax/` — 实验性功能（allocators, memory resources）
|
||||
- `cub/cub/device/dispatch/tuning/` — 27 个硬件特化 tuning 文件（17000+ 行）
|
||||
|
||||
### 4.3 Label 体系已建立
|
||||
|
||||
仓库上已创建 16 个 label：基本功能、thinking、采样参数、max_tokens、基础能力、结构化输出、多语言、多模态、参数校验、截断测试、效果测试、性能测试、infra、vllm、天垓100、Qwen3.6
|
||||
|
||||
### 4.4 Project 6 里有 15 个遗留 Draft Issue 需要清理
|
||||
|
||||
这些是早期用 addProjectV2DraftIssue 创建的，没有 repo 关联、没有 label。应该从 Project 面板里手动删除。
|
||||
|
||||
## 五、还没做的（下一步）
|
||||
|
||||
1. ~~muh 语言 PRD 设计~~ ✅ Done — muh是C++ header-only lib，不是独立语言
|
||||
2. ~~从 CCCL tuning_*.cuh 提取参数空间~~ ✅ Done — 6个算法的policy_selector已实现
|
||||
3. **在BI-V100上跑benchmark** — 用实测数据替换bi100_*中的SM100复制值
|
||||
4. **获取 enginex-vllm-bi100-qwen36 的实际代码** — 需要在 Phanthy Cloud 开发环境里操作
|
||||
5. **设计 muh → vllm kernel 的 codegen 管道**
|
||||
6. **实际在天垓100 上跑 benchmark**
|
||||
|
||||
## 六、参考项目
|
||||
|
||||
- **NVIDIA CCCL Project #6**: github.com/orgs/NVIDIA/projects/6（1990 items，Issue-first 模式，label 做模块分类）
|
||||
- **pub/sub-loop Project #4**: github.com/users/dylanyunlon/projects/4（1632 items，Draft-first 模式，已验证 1111 个有真实测试步骤，154 个有"按AC验证"占位符）
|
||||
- **PND 测试库**: 818 条车载软件测试用例，作为 PRD 测试用例质量基准
|
||||
|
||||
## 七、关键文件路径
|
||||
|
||||
```
|
||||
project_6/
|
||||
├── cccl_upstream/ # NVIDIA CCCL 完整代码
|
||||
│ ├── cub/cub/device/dispatch/tuning/ # 27 个 tuning policy 文件
|
||||
│ ├── cub/cub/warp/ # warp-level 原语
|
||||
│ ├── cub/cub/block/ # block-level 原语
|
||||
│ ├── thrust/examples/ # 60 个优化模式示例
|
||||
│ └── cudax/...allocators/ # 内存分配器
|
||||
├── Dockerfile # TODO: 待创建
|
||||
├── computility-run.yaml # TODO: 待创建
|
||||
└── muh/ # TODO: muh 语言实现
|
||||
```
|
||||
|
||||
## 八、竞赛关键参数（来自 computility-run.yaml 参考）
|
||||
|
||||
```yaml
|
||||
concurrency: 1
|
||||
command:
|
||||
- python3 -m vllm.entrypoints.openai.api_server
|
||||
- --model /model
|
||||
- --served-model-name llm
|
||||
- --max-model-len 100000
|
||||
- --gpu-memory-utilization 0.9
|
||||
- -tp 4
|
||||
- --max-num-seqs 1
|
||||
- --max-num-batched-tokens 8192
|
||||
- --enable-chunked-prefill
|
||||
- --max-seq-len-to-capture 32768
|
||||
- --enable-auto-tool-choice
|
||||
- --tool-call-parser qwen3_coder
|
||||
- --reasoning-parser qwen3
|
||||
- --enable-prefix-caching
|
||||
env:
|
||||
- name: VLLM_ENGINE_ITERATION_TIMEOUT_S
|
||||
value: 3600
|
||||
```
|
||||
|
||||
基础镜像: `harbor.4pd.io/modelhubxc/enginex-iluvatar/bi100-3.2.3-x86-ubuntu20.04-py3.10-poc-llm-infer:v1.2.3`
|
||||
|
||||
## 九、CCCL Tuning 文件全量模型输入记录
|
||||
|
||||
**所有 27 个 tuning_*.cuh 文件的完整源码已在本 context 中作为模型输入读取。** 关键发现：
|
||||
|
||||
### policy_selector 统一模式
|
||||
|
||||
每个算法都有一个 `policy_selector` struct，接受 `::cuda::compute_capability cc` 参数，内部按 SM 版本做 if-else 分支：
|
||||
|
||||
```
|
||||
if (cc >= {10, 0}) → sm100 tuning (Blackwell)
|
||||
if (cc >= {9, 0}) → sm90 tuning (Hopper)
|
||||
if (cc >= {8, 0}) → sm80 tuning (Ampere)
|
||||
if (cc >= {7, 0}) → sm70 tuning (Volta)
|
||||
if (cc >= {6, 0}) → sm60 tuning (Pascal)
|
||||
fallback → sm50 tuning
|
||||
```
|
||||
|
||||
**muh 的核心工作就是给每个 policy_selector 添加一个 `cc == {iluvatar, 100}` 分支，填入在天垓100 上跑出的最优 benchmark 数据。**
|
||||
|
||||
### 各算法提取的参数维度
|
||||
|
||||
| 算法 | 文件 | 行数 | 参数维度 |
|
||||
|------|------|------|---------|
|
||||
| reduce | tuning_reduce.cuh | 478 | threads, items, vec_size, reduce_algorithm, load_modifier, determinism |
|
||||
| scan | tuning_scan.cuh | 1525 | threads, items, load_algo, load_mod, store_algo, scan_algo, delay_policy + lookahead variant |
|
||||
| radix_sort | tuning_radix_sort.cuh | 2381 | histogram(threads,items,partitions,radix_bits) + exclusive_sum + onesweep(threads,items,store,rank,scan,partitions,radix_bits) + downsweep + upsweep + single_tile |
|
||||
| reduce_by_key | tuning_reduce_by_key.cuh | 1735 | threads, items, load_algo, load_mod, scan_algo, delay_policy |
|
||||
| select_if | tuning_select_if.cuh | 2729 | threads, items, load_algo, load_mod, scan_algo, delay_policy |
|
||||
| histogram | tuning_histogram.cuh | 363 | threads, pixels_per_thread, vec_size, load_algo, load_mod, rle_compress, mem_preference, work_stealing |
|
||||
| topk | tuning_topk.cuh | 121 | threads, items (simple, no SM-specific tuning yet) |
|
||||
| batched_topk | tuning_batched_topk.cuh | 186 | worker_policy array × 6 tiers + multi_worker_policy |
|
||||
| merge | tuning_merge.cuh | 180 | threads, items, load_mod, store_algo, bulk_copy_keys, bulk_copy_values |
|
||||
| merge_sort | tuning_merge_sort.cuh | 193 | threads, items, load_algo, load_mod, store_algo |
|
||||
| transform | tuning_transform.cuh | 549 | threads, items, load_algo, store_algo, load_mod |
|
||||
| rle_encode | tuning_rle_encode.cuh | 626 | threads, items, load_algo, load_mod, scan_algo, delay_policy |
|
||||
| rle_non_trivial | tuning_rle_non_trivial_runs.cuh | 691 | threads, items, load_algo, load_mod, store_time_slicing, scan_algo, delay |
|
||||
| adjacent_diff | tuning_adjacent_difference.cuh | 118 | threads, items, load_algo, load_mod, store_algo (single policy, no SM branching) |
|
||||
| for | tuning_for.cuh | 78 | threads, items (trivial, 256×2) |
|
||||
| find | tuning_find.cuh | 90 | threads, items, vec_size, load_mod |
|
||||
| batch_memcpy | tuning_batch_memcpy.cuh | 227 | small_buffer + large_buffer sub-policies |
|
||||
| scan_by_key | tuning_scan_by_key.cuh | ~2000 | same as reduce_by_key pattern |
|
||||
| unique_by_key | tuning_unique_by_key.cuh | ~1500 | same pattern |
|
||||
| three_way_partition | tuning_three_way_partition.cuh | ~780 | same pattern |
|
||||
| segmented_* | 4 files | ~1300 total | segmented variants of reduce/scan/sort |
|
||||
|
||||
### Benchmark 注释格式
|
||||
|
||||
每个 sm100 tuning 都有注释格式：
|
||||
```
|
||||
// ipt_22.tpb_384.ns_1904.dcid_6.l2w_830.trp_1.ld_0 1.148442 0.997167 1.139902 1.462651
|
||||
```
|
||||
- `ipt` = items_per_thread
|
||||
- `tpb` = threads_per_block
|
||||
- `ns` = delay nanoseconds
|
||||
- `dcid` = delay constructor ID
|
||||
- `l2w` = L2 cache window
|
||||
- `trp` = transpose (0=DIRECT, 1=WARP_TRANSPOSE)
|
||||
- `ld` = load modifier (0=DEFAULT, 1=LDG, 2=CA)
|
||||
- 4 个数字 = 4 种 problem size 下的加速比 (vs 前代 SM)
|
||||
@@ -1,85 +0,0 @@
|
||||
# muh Tuning Gap Analysis — CCCL vs BI-V100 适配
|
||||
## 2026-08-07
|
||||
|
||||
### 方法论
|
||||
|
||||
直接读取 CCCL 源码（26 个 tuning_*.cuh），提取竞赛相关的 benchmark annotations，
|
||||
对比 muh 已有的 BI-V100 struct 值。每个算法的优先级由竞赛评分公式决定：
|
||||
|
||||
```
|
||||
Score = Output_TPS × 16.796 + Input_TPS × 2.799 + Cache_TPS × 0.56
|
||||
```
|
||||
|
||||
Output TPS = 83%, Input TPS = 14%, Cache TPS = 3%
|
||||
|
||||
---
|
||||
|
||||
### P0: 直接影响竞赛评分的算法
|
||||
|
||||
#### 1. REDUCE (Output TPS 83%) — ★★★★★
|
||||
- **竞赛路径**: paged_attention score reduction, float32, plus
|
||||
- **CCCL SM100**: `ipt_16.tpb_512.ipv_2 → 1.061/1.000/1.065/1.167`
|
||||
- **muh BI-V100**: `bi100_plus_float32_o4 {512, 24, 2}` — tile=12288 (1.5× SM100)
|
||||
- **状态**: ✅ 完成 (62% 行覆盖)
|
||||
- **待定**: SM=16 items 适配 (P0 BUG)、LOAD_LDG vs LOAD_DEFAULT benchmark
|
||||
|
||||
#### 2. SCAN (Output TPS 83%) — ★★★★☆
|
||||
- **竞赛路径**: softmax denominator prefix sum, float32, plus
|
||||
- **CCCL SM100**: `ipt_22.tpb_384.ns_1904.dcid_6.l2w_830 → 1.148/0.997/1.140/1.463`
|
||||
- **muh BI-V100**: `bi100_lookback_4B_o4 {384, 22}` — 与 SM100 同 tile
|
||||
- **状态**: ✅ 核心完成 (39% 行覆盖，lookback + SM90 fallback)
|
||||
- **待定**: Lookback delay 参数需实测校准、8B structs 99% SMEM 需验证
|
||||
|
||||
#### 3. TRANSFORM (Input TPS 14% + all activations) — ★★★★☆
|
||||
- **竞赛路径**: SiLU/GeLU/RMSNorm, bfloat16
|
||||
- **CCCL**: bytes_in_flight 是核心参数, B200=64KB, H100=48KB
|
||||
- **muh BI-V100**: bytes_in_flight=64KB (confirmed by babelstream bench)
|
||||
- **状态**: ✅ 核心完成
|
||||
- **待定**: Vectorized vs prefetch algorithm 选择需实测
|
||||
|
||||
---
|
||||
|
||||
### P1: 间接影响性能的算法
|
||||
|
||||
#### 4. TOPK (sampling, Output TPS) — ★★★☆☆
|
||||
- **竞赛路径**: logit sampling, float32 keys
|
||||
- **CCCL**: bits_per_pass, thread count, BLOCK_SCAN_WARP_SCANS
|
||||
- **muh BI-V100**: 有 inline tuning (threads=512, bits_per_pass=11)
|
||||
- **状态**: ✅ 基本完成
|
||||
- **待定**: Onesweep vs multi-sweep 选择
|
||||
|
||||
#### 5. SELECT_IF (MoE routing) — ★★☆☆☆
|
||||
- **竞赛路径**: expert selection, float32, not_flagged, no_rejects, offset_4
|
||||
- **CCCL SM80**: `{threads=256, items=18, WARP_TRANSPOSE, no_delay=1130}`
|
||||
- **muh BI-V100**: 零 bi100 structs, 用 get_sm100_adapted() inline 计算
|
||||
- **状态**: ⚠️ 只需 1/77 个 specialization, 但完全缺失
|
||||
- **待定**: 需添加 bi100_select_float32_nf_nr_o4 struct
|
||||
|
||||
#### 6. RADIX_SORT (topk helper) — ★★☆☆☆
|
||||
- **竞赛路径**: float32 key sort for sampling
|
||||
- **CCCL**: 2381 行, onesweep + histogram, SM100 有复杂分支
|
||||
- **muh BI-V100**: 222 行 (9% 覆盖)
|
||||
- **状态**: ⚠️ 需要 onesweep 路径
|
||||
- **待定**: bits_per_pass 和 histogram SMEM
|
||||
|
||||
---
|
||||
|
||||
### P2: 理论覆盖但不直接影响评分
|
||||
|
||||
| 算法 | CCCL 行数 | muh 行数 | 覆盖率 | 竞赛影响 |
|
||||
|------|----------|---------|-------|---------|
|
||||
| reduce_by_key | 1735 | 217 | 13% | 低 |
|
||||
| scan_by_key | 2008 | 161 | 8% | 低 |
|
||||
| unique_by_key | 1510 | 179 | 12% | 低 |
|
||||
| three_way_partition | 708 | 67 | 9% | 低 |
|
||||
| segmented_reduce | 471 | 112 | 24% | 低 |
|
||||
| 其余 14 个 | ~4000 | ~800 | ~20% | 无 |
|
||||
|
||||
---
|
||||
|
||||
### 关键差距总结
|
||||
|
||||
1. **gen_patch.py 管道断裂** — 产出零 patch。已被 gen_config.py 替代。
|
||||
2. **muh headers 20% 完成** — 但竞赛相关的 5 个算法 (reduce/scan/transform/topk/select_if) 核心参数已就位。
|
||||
3. **缺 benchmark 验证** — 所有 BI-V100 speedup 标 TBD，需要在 Phanthy Cloud 上跑。
|
||||
4. **Python layer 是真正的注入点** — 已在 triton_flash_attention.py 添加 8 个 BI-V100 configs, prefix_prefill.py 修 BLOCK=64, _custom_ops.py 修 SMEM=48KB。gen_config.py 又发现 19 个新候选 configs。
|
||||
@@ -1,50 +0,0 @@
|
||||
# muh Pipeline Ground Truth — 2026-08-07
|
||||
|
||||
## 管道实际状态（不是设计稿，是已部署代码的真实描述）
|
||||
|
||||
### scale_mem_bound: FULL PARITY ✓
|
||||
11/11测试用例与CCCL `cub::detail::scale_mem_bound` 完全匹配。
|
||||
返回值顺序 `{items_per_thread, threads_per_block}` — items-first，与CCCL一致。
|
||||
|
||||
### C++ Tuning Headers: 27/27 ✓
|
||||
所有26个算法（+common）都有bi100 header，`policy_selector::operator()` 接受
|
||||
`hardware_capability` 参数。SMEM overflow保护覆盖所有type_size。
|
||||
|
||||
### Injection现状（enginex没有.cu源码）
|
||||
|
||||
| 注入位置 | 状态 | 值 | commit |
|
||||
|---------|------|-----|--------|
|
||||
| prefix_prefill.py BLOCK | ✓ 已手动修改 | BLOCK=64, WARPS=4 | 多个commit |
|
||||
| paged_attn.py _PARTITION_SIZE | ✓ 保持默认 | 512 | — |
|
||||
| paged_attn.py V1/V2 dispatch | ✓ 已手动修改 | use_v1 threshold | cbd1f08 |
|
||||
| _custom_ops.py SMEM | ✓ 已手动修改 | 48KB | 16f0b30 |
|
||||
| triton_flash_attention.py | ✓ 已添加BI-V100 configs | BLOCK=32/64 | 多个commit |
|
||||
| protocol.py 兼容性 | ✓ 已修复 | max_completion_tokens等 | 2c353da |
|
||||
|
||||
### gen_patch.py 角色
|
||||
设计时期望: C++ header → unified diff → vllm .cu文件
|
||||
实际情况: enginex只有Python + .so, 没有.cu源码
|
||||
当前角色: 文档工具 + 验证（确认header值与已部署Python代码一致）
|
||||
|
||||
### CCCL SM100 Benchmark数据（从源码提取，已存入cccl_sm100_benchmark_values.json）
|
||||
|
||||
**Reduce** (paged_attention score reduction, Output TPS 83%权重):
|
||||
- float32+plus: items=16, threads=512, vec=2, speedup=[1.061, 1.000, 1.065, 1.167]
|
||||
- float64+plus: items=16, threads=640, vec=1, speedup=[1.018, 1.000, 1.016, 1.057]
|
||||
|
||||
**Scan** (softmax prefix-sum):
|
||||
- 4B lookback: items=22, threads=384, delay=1904ns/dcid=6/l2w=830, speedup=[1.148, 0.997, 1.140, 1.463]
|
||||
- 8B lookback: items=23, threads=416, delay=772ns/dcid=5/l2w=710, speedup=[1.089, 1.016, 1.086, 1.265]
|
||||
|
||||
**muh BI-V100适配**:
|
||||
- reduce float32: items=24(+50%), threads=512(=), vec=2(=) → 补偿16 SMs
|
||||
- scan 4B: 通过scale_mem_bound自动适配(items=22 @4B安全, @8B降级到16)
|
||||
- delay参数: ns×0.5, l2w×0.6 (启发式, 待实测)
|
||||
|
||||
### 竞赛门槛
|
||||
- 功能测试: 50+ TC, 项目看板14个FEA item覆盖
|
||||
- 效果测试: benchmark偏差 ≤ ±4%
|
||||
- 性能测试: Token吞吐加权值 ≥ 8000
|
||||
- Output TPS × 16.796 (83%) → reduce/scan/topk
|
||||
- Input TPS × 2.799 (14%) → scan/transform
|
||||
- Cache TPS × 0.56 (3%) → batch_memcpy
|
||||
@@ -1,71 +0,0 @@
|
||||
# muh 管道现实检查 — 2026-08-07
|
||||
|
||||
## 核心发现
|
||||
|
||||
### 1. gen_patch.py 输出为零
|
||||
|
||||
```
|
||||
$ python3 muh/gen_patch.py --dry-run
|
||||
READ reduce: bi100_plus_float32_o4 → {items: 24, threads: 512, vec: 2}
|
||||
READ scan: bi100_sm90_float32 → {threads: 128, items: 24}
|
||||
...
|
||||
No patches generated.
|
||||
```
|
||||
|
||||
原因: `VLLM_INJECTION_POINTS` 的 key `('reduce', 'partition_size')` 和 struct 提取出的 field `items`/`threads`/`vec` 不匹配。gen_patch 的"读"和"写"两端从未对齐。
|
||||
|
||||
### 2. 注入目标是 Python 不是 C++
|
||||
|
||||
enginex-vllm-bi100 **没有 `.cu` 源码**。所有 CUDA kernel 是预编译的 ixformer `.so`。
|
||||
|
||||
实际可调的全部是 Python 层:
|
||||
|
||||
| 文件 | 可调参数 | 竞赛影响 |
|
||||
|------|---------|---------|
|
||||
| `paged_attn.py` | `_PARTITION_SIZE=512`, V1/V2 dispatch logic | Output TPS (83%) |
|
||||
| `prefix_prefill.py` | `BLOCK=64`, `BLOCK_N=64`, `NUM_WARPS=4` | Input TPS (14%) |
|
||||
| `vllm/attention/ops/triton_flash_attention.py` | 17 个 autotune configs | Prefill throughput |
|
||||
| `vllm/_custom_ops.py` | `return 49152` (SMEM fix) | 所有 Triton kernels |
|
||||
| `computility-run.yaml` | `--max-num-seqs`, `--gpu-memory-utilization` | 调度效率 |
|
||||
|
||||
gen_patch.py 中的 `csrc/*.cu` 注入点全部是 dead code (注释已标注)。
|
||||
|
||||
### 3. muh C++ headers 的实际价值
|
||||
|
||||
muh 的 26 个 tuning headers 和 `scale_mem_bound` 实现是正确的理论分析工具。它们的价值不在于直接注入 vllm，而在于:
|
||||
|
||||
- 推导 SMEM 约束 (Triton `BLOCK_M × head_dim × elem_size` 上限)
|
||||
- 推导 occupancy 模型 (BI-V100 16 SMs 的 wave efficiency)
|
||||
- 推导 bytes_in_flight (56 GB/s per-SM → 64KB prefetch window → `num_stages=2`)
|
||||
- 为 CCCL benchmark 验证提供 ground truth
|
||||
|
||||
这些推导已经手工应用到了 Python 代码中:
|
||||
- `triton_flash_attention.py` 的 8 个 BI-V100 configs 引用了 CCCL babelstream/scan 分析
|
||||
- `prefix_prefill.py` 的 BLOCK_N=64 推导基于 48KB SMEM 约束
|
||||
- `_custom_ops.py` 的 49152 来自 hardware.cuh
|
||||
|
||||
### 4. 管道闭环的正确路径
|
||||
|
||||
```
|
||||
CCCL tuning analysis Python layer injection Triton autotune
|
||||
(理论推导) (参数修改) (运行时选择)
|
||||
│ │ │
|
||||
▼ ▼ ▼
|
||||
muh headers paged_attn.py triton.Config([...])
|
||||
common.cuh prefix_prefill.py autotune picks best
|
||||
hardware.cuh _custom_ops.py at runtime
|
||||
│ │ │
|
||||
└───────────────────────┴───────────────────────┘
|
||||
│
|
||||
竞赛评测得分
|
||||
```
|
||||
|
||||
不是: `muh headers → gen_patch → #define injection → recompile`
|
||||
而是: `muh analysis → Python config → Triton autotune → runtime perf`
|
||||
|
||||
## 下一步
|
||||
|
||||
1. 删除 gen_patch.py 中所有 dead `csrc/*.cu` 注入点
|
||||
2. 重写 gen_patch 为 `gen_config.py`: 从 muh headers 推导 → 直接输出 Python patch
|
||||
3. 用 CCCL benchmarks 验证: reduce/sum.cu, scan/exclusive/sum.cu, topk/keys.cu
|
||||
4. 扩展 triton_flash_attention.py autotune 搜索空间 (当前 17 configs, 可加到 30+)
|
||||
@@ -1,86 +0,0 @@
|
||||
# muh Pipeline Status — Ground Truth
|
||||
|
||||
**Last verified**: 2026-08-07T01:45:31Z by automated analysis
|
||||
|
||||
## Architecture Summary
|
||||
|
||||
```
|
||||
CCCL policy_selector(compute_capability) → ReducePolicy{threads, items, vec, algo, load_mod}
|
||||
↕ mirrors
|
||||
muh policy_selector(hardware_capability) → same struct types, BI-V100 values
|
||||
↕ gen_patch.py extracts bi100_* values
|
||||
vllm patch_ops.sh → full-file Python replacements with tuning values baked in
|
||||
```
|
||||
|
||||
## Injection Reality
|
||||
|
||||
### What gen_patch.py THINKS (csrc/*.cu — DEAD)
|
||||
```
|
||||
tuning_reduce.cuh → csrc/attention/attention_kernels.cu NUM_THREADS ← NO .cu SOURCE
|
||||
tuning_scan.cuh → csrc/attention/paged_attention_v1.cu SCAN_BLOCK_SIZE ← NO .cu SOURCE
|
||||
tuning_topk.cuh → csrc/sampling/sampling_kernels.cu SAMPLING_BLOCK_SIZE ← NO .cu SOURCE
|
||||
```
|
||||
|
||||
### What ACTUALLY happens (Python runtime — ALIVE)
|
||||
```
|
||||
_custom_ops.py → SMEM 49152 (was 32768) ← DEPLOYED ✓
|
||||
paged_attn.py → _PARTITION_SIZE=512 ← DEPLOYED ✓ (V2 partition, NOT CTA tile)
|
||||
xformers.py → _Q_CHUNK=256, sdpa_fallback ← DEPLOYED ✓
|
||||
sampler.py → torch.topk fast path ← DEPLOYED ✓
|
||||
prefix_prefill.py → Triton BLOCK_M/N/warps ← DEPLOYED ✓ (but Triton not available)
|
||||
computility-run.yaml → vllm server args ← DEPLOYED ✓
|
||||
```
|
||||
|
||||
### The Gap
|
||||
muh C++ headers define precise per-type-per-op tuning values (14 reduce structs, 22 scan structs).
|
||||
But the vllm engine on BI-V100 runs ixformer .so (precompiled, not tunable) + Python fallbacks.
|
||||
The C++ headers' values cannot be injected into the precompiled .so.
|
||||
They CAN inform:
|
||||
1. Python fallback implementations (paged_attn.py, xformers.py) — tile sizes, chunk sizes
|
||||
2. Triton JIT configs — if Triton were available (it's not on BI-V100 base image)
|
||||
3. Future EngineX releases that expose tuning knobs
|
||||
|
||||
## Asset Inventory
|
||||
|
||||
| Asset | Count | Status |
|
||||
|-------|-------|--------|
|
||||
| CCCL tuning headers (upstream) | 27 | Complete |
|
||||
| muh BI-V100 headers | 27 | Complete (14 reduce + 22 scan + others) |
|
||||
| muh schema YAMLs | 27 | Complete |
|
||||
| CUB benchmarks | 91 | Synced to NVIDIA/cccl main |
|
||||
| CUB tests | 243 | Complete |
|
||||
| CUB examples | 18 | Complete |
|
||||
| Thrust examples | 60 | Complete |
|
||||
| Deployed patches | 15 files | Via patch_ops.sh full replacement |
|
||||
| bench_bi100.py search spaces | 5 algos | Defined, needs BI-V100 hardware to run |
|
||||
|
||||
## Tool Chain Status
|
||||
|
||||
| Tool | Input | Output | Status |
|
||||
|------|-------|--------|--------|
|
||||
| parse.py | baseline.muh | JSON config | ✓ Working |
|
||||
| gen_patch.py | tuning_*.cuh | Patch report | ⚠ Reports structs but generates 0 patches (injection mapping mismatch) |
|
||||
| gen_yaml.py | baseline.muh | computility-run.yaml | ✓ Working |
|
||||
| bench_bi100.py | algo+dtype | CCCL-format speedup data | Needs BI-V100 hardware |
|
||||
| patch_ops.sh | qwen3_6_scripts/ | Docker vllm patches | ✓ Working |
|
||||
| muh_dispatch.py | hw+dtype+head_dim | AttentionConfig | ✓ Working (needs torch) |
|
||||
| scale_mem_bound | (threads, items, type_size) | (items, threads) | ✓ CCCL parity verified |
|
||||
|
||||
## Critical Numbers
|
||||
|
||||
| Metric | Competition Threshold | Current Status |
|
||||
|--------|----------------------|----------------|
|
||||
| Functional tests | 50+ pass | 13 items In Progress (all FEA) |
|
||||
| Effect deviation | ≤ ±4% | Untested (needs hardware) |
|
||||
| Token throughput weighted | ≥ 8000 | Untested |
|
||||
| Output TPS weight | 83% (×16.796) | Reduce/scan/topk optimization focus |
|
||||
| SMEM limit | 49152 bytes | All 36 scan+reduce structs verified ✓ |
|
||||
| SM count | 16 (confirmed) | All headers updated |
|
||||
|
||||
## Next Actions (Ranked by Competition Impact)
|
||||
|
||||
1. **Run bench_bi100.py on BI-V100** → get real speedup data for reduce/scan/topk
|
||||
2. **Backfill speedup data to muh headers** → replace TBD/theoretical values
|
||||
3. **Optimize Python fallback tile sizes** → paged_attn.py, xformers.py Q_CHUNK
|
||||
4. **Tune computility-run.yaml** → max-num-seqs, max-batched-tokens, gpu-mem-util
|
||||
5. **Enable prefix caching benchmark** → cached_tokens > 0 for repeat prompts
|
||||
6
PRD.md
6
PRD.md
@@ -8,3 +8,9 @@
|
||||
- 效果偏差 ≤±4%
|
||||
- 性能门槛 Token 吞吐加权值 ≥8000
|
||||
- Output TPS 权重占 83%（decode kernel 优化投入产出比最高）
|
||||
|
||||
AllReduce 大概占 10ms。剩下的 36ms 是 Python dispatch。1400 次 PyTorch 函数调用 × 25 微秒。
|
||||
|
||||
这台机器有没有 NVLink 改变不了 Python 每次调用花 25 微秒的事实。NVIDIA 上用 CUDA Graph 一次性录制所有 kernel launch，replay 时零 Python 开销。但 BI-V100 CUDA 10.2 对 Graph 支持有限。
|
||||
|
||||
最大的问题是 太多小 kernel 走 Python dispatch。减少 launch 次数比优化任何单个 kernel 都有效。
|
||||
@@ -4,119 +4,3 @@
|
||||
天垓100 (BI-V100) 推理引擎竞赛，在 4×BI-V100 上运行 Qwen3.5-27B 推理服务。
|
||||
竞赛目标：Token吞吐加权值 ≥ 8000（Output TPS × 83% + Input TPS × 14% + Cache TPS × 3%）
|
||||
|
||||
## 技术栈
|
||||
- Base image: bi100-3.2.3-x86-ubuntu20.04-py3.10-poc-llm-infer:v1.2.3
|
||||
- vLLM 0.6.3 (base) + serving层patch
|
||||
- ixformer (CoreX SDK, 含 flash_attn / paged_attention / silu_and_mul 等)
|
||||
- Tensor Parallel = 4, enforce_eager=True
|
||||
|
||||
## 文件结构
|
||||
|
||||
```
|
||||
project_6/
|
||||
├── PRD.md # 竞赛需求 + CCCL→base映射
|
||||
├── SYSTEM_DESIGN.md # 架构设计: Docker/Build/Runtime/GDN dispatch
|
||||
├── Dockerfile # Docker构建
|
||||
├── computility-run.yaml # vLLM启动参数
|
||||
├── qwen3_6_scripts/ # serving层 + model patches (部署到vllm)
|
||||
│ ├── qwen3_5.py (2040行) 模型代码: GDN + MoE + Attention
|
||||
│ ├── serving_chat.py OpenAI API处理核心
|
||||
│ ├── protocol.py 请求/响应模型
|
||||
│ ├── api_server.py FastAPI入口
|
||||
│ ├── patch_ops.sh 部署脚本 (全部patch的安装器)
|
||||
│ ├── flash_qla_sm70/ GDN CUDA kernel (gdn_forward.cu 1919行)
|
||||
│ └── ... 其他patches
|
||||
├── ex_engine/ # EX引擎: 算法因子置换层
|
||||
│ ├── csrc/
|
||||
│ │ ├── ix_full_bridge.cpp (331行) pybind11桥接→ixformer::infer 14个C++函数
|
||||
│ │ ├── ix_moe_bridge.cpp (258行) MoE-only子集桥接
|
||||
│ │ └── moe_topk_softmax_v3.cu (148行) 独立CUDA topk kernel
|
||||
│ ├── python/
|
||||
│ │ ├── corex_moe.py (196行) MoE分发: ix_bridge→ixformer::infer 7步pipeline
|
||||
│ │ ├── corex_gdn.py (217行) GDN分发: chunked delta rule + decode
|
||||
│ │ ├── corex_fa2.py (228行) FA2分发: packed/paged/chunked三模式
|
||||
│ │ ├── ix_bridge.py (162行) ix_full_bridge.so加载器
|
||||
│ │ └── moe_topk.py CUDA topk Python wrapper
|
||||
│ ├── build.sh 编译脚本 (corex clang/16)
|
||||
│ └── include/ C++ headers
|
||||
├── cccl_upstream/ (8900文件) NVIDIA CCCL strategic subset
|
||||
│ ├── cub/ tuning headers + benchmarks + tests
|
||||
│ ├── thrust/ examples + tests
|
||||
│ └── libcudacxx/ C++ STL headers
|
||||
├── muh/ muh工具链: BI-V100 tuning parameter生成
|
||||
│ ├── include/muh/tuning/ 27个BI-V100 policy_selector headers
|
||||
│ └── gen_patch.py C++ header → vllm unified diff
|
||||
├── upstream_ref/ 上游参考代码
|
||||
│ ├── ds_vllm/ ds-vllm (vllm fork, 含topk_softmax_kernels.cu)
|
||||
│ └── xllm/ xllm (ILU backend: kernels/ilu + layers/ilu)
|
||||
├── vllm/ vllm源码副本 (参考用)
|
||||
└── docs/ 分析文档
|
||||
```
|
||||
|
||||
## 关键文件说明
|
||||
|
||||
### ex_engine/csrc/ix_full_bridge.cpp
|
||||
- `ix_topk_softmax()` → `ixformer::infer::topk_softmax`
|
||||
- `ix_moe_gen_idx()` → `ixformer::infer::moe_compute_token_index_api`
|
||||
- `ix_moe_expand_input()` → `ixformer::infer::moe_expand_input`
|
||||
- `ix_group_gemm()` → `ixformer::infer::moe_w16a16_group_gemm`
|
||||
- `ix_silu_and_mul()` → `ixformer::infer::silu_and_mul`
|
||||
- `ix_moe_combine_result()` → `ixformer::infer::moe_output_reduce_sum`
|
||||
- `ix_fused_moe_forward()` — 以上6步组合, 一次C++调用完成整个MoE
|
||||
- `ix_paged_attention()` → `ixformer::infer::xllm_paged_attention`
|
||||
- `ix_flash_attn_prefill()` → `ixformer::infer::ixinfer_flash_attn_unpad_with_block_tables`
|
||||
- `ix_rms_norm()` / `ix_fused_add_rms_norm()` / `ix_rotary_embedding()` / `ix_reshape_and_cache()`
|
||||
|
||||
### ex_engine/python/corex_moe.py
|
||||
- `moe_forward()` — 3级分发: ix_bridge全C++ → ix_bridge逐步 → Python loop
|
||||
- `topk_softmax()` — ix_bridge优先, fallback到Python softmax+topk
|
||||
- `moe_prefill()` / `moe_decode()` — 日志匹配comp 168格式
|
||||
|
||||
### qwen3_6_scripts/qwen3_5.py
|
||||
- `GatedDeltaNet.forward()` — GDN层: corex_gdn dispatch
|
||||
- `Qwen3_5MoE.forward()` — MoE层: Tier 0-3分发 (ix_fused_moe → ix_bridge → corex_moe → PyTorch)
|
||||
|
||||
## 当前状态
|
||||
- 370+ commits, 67 GitHub issues (63 open, 4 closed)
|
||||
- GitHub Project #6: 149 items (121 draft issues + 28 real issues)
|
||||
- CCCL upstream (5205 files) 作为工程基座, tuning/dispatch pattern 1:1映射
|
||||
- 真机 comp 168 日志已完整分析: 3个致命bug已定位并修复
|
||||
- 可提交竞赛平台测试
|
||||
|
||||
## 本次任务完成内容
|
||||
comp 168 docker日志 + upstream_ref 系统设计分析 → 三个致命bug修复:
|
||||
|
||||
1. **OOM修复**: computility-run.yaml max_model_len 256000→80000
|
||||
- comp 168日志: `torch.cuda.OutOfMemoryError: Tried to allocate 32.00 MiB`
|
||||
- 引擎OOM→崩溃→replay_tencent 881请求中704个 Connection refused
|
||||
- BI-V100 KV cache容量~88112 blocks, 256000远超上限
|
||||
|
||||
2. **topk_softmax ERROR日志消除**: _custom_ops.py silent fallback
|
||||
- comp 168日志: `ixformer.functions has no attribute vllm_moe_topk_softmax` × 500+次
|
||||
- 从 ixformer.h 确认 `ixformer::infer::topk_softmax` 在C++层存在但Python binding缺失
|
||||
- 新代码: 尝试 ixformer._C.topk_softmax → 安静 PyTorch fallback
|
||||
|
||||
3. **_custom_ops.py 部署**: patch_ops.sh 添加部署步骤
|
||||
- 之前标记为 "DO NOT deploy", 现在修复后部署
|
||||
|
||||
关键发现 (from upstream_ref/xllm):
|
||||
- xllm/core/kernels/ilu/ixformer.h: 完整的 ixformer::infer API (14函数)
|
||||
- xllm/core/layers/ilu/fused_moe.cpp: 生产级7步MoE pipeline (797行)
|
||||
- xllm/core/kernels/ilu/fused_moe.cpp: topk_softmax + gen_idx + expand + combine
|
||||
- 这些代码在 upstream_ref 中已存在, 接口与我们的 ix_full_bridge.cpp 完全一致
|
||||
|
||||
## 历史任务摘要
|
||||
- comp 168 三个致命bug修复 (OOM + topk_softmax + _custom_ops部署)
|
||||
- corex_moe/corex_gdn/corex_fa2 dlopen模块重写 (ixformer::infer dispatch chain)
|
||||
- CCCL upstream导入(5205文件) + 27/27 muh tuning headers + CCCL→vllm pattern mapping
|
||||
- ix_full_bridge.cpp 14函数桥接 + moe_topk_softmax_v3.cu
|
||||
- GDN dtype guard + NaN clamp修复
|
||||
- serving层部署(protocol/serving_chat/api_server等) + Sub508/509功能修复
|
||||
- 67 GitHub issues + 121 draft issues + PRD/SYSTEM_DESIGN文档
|
||||
|
||||
## 遗留问题/下次继续
|
||||
1. **GDN NaN (P0)** — prefill GDN 99.98% NaN, 替换为zeros=模型质量归零; 需要参考 xllm/npu_torch/qwen3_gated_delta_net_base.cpp 做 fp32 accumulation
|
||||
2. **真机编译ix_full_bridge.cpp** — JIT编译后MoE走Tier 0 (C++ 7步) 取代 Python loop
|
||||
3. **MoE性能** — 当前全走PyTorch for循环 (64 experts × 每token), Output TPS=11.86
|
||||
4. **121个draft issues→真issue** — GitHub API批量转换
|
||||
5. **提交竞赛平台** — 当前修复应能通过functional_acceptance基本测试, 不再OOM崩溃
|
||||
|
||||
@@ -1,127 +0,0 @@
|
||||
# 动态链接库完整清单与调用链
|
||||
|
||||
## 1. 已有预编译 .so（22 个）→ 调用链状态
|
||||
|
||||
### A. 已接入模型调用链（15 个）
|
||||
|
||||
| .so | 来源 | 模型中的环境变量 | 状态 |
|
||||
|-----|------|-----------------|------|
|
||||
| corex_gdn_causal_conv | 自研 CUDA | `BI100_GDN_COREX_CAUSAL_CONV` (default=True) | ✅ 代码引用 4 处 |
|
||||
| corex_gdn_gated_norm | 自研 CUDA | `BI100_GDN_COREX_GATED_NORM` (default=True) | ✅ 代码引用 4 处 |
|
||||
| corex_gdn_beta_decay | 自研 CUDA | `BI100_GDN_COREX_BETA_DECAY` (default=True) | ✅ 代码引用 4 处 |
|
||||
| corex_gdn_qk_map | 自研 CUDA | `BI100_GDN_COREX_QK_MAP` (default=True) | ✅ 代码引用 4 处 |
|
||||
| corex_gdn_packed_decode | 自研 CUDA | `BI100_GDN_COREX_PACKED_DECODE` (default=False) | ✅ yaml 已开 |
|
||||
| corex_gdn_chunk_recurrent | 自研 CUDA | 自动检测 | ✅ 代码引用 4 处 |
|
||||
| corex_attn_head_rms_norm | 自研 CUDA | `BI100_ATTN_COREX_HEAD_RMS_NORM` (default=True) | ✅ 代码引用 5 处 |
|
||||
| corex_moe_direct_routed | 自研 CUDA | `BI100_MOE_COREX_DIRECT_ROUTED` (default=False) | ✅ yaml 已开 |
|
||||
| corex_moe_exact_reduce | 自研 CUDA | `BI100_MOE_COREX_EXACT_REDUCE` (default=True) | ✅ 代码引用 4 处 |
|
||||
| corex_moe_weight_gather | 自研 CUDA | `BI100_MOE_COREX_WEIGHT_GATHER` (default=True) | ✅ 代码引用 4 处 |
|
||||
| corex_moe_topk_softmax | 自研 CUDA | `BI100_MOE_COREX_TOPK_SOFTMAX` (default=True) | ✅ yaml 已开 |
|
||||
| corex_moe_index_combine | 自研 CUDA | `BI100_MOE_COREX_INDEX_COMBINE` (default=True) | ✅ 代码引用 4 处 |
|
||||
| xllm_moe | 搬自 xllm upstream | `BI100_MOE_XLLM` (default=True) | ✅ 代码引用 7 处 |
|
||||
| xllm_activation | 搬自 xllm upstream | 无直接 env | ❌ 编了但没接入 |
|
||||
| xllm_norm | 搬自 xllm upstream | 无直接 env | ❌ 编了但没接入 |
|
||||
|
||||
### B. 已编译但未接入（7 个） — 需要修复
|
||||
|
||||
| .so | 来源 | 提供的函数 | 为什么没接入 | 接入方案 |
|
||||
|-----|------|-----------|------------|---------|
|
||||
| **ix_full_bridge** | ix_full_bridge.cpp → ixformer::infer | silu_and_mul, rms_norm, fused_add_rms_norm, ix_linear, ix_linear_ex | qwen3_5.py 没有 import | patch_vllm_ops.py 已写好（最新 commit），通过 ix_startup_patch.py 自动 hook |
|
||||
| **xllm_activation** | xllm activation.cu | silu_and_mul, gelu_and_mul, act_and_mul | 与 _custom_ops→ixf_F 冗余 | 作为 backup，当 ixf_F 不可用时走 xllm kernel |
|
||||
| **xllm_norm** | xllm norm.cu | rms_norm, fused_add_rms_norm | 与 _custom_ops→ixf_F 冗余 | 同上 |
|
||||
| **xllm_rope** | xllm rope.cu | rotary_embedding | 与 _custom_ops→ixf_F 冗余 | 同上 |
|
||||
| **xllm_cache** | xllm reshape_paged_cache.cu | reshape_paged_cache | paged_attn.py 没有调用 | 需要在 cache 写入路径接入 |
|
||||
| **corex_fused_paged_prefill** | 自研 CUDA | fused prefill attention | paged_attn.py 有代码但 env 没开 | computility-run.yaml 加 `BI100_ATTN_COREX_FUSED_PAGED_PREFILL=1` |
|
||||
| **corex_paged_kv_gather** | 自研 CUDA | paged KV gather | paged_attn.py 有代码但 env 没开 | 同上 |
|
||||
| **corex_block_major_kv_transfer** | 自研 CUDA | block-major KV copy | 完全没有调用点 | 需要在 worker/cache_engine 接入 |
|
||||
|
||||
## 2. 需要从 upstream 搬过来编译的代码
|
||||
|
||||
### 来源: upstream_ref/xllm/xllm/core/kernels/cuda/
|
||||
|
||||
| 文件 | 功能 | 对应 .so | 优先级 |
|
||||
|------|------|---------|--------|
|
||||
| xattention/decoder_reshape_and_cache.cu | fused KV cache write | xllm_xattn_cache | P0 |
|
||||
| xattention/prefill_reshape_and_cache.cu | prefill cache write | xllm_xattn_cache | P0 |
|
||||
| xattention/cache_select.cu | cache select | xllm_xattn_cache | P1 |
|
||||
| xattention/lse_combine.cu | LSE combine | xllm_xattn_cache | P1 |
|
||||
| fused_qknorm_rope.cu | fused QK norm + RoPE | xllm_fused_qknorm_rope | P0（每层省 4 kernel launch） |
|
||||
| matmul.cpp | ixformer GEMM wrapper | 已在 ilu/matmul.cpp | ✅ 已搬 |
|
||||
| fp8_quant.cu | FP8 quantization | xllm_fp8 | P2 |
|
||||
|
||||
### 来源: upstream_ref/xllm/xllm/core/kernels/ilu/
|
||||
|
||||
**全部已搬到 ex_engine/xllm_kernels/ilu/**（对比确认只差 CMakeLists.txt）
|
||||
|
||||
### 来源: upstream_ref/ds_vllm/csrc/libtorch_stable/
|
||||
|
||||
| 文件 | 功能 | 可用性 |
|
||||
|------|------|--------|
|
||||
| attention/paged_attention_v1.cu | paged attention v1 | SM70 兼容，但依赖 vllm C++ build |
|
||||
| attention/paged_attention_v2.cu | paged attention v2 | 同上 |
|
||||
| layernorm_kernels.cu | RMSNorm kernel | SM70 兼容 |
|
||||
| activation_kernels.cu | SiLU kernel | SM70 兼容 |
|
||||
| pos_encoding_kernels.cu | RoPE kernel | SM70 兼容 |
|
||||
| moe/topk_softmax_kernels.cu | topk+softmax fused | SM70 兼容 |
|
||||
| moe/moe_align_sum_kernels.cu | MoE align+sum | SM70 兼容 |
|
||||
|
||||
## 3. ixformer::infer 可用 API（base 镜像已有）
|
||||
|
||||
来自 `upstream_ref/xllm_latest/core/kernels/ilu/ixformer.h`:
|
||||
|
||||
```
|
||||
ixformer::infer::silu_and_mul(input, output)
|
||||
ixformer::infer::rms_norm(input, weight, output, bias, eps)
|
||||
ixformer::infer::residual_rms_norm(input, residual, weight, output, residual_out, bias, alpha, eps, is_post)
|
||||
ixformer::infer::ixformer_linear(input, weight, act_type, bias, out, persistent)
|
||||
ixformer::infer::ixformer_linear_ex(input, weight, bias, out)
|
||||
ixformer::infer::xllm_rotary_embedding(positions, query, key, head_size, cos_sin_cache, is_neox)
|
||||
ixformer::infer::xllm_reshape_and_cache(key, value, key_cache, value_cache, slot_mapping, key_stride, value_stride)
|
||||
ixformer::infer::xllm_paged_attention(out, query, key_cache, value_cache, ...)
|
||||
ixformer::infer::ixinfer_flash_attn_unpad_with_block_tables(query, key_cache, value_cache, ...)
|
||||
ixformer::infer::topk_softmax(weights, indices, token_expert_indices, gating_output, renormalize)
|
||||
ixformer::infer::moe_compute_token_index_api(topk_ids, src_dst, dst_src, expert_sizes, ...)
|
||||
ixformer::infer::moe_expand_input(output, input, dst_to_src, src_to_dst, dst_tokens, expand_factor)
|
||||
ixformer::infer::moe_w16a16_group_gemm(output, input, weights, tokens_per_experts, ...)
|
||||
ixformer::infer::moe_output_reduce_sum(output, input, weight, mask, extra_residual, scaling)
|
||||
```
|
||||
|
||||
这些函数通过 `ix_full_bridge.so` pybind11 暴露给 Python 侧。
|
||||
|
||||
## 4. 调用链完整性检查
|
||||
|
||||
### 当前断裂点:
|
||||
|
||||
1. **ix_full_bridge.so 的 group_gemm → MoE Python for-loop**
|
||||
- `ixformer::infer::moe_w16a16_group_gemm` 在 ix_full_bridge.so 中可用
|
||||
- 但 qwen3_5.py MoE prefill 路径 (L1813-1825) 还是 `F.linear` per-expert loop
|
||||
- 需要: ix_fused_moe.py 的 7 步 pipeline 走 group_gemm 而非 per-expert linear
|
||||
|
||||
2. **corex_fused_paged_prefill → paged_attn.py env 没开**
|
||||
- .so 已编译已部署
|
||||
- paged_attn.py 已有完整调用代码 (L2030)
|
||||
- computility-run.yaml 缺少 `BI100_ATTN_COREX_FUSED_PAGED_PREFILL=1`
|
||||
|
||||
3. **xllm_cache → reshape_and_cache 没接入**
|
||||
- base 镜像 ixformer 已有 `xllm_reshape_and_cache`
|
||||
- vllm 的 cache_ops 走的是另一条路径
|
||||
|
||||
## 5. 需要编出的新 .so
|
||||
|
||||
| 目标 .so | 源文件 | 编译方式 | 依赖 |
|
||||
|---------|--------|---------|------|
|
||||
| xllm_fused_qknorm_rope.so | upstream fused_qknorm_rope.cu + bind | corex clang --cuda-gpu-arch=ivcore10 | libcudart, torch |
|
||||
| xllm_xattn_cache.so | upstream xattention/*.cu + bind | 同上 | 同上 |
|
||||
|
||||
## 6. computility-run.yaml 需要补全的 env
|
||||
|
||||
```yaml
|
||||
- name: BI100_ATTN_COREX_FUSED_PAGED_PREFILL
|
||||
value: '1'
|
||||
- name: BI100_ATTN_COREX_PAGED_KV_GATHER
|
||||
value: '1'
|
||||
- name: IX_OPS_AUTO_PATCH
|
||||
value: '1'
|
||||
- name: PYTORCH_CUDA_ALLOC_CONF
|
||||
value: 'expandable_segments:True'
|
||||
```
|
||||
@@ -1,141 +0,0 @@
|
||||
# Sub509 深度诊断 — 基于CCCL源码阅读的系统级分析
|
||||
|
||||
## 一、Sub509 vs Sub168 关键数据对比
|
||||
|
||||
| 测试 | 对手Sub168 | 我们Sub509 | 差距分析 |
|
||||
|------|-----------|-----------|---------|
|
||||
| d01_basic_nostream | 8.49s, content[11] tok=139 | 95.85s, content[0] reasoning[1102] tok=1085 | 11x慢; 我们产了1085个token全是reasoning |
|
||||
| d02_stream_usage | 2.75s, chunks=53 | 1.84s, chunks=9 | 我们居然更快(但只产了9个chunks vs 53) |
|
||||
| d03_tool_call | 2.12s, tool=get_weather | **49.04s, tools=0 finish=stop** | **致命**: 模型不输出<tool_call> XML |
|
||||
| d04_reasoning | 17.78s, content[181] reasoning[1011] | 128.74s, content[0] reasoning[1447] | 7x慢; 我们有reasoning但没有content |
|
||||
|
||||
## 二、三大根因(按严重程度排序)
|
||||
|
||||
### 根因1: GatedDeltaNet每层产NaN → 模型"智力"丧失
|
||||
|
||||
docker日志证据:
|
||||
```
|
||||
WARNING qwen3_5.py:445] NaN in prefill GatedDeltaNet layer 0 (frac=0.9998)
|
||||
WARNING qwen3_5.py:445] NaN in prefill GatedDeltaNet layer 1 (frac=0.9997)
|
||||
WARNING qwen3_5.py:445] NaN in prefill GatedDeltaNet layer 2 (frac=1.0000)
|
||||
WARNING qwen3_5.py:445] NaN in prefill GatedDeltaNet layer 4 (frac=1.0000)
|
||||
```
|
||||
|
||||
**99.98%-100% NaN率**。`nan_to_num(result, nan=0.0)` 将这些NaN替换为零，等于整个DeltaNet层输出全是零。
|
||||
这是一种"活着但脑死亡"的状态——前向传播不报错，但模型失去了DeltaNet层的能力。
|
||||
|
||||
**NaN来源追踪**:
|
||||
1. `_torch_chunk_gated_delta_rule` 中 `g.cumsum(dim=-1)` → 累积值可能极大
|
||||
2. `g.clamp(-20,20)` 后 `g.exp()` → 最大 ~5e8，但这些值进入矩阵乘法后仍可能溢出
|
||||
3. `decay_mask = (g_diff).tril().exp()` → 即使单个exp不溢出，大矩阵乘法的累加也可能溢出
|
||||
4. `_forward_sub_lower` 中的前向替代: `x[i] = rhs[i] + A[i,:i] @ x[:i]`，如果A中有大值，误差逐行放大
|
||||
|
||||
**对手为什么没有这个问题**: 对手可能用的是不同的模型架构(不含DeltaNet)，或者在NVIDIA GPU上float32精度够高不会溢出。
|
||||
|
||||
### 根因2: FusedMoE完全fallback → 性能灾难
|
||||
|
||||
```
|
||||
ERROR _custom_ops.py:58] module 'ixformer.functions' has no attribute 'vllm_moe_topk_softmax'
|
||||
WARNING qwen3_5.py:913] FusedMoE native kernel failed, falling back to pure PyTorch experts permanently.
|
||||
```
|
||||
|
||||
BI-V100的ixformer没有MoE kernel，所有MoE层都用纯PyTorch:
|
||||
- 256个expert × top_k=8 → 最多256次F.linear调用(prefill)
|
||||
- 每次decode也需要top_k=8次expert forward
|
||||
- 对比native kernel的1次fused launch，这是数量级的差距
|
||||
|
||||
### 根因3: computility-run.yaml vs 实际参数不一致
|
||||
|
||||
yaml写的: `--max-model-len 256000 --max-num-seqs 2 --gpu-memory-utilization 0.95`
|
||||
docker日志: `max_seq_len=100000, max_num_seqs=1, gpu_memory_utilization=0.9`
|
||||
|
||||
**可能原因**: 部署时还在用旧的配置。需要确认yaml是否真的被用于部署。
|
||||
|
||||
## 三、d03_tool_call为什么FAIL
|
||||
|
||||
d03日志: `tools=0 finish=stop reasoning[0] (tool_choice=auto) (49.04s)`
|
||||
|
||||
**reasoning[0]说明enable_thinking=False确实生效了**。但模型仍然不输出`<tool_call>` XML。
|
||||
|
||||
analysis:
|
||||
1. enable_thinking=False → 模型不产生`<think>...</think>`块 ✓
|
||||
2. 但模型的输出内容不包含`<tool_call><function=get_weather>...` 格式
|
||||
3. tool parser `Qwen3CoderToolParser` 在输出中找不到 `<function=` → tools_called=False
|
||||
4. 49.04s意味着模型在漫长生成纯文本回答(可能是口头描述天气而不是调用tool)
|
||||
|
||||
**核心问题: GatedDeltaNet的NaN导致模型质量太差,不能正确follow tool_call格式**
|
||||
|
||||
这不是serving_chat.py的问题。serving_chat.py和protocol.py中的tool_call thinking禁用逻辑是正确的。问题在模型本身。
|
||||
|
||||
## 四、对手Sub168分析
|
||||
|
||||
对手最终得分60194.6:
|
||||
- functional: ~48/52 PASS
|
||||
- case_truncation: score=1.0
|
||||
- replay_tencent: score=60194 (94/881成功, tps avg 11.86)
|
||||
- opencompass: 0.0 (server也崩了)
|
||||
|
||||
**对手server在replay后期也崩溃了**(704个connection refused)。
|
||||
**对手的replay也只有94/881成功(10.7%)**。
|
||||
|
||||
但对手赢在:
|
||||
1. functional高通过率 → 基础分
|
||||
2. case_truncation通过 → 引擎稳定
|
||||
3. replay中94个成功请求 × tps → 得到分数
|
||||
|
||||
## 五、修复路径(按投入产出比排序)
|
||||
|
||||
### 修复1: NaN问题 — 强制float32精度 + 更激进的clamp
|
||||
|
||||
当前: `g.clamp(-20, 20)` 不够。cumsum后再clamp太晚了。
|
||||
需要: 在cumsum之前就对g的原始值做clamp。
|
||||
|
||||
在 `_torch_chunk_gated_delta_rule`:
|
||||
```python
|
||||
# 现在: g = g.cumsum(dim=-1).clamp(-20, 20)
|
||||
# 改为: g = g.clamp(-5, 5).cumsum(dim=-1).clamp(-15, 15)
|
||||
```
|
||||
|
||||
在 GatedDeltaNet.forward 的 prefill path:
|
||||
```python
|
||||
# 现在: _A_safe = self.A_log.float().clamp(-20.0, 20.0)
|
||||
# 改为更窄: _A_safe = self.A_log.float().clamp(-10.0, 10.0)
|
||||
```
|
||||
|
||||
### 修复2: MoE性能 — 尝试真正使用native kernel
|
||||
|
||||
docker日志说 `vllm_moe_topk_softmax` 不存在，但 _custom_ops.py 里应该有PyTorch fallback。
|
||||
问题是 `_hw_policy.moe_native_align` 和 `_hw_policy.moe_native_invoke` 也是False。
|
||||
如果这两个真的不存在，那PyTorch fallback就是唯一选择。
|
||||
|
||||
**性能改进**: 在 `_pure_pytorch_experts` 的 prefill path 中:
|
||||
- 现在: for-loop over experts, 每个一次F.linear
|
||||
- 改为: 按expert batch size排序，大batch的expert合并成一个大F.linear (CCCL histogram pattern已经实现了，但可以更激进)
|
||||
|
||||
### 修复3: computility-run.yaml 参数对齐
|
||||
|
||||
确保部署时真的用了yaml里的参数。max-num-seqs=2让n=2请求不会崩溃。
|
||||
|
||||
### 修复4: d01/d04速度
|
||||
|
||||
d01: 95.85s产了1085个token，约11.3 tok/s — 其实tps不太差
|
||||
对手d01: 8.49s产了139个token，约16.4 tok/s
|
||||
|
||||
**关键差异不是tps，是产了多少token!** 我们1085 vs 对手139。
|
||||
我们的模型在thinking里产了大量token。
|
||||
d01是basic_nostream，没有tool，所以enable_thinking=True是默认的。
|
||||
thinking产了1102个reasoning token + 0个content token。
|
||||
|
||||
**问题: d01测试的content[0]意味着没有实际内容输出!**
|
||||
对手content[11]说明他输出了内容。
|
||||
|
||||
这又回到了GatedDeltaNet NaN → 模型质量差的问题。
|
||||
|
||||
## 六、CCCL启示
|
||||
|
||||
CCCL在处理数值稳定性方面的核心设计:
|
||||
1. `overflow_cast_t<T>` — 在可能溢出的地方用更高精度的中间类型
|
||||
2. `cc_dispatch` — 不同硬件不同策略，不硬编码
|
||||
3. `policy_selector` — 基于benchmark数据选择参数，不拍脑袋
|
||||
|
||||
我们的DeltaNet实现缺少CCCL级别的数值稳定性保证。
|
||||
@@ -1,48 +0,0 @@
|
||||
# Sub508/509 完整诊断报告
|
||||
|
||||
## 修复提交记录
|
||||
|
||||
| Commit | 修复 | 影响 |
|
||||
|--------|------|------|
|
||||
| e0344b1 | 禁用 tool_call 请求的 thinking | d03 FAIL → 预计 PASS |
|
||||
| c241764 | get_scheduler_config try-catch | 防止引擎崩溃 |
|
||||
| 994c657 | clamp n>1 to 1 | 防止 t2_n_2 级联崩溃 (19 个测试) |
|
||||
|
||||
## Sub508 完整测试结果 (56 tests)
|
||||
|
||||
### 实际结果: PASS=21, FAIL=30, SKIP=5
|
||||
|
||||
### 级联崩溃 (19 个 FAIL 来自 t2_n_2 引擎崩溃)
|
||||
t2_n_2 → HTTP 500 → 引擎死亡 → t3_max_tokens_none/1/64/mid/max/neg1/over,
|
||||
t4a/4b, t5, t6, t7, t8, t9, t10, t12_chinese/japanese/emoji 全部 HTTP 500
|
||||
|
||||
### 修复后预期: PASS ≈ 40+, FAIL ≈ 10-
|
||||
|
||||
### 真正的功能性 FAIL (非级联)
|
||||
|
||||
| 测试 | 状态 | 根因 | 可修 |
|
||||
|------|------|------|------|
|
||||
| d03_tool_call | tools=0 finish=stop | ✅ 已修复 thinking budget | 是 |
|
||||
| d05_multimodal | HTTP 400 | multimodal 请求格式 | 需查 |
|
||||
| d07_reasoning+content | content[0] | 模型 think 后不产 content | 否(模型) |
|
||||
| d10_thinking_disable_ctk | 乱码 content | 模型质量 | 否(模型) |
|
||||
| t1a_thinking_true | reasoning[0] | 模型跳过 thinking | 否(模型) |
|
||||
| t1c_thinking_default | reasoning[0] | 同上 | 否(模型) |
|
||||
| t2_n_2 | HTTP 500 → cascade | ✅ 已修复 clamp n | 是(防崩) |
|
||||
|
||||
## 对手 Sub168 对比
|
||||
|
||||
| 维度 | 对手 | 我们 |
|
||||
|------|------|------|
|
||||
| functional PASS | ~50/56 | 21/56 → 修后 ~40/56 |
|
||||
| d01 速度 | 8.49s | 95.87s |
|
||||
| d04 速度 | 17.78s | 129.19s |
|
||||
| replay max_completion_tokens | ✗ 400 rejected (30+次) | ✓ 已支持 (extra=ignore) |
|
||||
| replay tool_calls content=None | ✗ 400 rejected | ✓ 已支持 (normalize) |
|
||||
| decode TPS | ~16 tok/s | ~11 tok/s |
|
||||
|
||||
## 我们 vs 对手的优势
|
||||
1. `max_completion_tokens` 支持 — 对手 replay 有 30+ 个 400 错误
|
||||
2. `tool_calls` content=None 支持 — 对手 replay preflight 失败
|
||||
3. `reasoning_effort` 字段容忍 — 对手被拒
|
||||
4. prefix caching 工作 (d06 PASS) — 对手 d06 FAIL
|
||||
Reference in New Issue
Block a user