From 882669b0bc726e07cf5a1a080a5004c713519c09 Mon Sep 17 00:00:00 2001 From: root Date: Thu, 20 Aug 2026 02:28:32 +0000 Subject: [PATCH] delete old md --- CCCL_INTEGRATION_STATUS.md | 96 ------------- CCCL_MUH_GAP_ANALYSIS.md | 120 ---------------- CCCL_MUH_PARITY_AUDIT.md | 134 ----------------- CCCL_PATTERN_MAP.md | 104 -------------- CCCL_TUNING_GAP_REPORT.md | 89 ------------ CODEPATH_MAP.md | 193 ------------------------- COMP168_DIAGNOSIS.md | 165 --------------------- COMPETITIVE_ANALYSIS_AND_FIX_PLAN.md | 121 ---------------- ENGINEX_INJECTION_MAP.md | 94 ------------ ENGINE_CODEPATH_TIMELINE.md | 150 ------------------- GROUND_TRUTH_STATUS.md | 124 ---------------- GROUND_TRUTH_STATUS_v2.md | 101 ------------- MOE_SYMBOL_TRUTH.md | 68 --------- MUH_PROJECT_CHECKPOINT.md | 208 --------------------------- MUH_TUNING_GAP_ANALYSIS.md | 85 ----------- PIPELINE_GROUND_TRUTH.md | 50 ------- PIPELINE_REALITY_CHECK.md | 71 --------- PIPELINE_STATUS.md | 86 ----------- PRD.md | 6 + PROJECT_SUMMARY.md | 116 --------------- SO_BUILD_MANIFEST.md | 127 ---------------- SUB509_DEEP_DIAGNOSIS.md | 141 ------------------ SUB509_DIAGNOSIS.md | 48 ------- 23 files changed, 6 insertions(+), 2491 deletions(-) delete mode 100644 CCCL_MUH_GAP_ANALYSIS.md delete mode 100644 CCCL_MUH_PARITY_AUDIT.md delete mode 100644 CCCL_PATTERN_MAP.md delete mode 100644 CCCL_TUNING_GAP_REPORT.md delete mode 100644 CODEPATH_MAP.md delete mode 100644 COMP168_DIAGNOSIS.md delete mode 100644 COMPETITIVE_ANALYSIS_AND_FIX_PLAN.md delete mode 100644 ENGINEX_INJECTION_MAP.md delete mode 100644 ENGINE_CODEPATH_TIMELINE.md delete mode 100644 GROUND_TRUTH_STATUS.md delete mode 100644 GROUND_TRUTH_STATUS_v2.md delete mode 100644 MOE_SYMBOL_TRUTH.md delete mode 100644 MUH_PROJECT_CHECKPOINT.md delete mode 100644 MUH_TUNING_GAP_ANALYSIS.md delete mode 100644 PIPELINE_GROUND_TRUTH.md delete mode 100644 PIPELINE_REALITY_CHECK.md delete mode 100644 PIPELINE_STATUS.md delete mode 100644 SO_BUILD_MANIFEST.md delete mode 100644 SUB509_DEEP_DIAGNOSIS.md delete mode 100644 SUB509_DIAGNOSIS.md diff --git a/CCCL_INTEGRATION_STATUS.md b/CCCL_INTEGRATION_STATUS.md index 5658b773..90963d93 100644 --- a/CCCL_INTEGRATION_STATUS.md +++ b/CCCL_INTEGRATION_STATUS.md @@ -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 硬件约束下的最优解） - -竞赛成绩 = 算法因子置换的精度 × 硬件实测标定的覆盖度。 diff --git a/CCCL_MUH_GAP_ANALYSIS.md b/CCCL_MUH_GAP_ANALYSIS.md deleted file mode 100644 index 590fbfc5..00000000 --- a/CCCL_MUH_GAP_ANALYSIS.md +++ /dev/null @@ -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 上跑竞赛功能验证 diff --git a/CCCL_MUH_PARITY_AUDIT.md b/CCCL_MUH_PARITY_AUDIT.md deleted file mode 100644 index 1186efca..00000000 --- a/CCCL_MUH_PARITY_AUDIT.md +++ /dev/null @@ -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_.tpb_.ns_.dcid_.l2w_.trp_.ld_ - 后跟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的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需要相应增大 diff --git a/CCCL_PATTERN_MAP.md b/CCCL_PATTERN_MAP.md deleted file mode 100644 index d6681176..00000000 --- a/CCCL_PATTERN_MAP.md +++ /dev/null @@ -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) -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) diff --git a/CCCL_TUNING_GAP_REPORT.md b/CCCL_TUNING_GAP_REPORT.md deleted file mode 100644 index fd374012..00000000 --- a/CCCL_TUNING_GAP_REPORT.md +++ /dev/null @@ -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% | diff --git a/CODEPATH_MAP.md b/CODEPATH_MAP.md deleted file mode 100644 index b7625f2f..00000000 --- a/CODEPATH_MAP.md +++ /dev/null @@ -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 — 模型不能输出 XML -├─ d07_reasoning_plus_content: content[0] — 模型不输出 -├─ 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%) -``` diff --git a/COMP168_DIAGNOSIS.md b/COMP168_DIAGNOSIS.md deleted file mode 100644 index 9a85b7fb..00000000 --- a/COMP168_DIAGNOSIS.md +++ /dev/null @@ -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" 的硬要求落实到代码里 diff --git a/COMPETITIVE_ANALYSIS_AND_FIX_PLAN.md b/COMPETITIVE_ANALYSIS_AND_FIX_PLAN.md deleted file mode 100644 index 296d411f..00000000 --- a/COMPETITIVE_ANALYSIS_AND_FIX_PLAN.md +++ /dev/null @@ -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 中耗尽了时间，没有产生 `` 标签 -- 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 是否正确注入了 `` 标签。 - -### 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. 提交部署，跑测试 diff --git a/ENGINEX_INJECTION_MAP.md b/ENGINEX_INJECTION_MAP.md deleted file mode 100644 index a2de7440..00000000 --- a/ENGINEX_INJECTION_MAP.md +++ /dev/null @@ -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 参数 -``` diff --git a/ENGINE_CODEPATH_TIMELINE.md b/ENGINE_CODEPATH_TIMELINE.md deleted file mode 100644 index 2e533fdf..00000000 --- a/ENGINE_CODEPATH_TIMELINE.md +++ /dev/null @@ -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()` diff --git a/GROUND_TRUTH_STATUS.md b/GROUND_TRUTH_STATUS.md deleted file mode 100644 index fdce20ca..00000000 --- a/GROUND_TRUTH_STATUS.md +++ /dev/null @@ -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) -``` diff --git a/GROUND_TRUTH_STATUS_v2.md b/GROUND_TRUTH_STATUS_v2.md deleted file mode 100644 index 1fdf7a96..00000000 --- a/GROUND_TRUTH_STATUS_v2.md +++ /dev/null @@ -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+ 用例 diff --git a/MOE_SYMBOL_TRUTH.md b/MOE_SYMBOL_TRUTH.md deleted file mode 100644 index 087e1352..00000000 --- a/MOE_SYMBOL_TRUTH.md +++ /dev/null @@ -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 -``` diff --git a/MUH_PROJECT_CHECKPOINT.md b/MUH_PROJECT_CHECKPOINT.md deleted file mode 100644 index 3995abc8..00000000 --- a/MUH_PROJECT_CHECKPOINT.md +++ /dev/null @@ -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) diff --git a/MUH_TUNING_GAP_ANALYSIS.md b/MUH_TUNING_GAP_ANALYSIS.md deleted file mode 100644 index 610827eb..00000000 --- a/MUH_TUNING_GAP_ANALYSIS.md +++ /dev/null @@ -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。 diff --git a/PIPELINE_GROUND_TRUTH.md b/PIPELINE_GROUND_TRUTH.md deleted file mode 100644 index b3ddee64..00000000 --- a/PIPELINE_GROUND_TRUTH.md +++ /dev/null @@ -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 diff --git a/PIPELINE_REALITY_CHECK.md b/PIPELINE_REALITY_CHECK.md deleted file mode 100644 index 006c37fc..00000000 --- a/PIPELINE_REALITY_CHECK.md +++ /dev/null @@ -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+) diff --git a/PIPELINE_STATUS.md b/PIPELINE_STATUS.md deleted file mode 100644 index eabf151d..00000000 --- a/PIPELINE_STATUS.md +++ /dev/null @@ -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 diff --git a/PRD.md b/PRD.md index 7d33ef4b..2a0e6415 100644 --- a/PRD.md +++ b/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 都有效。 \ No newline at end of file diff --git a/PROJECT_SUMMARY.md b/PROJECT_SUMMARY.md index 14c4c044..8c3996a7 100644 --- a/PROJECT_SUMMARY.md +++ b/PROJECT_SUMMARY.md @@ -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崩溃 diff --git a/SO_BUILD_MANIFEST.md b/SO_BUILD_MANIFEST.md deleted file mode 100644 index 246a57f8..00000000 --- a/SO_BUILD_MANIFEST.md +++ /dev/null @@ -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' -``` diff --git a/SUB509_DEEP_DIAGNOSIS.md b/SUB509_DEEP_DIAGNOSIS.md deleted file mode 100644 index 72d8a602..00000000 --- a/SUB509_DEEP_DIAGNOSIS.md +++ /dev/null @@ -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** | **致命**: 模型不输出 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确实生效了**。但模型仍然不输出`` XML。 - -analysis: -1. enable_thinking=False → 模型不产生`...`块 ✓ -2. 但模型的输出内容不包含`...` 格式 -3. tool parser `Qwen3CoderToolParser` 在输出中找不到 `` — 在可能溢出的地方用更高精度的中间类型 -2. `cc_dispatch` — 不同硬件不同策略，不硬编码 -3. `policy_selector` — 基于benchmark数据选择参数，不拍脑袋 - -我们的DeltaNet实现缺少CCCL级别的数值稳定性保证。 diff --git a/SUB509_DIAGNOSIS.md b/SUB509_DIAGNOSIS.md deleted file mode 100644 index c1aee0b9..00000000 --- a/SUB509_DIAGNOSIS.md +++ /dev/null @@ -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