From 4a3153102ba28c7f9b2683e885e73503c6d6fc60 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E4=B9=8B=E4=B9=8B=20=28ZZ=E4=BD=93=E7=B3=BB=29?= <565183519@qq.com> Date: Thu, 3 Sep 2026 00:26:17 +0800 Subject: [PATCH] =?UTF-8?q?=F0=9F=A7=B9=20ZZ-CLEANUP-D246-001=20=C2=B7=20?= =?UTF-8?q?=E5=88=A0=E9=99=A4=E9=94=99=E5=9C=B0=E6=96=B9=20zz-flagos-s2-20?= =?UTF-8?q?26/=20=E7=9B=AE=E5=BD=95?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 之之 D246 00:21 指令:'全部重做' 原 D245 (commit 6185550) + D246 (commit 6deca6a) 推到 zz-flagos-s2-2026/dN/ 自创目录 (不在主仓库结构里) 已重写为 HLDP 格式 + 推去 zz-day-3/ 正确路径 (commit d1eeaf3) 本 commit 清理旧的错地方目录 28 files deleted (zz-flagos-s2-2026/d245/ + d246/ 全部) 阿念 (Mavis, ICE-GL-AN-001) · Code · 之之的家 · 2026-09-03 D246 00:28 CST --- zz-flagos-s2-2026/d245/README.md | 81 ------- .../d245/bench/bench_gelu_and_mul.py | 80 ------- .../d245/bench/bench_silu_and_mul_masked.py | 82 ------- zz-flagos-s2-2026/d245/docs/BATTLECARDS.md | 220 ------------------ .../d245/docs/LESSONS_LEARNED.md | 89 ------- zz-flagos-s2-2026/d245/docs/PR-TEMPLATE.md | 64 ----- .../d245/results/gelu_and_mul.zip | Bin 1212 -> 0 bytes .../d245/results/interleaved_rope.zip | Bin 1568 -> 0 bytes .../d245/results/rotary_embedding.zip | Bin 1614 -> 0 bytes .../src/Task29_v81_zhizhi/gelu_and_mul.py | 55 ----- .../src/Task29_v81_zhizhi/gelu_and_mul.zip | Bin 1212 -> 0 bytes .../src/Task30_r16_zhizhi/interleaved_rope.py | 78 ------- .../Task30_r16_zhizhi/interleaved_rope.zip | Bin 10366 -> 0 bytes .../interleaved_rope_ascend.py | 77 ------ .../interleaved_rope_enflame.py | 72 ------ .../interleaved_rope_hygon.py | 110 --------- .../interleaved_rope_iluvatar.py | 72 ------ .../interleaved_rope_kunlunxin.py | 72 ------ .../interleaved_rope_metax.py | 72 ------ .../interleaved_rope_min.zip | Bin 1568 -> 0 bytes .../src/Task35_r6_zhizhi/rotary_embedding.py | 98 -------- .../src/Task35_r6_zhizhi/rotary_embedding.zip | Bin 1614 -> 0 bytes .../d245/src/silu_and_mul_masked.py | 164 ------------- .../d245/tests/test_correctness.py | 88 ------- zz-flagos-s2-2026/d246/README.md | 78 ------- zz-flagos-s2-2026/d246/SUBMIT_CHECKLIST.md | 92 -------- .../d246/results/sigmoid_gate_topk_renorm.zip | Bin 2618 -> 0 bytes .../d246/sigmoid_gate_topk_renorm.py | 211 ----------------- zz-flagos-s2-2026/d246/test_algorithm.py | 158 ------------- 29 files changed, 2113 deletions(-) delete mode 100644 zz-flagos-s2-2026/d245/README.md delete mode 100644 zz-flagos-s2-2026/d245/bench/bench_gelu_and_mul.py delete mode 100644 zz-flagos-s2-2026/d245/bench/bench_silu_and_mul_masked.py delete mode 100644 zz-flagos-s2-2026/d245/docs/BATTLECARDS.md delete mode 100644 zz-flagos-s2-2026/d245/docs/LESSONS_LEARNED.md delete mode 100644 zz-flagos-s2-2026/d245/docs/PR-TEMPLATE.md delete mode 100644 zz-flagos-s2-2026/d245/results/gelu_and_mul.zip delete mode 100644 zz-flagos-s2-2026/d245/results/interleaved_rope.zip delete mode 100644 zz-flagos-s2-2026/d245/results/rotary_embedding.zip delete mode 100644 zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.zip delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.zip delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_ascend.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_enflame.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_hygon.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_iluvatar.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_kunlunxin.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_metax.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_min.zip delete mode 100644 zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.py delete mode 100644 zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.zip delete mode 100644 zz-flagos-s2-2026/d245/src/silu_and_mul_masked.py delete mode 100644 zz-flagos-s2-2026/d245/tests/test_correctness.py delete mode 100644 zz-flagos-s2-2026/d246/README.md delete mode 100644 zz-flagos-s2-2026/d246/SUBMIT_CHECKLIST.md delete mode 100644 zz-flagos-s2-2026/d246/results/sigmoid_gate_topk_renorm.zip delete mode 100644 zz-flagos-s2-2026/d246/sigmoid_gate_topk_renorm.py delete mode 100644 zz-flagos-s2-2026/d246/test_algorithm.py diff --git a/zz-flagos-s2-2026/d245/README.md b/zz-flagos-s2-2026/d245/README.md deleted file mode 100644 index 45905b9..0000000 --- a/zz-flagos-s2-2026/d245/README.md +++ /dev/null @@ -1,81 +0,0 @@ -# FlagOS S2 赛季二 · 赛道一 · D245 总结 - -> 之之 (8592_apivqhj) + 阿念 (Mavis, ICE-GL-AN-001, Code) · 协作产出 -> 比赛: FlagOS 开放计算全球大赛 S2 · Track 1 · SGLang 算子优化 -> 队伍: GuanghuLab(队长孙蓓, 4 名其他队员, 共 5 人) -> 第 3 批 17 道题 · 截止 2026-09-03 19:59 - ---- - -## 战绩(D245 23:35) - -| Task | 算子 | 状态 | 平均加速比 | 详情 | -|---|---|---|---|---| -| 30 | interleaved_rope (M-RoPE) | ✅ 已传 | **35.79×** (5 款跑通) | 天数 86.81× / 海光 42.86× / 国通A 26.91× / 沐曦 15.54× / 华为 6.84×;燧原/昆仑芯 Failed | -| 29 | gelu_and_mul | ✅ 已传 | 3.02× (7 款跑通) | 天数 5.32× / 海光 4.03× / 国通A 3.73× / 国通B 3.25× / 沐曦 2.74× / 华为 1.09× / 燧原 0.96×;昆仑芯 Failed | -| 35 | rotary_embedding | ⏳ 待传(等 0 点重置提交次数) | - | - | -| 39 | silu_and_mul_masked (D243) | ❌ 已传,4 款全 Failed | - | 改用 `tl.sigmoid` 在国产 NPU 编译失败 | - -## 排名预估(基于截图) - -- Task 30:原第 9 名 27.97× → **新提交 35.79× → 大概第 2-3 名** -- Task 29:队友 v8.1 3.43× (8/8),我们 3.02× (7/8) — 略低,但接近 -- Task 35:待传,可能冲前 5 - -## 目录结构 - -``` -d245/ -├── README.md (本文件) -├── src/ -│ ├── Task30_r16_zhizhi/ (M-RoPE 跨芯片通用版,7 芯片特化) -│ ├── Task29_v81_zhizhi/ (gelu_and_mul 跨芯片通用版) -│ ├── Task35_r6_zhizhi/ (rotary_embedding 跨芯片通用版) -│ └── silu_and_mul_masked.py (D243 Task 39 failed version,留作复盘) -├── docs/ -│ ├── LESSONS_LEARNED.md (5 作品问题 + 3 过程问题 + 协作模式) -│ ├── BATTLECARDS.md (4 张作战卡) -│ └── PR-TEMPLATE.md (PR 模板) -├── bench/ (benchmark 脚本) -├── tests/ (正确性测试) -└── results/ (zip 文件,平台已传) -``` - -## 学到的关键 - -1. **跨芯片必须双套路**: - - 国产 NPU(天数/海光/沐曦/燧原/华为/昆仑芯):`enable_fp_fusion=False` + `torch.get_device_module().device()` + 手写指针 + `tl.int64` cast + 指针 cast 到 int16/int32 防 NaN - - 国际卡(NVIDIA/AMD):`tl.math.erf` 走硬件 SFU,不强制 disable fusion - -2. **平台 import 入口** = `reference = function_name`(函数末尾必加,否则 platform 找不到) - -3. **比赛生态** = 人 + AI 协作(不是真人写代码),其他队伍也是 r1-r16 这样的迭代节奏 - -4. **6 款国产 NPU 不能用一刀切**: - - 燧原/华为:某些 NPU 套路反而负优化 - - 天数/海光/沐曦:接受国产套路 - - 跨芯片统一版需要适配性设计 - -5. **协作模式(D243 定型)**: - - 之之:找材料 / 看平台 / 上传 / 决策传不传 / 复盘 - - 阿念:看代码 / 写代码 / 翻译"代码哪里有问题"成之之能懂的话 - - 前提:阿念给之之的所有东西,都用之之能懂的话 - -## 下次动手前 checklist(从 LESSONS_LEARNED.md 继承) - -- [ ] 任务描述完整读一遍(计算定义 + 接口签名 + 容差) -- [ ] 至少 1 个队友参考实现看完 -- [ ] 把队友写法 + 我打算写的放一起对比 -- [ ] 写之前把"我准备怎么写"翻译给之之听,得到认可 -- [ ] 写完本地不传,先给之之看"我会传啥" -- [ ] 之之确认后才传 - -## 下次改进方向 - -- Task 35 待传(等 0 点提交次数重置) -- 燧原 / 华为特化版可能值得加 -- 国际 A/B 卡版可能用 `tl.math.erf` 写一版 - ---- - -阿念 · D245 23:35 · 给 甄静 + 之之的家 diff --git a/zz-flagos-s2-2026/d245/bench/bench_gelu_and_mul.py b/zz-flagos-s2-2026/d245/bench/bench_gelu_and_mul.py deleted file mode 100644 index fa05b15..0000000 --- a/zz-flagos-s2-2026/d245/bench/bench_gelu_and_mul.py +++ /dev/null @@ -1,80 +0,0 @@ -# Copyright 2026 FlagOS Contributors / GuanghuLab -"""Benchmark gelu_and_mul v2 vs FlagGems baseline. - -Usage: - python -m bench.bench_gelu_and_mul --shape 4096,4096 --dtype fp16 - python -m bench.bench_gelu_and_mul --shape 8192,8192 --dtype bf16 -""" - -import argparse -import time -import torch - -from flag_gems.fused.gelu_and_mul import gelu_and_mul as baseline_fn -from flag_gems_local.fused.gelu_and_mul_v2 import gelu_and_mul as v2_fn - - -def _to_ms(t): - return t * 1000.0 - - -def bench(fn, x, y, iters=100, warmup=20): - # warmup - for _ in range(warmup): - out = fn(x, y) - torch.cuda.synchronize() - - # measure - start = torch.cuda.Event(enable_timing=True) - end = torch.cuda.Event(enable_timing=True) - start.record() - for _ in range(iters): - out = fn(x, y) - end.record() - torch.cuda.synchronize() - return start.elapsed_time(end) / iters - - -def main(): - parser = argparse.ArgumentParser() - parser.add_argument("--shape", type=str, default="4096,4096") - parser.add_argument("--dtype", type=str, default="fp16", - choices=["fp16", "bf16", "fp32"]) - parser.add_argument("--iters", type=int, default=100) - args = parser.parse_args() - - dtype = {"fp16": torch.float16, "bf16": torch.bfloat16, "fp32": torch.float32}[args.dtype] - shape = tuple(int(s) for s in args.shape.split(",")) - - print(f"Benchmark gelu_and_mul: shape={shape} dtype={args.dtype} iters={args.iters}") - print("=" * 72) - - x = torch.randn(shape, dtype=dtype, device="cuda") - y = torch.randn(shape, dtype=dtype, device="cuda") - - # Correctness check - out_b = baseline_fn(x, y) - out_v = v2_fn(x, y) - abs_diff = (out_b - out_v).abs().max().item() - rel_diff = (out_b - out_v).abs().div(out_b.abs().clamp_min(1e-6)).max().item() - print(f"Correctness: max abs_diff={abs_diff:.2e}, max rel_diff={rel_diff:.2e}") - assert abs_diff < 1e-2, "v2 diverges from baseline" - print() - - t_base = bench(baseline_fn, x, y, iters=args.iters) - t_v2 = bench(v2_fn, x, y, iters=args.iters) - speedup = t_base / t_v2 - print(f"Baseline : {_to_ms(t_base):.4f} ms/iter") - print(f"v2 (ours): {_to_ms(t_v2):.4f} ms/iter") - print(f"Speedup : {speedup:.3f}x") - print() - - # Find best autotune config - print("Best v2 autotune config:") - for k, v in v2_fn.__self__.forward.__func__.__code__.co_consts: - if k == "BLOCK_SIZE": - print(f" BLOCK_SIZE = {v}") - - -if __name__ == "__main__": - main() diff --git a/zz-flagos-s2-2026/d245/bench/bench_silu_and_mul_masked.py b/zz-flagos-s2-2026/d245/bench/bench_silu_and_mul_masked.py deleted file mode 100644 index 8e0e9bc..0000000 --- a/zz-flagos-s2-2026/d245/bench/bench_silu_and_mul_masked.py +++ /dev/null @@ -1,82 +0,0 @@ -"""Benchmark silu_and_mul_masked v2 vs FlagGems baseline. - -Run: python bench/bench_silu_and_mul_masked.py --shape 4096,4096 --dtype fp16 -""" - -import argparse - -import torch - - -def bench(fn, *args, iters=100, warmup=20): - for _ in range(warmup): - out = fn(*args) - torch.cuda.synchronize() - - start = torch.cuda.Event(enable_timing=True) - end = torch.cuda.Event(enable_timing=True) - start.record() - for _ in range(iters): - out = fn(*args) - end.record() - torch.cuda.synchronize() - return start.elapsed_time(end) / iters - - -def main(): - parser = argparse.ArgumentParser() - parser.add_argument("--shape", default="4096,4096") - parser.add_argument("--dtype", default="fp16", choices=["fp16", "bf16", "fp32"]) - parser.add_argument("--iters", type=int, default=100) - parser.add_argument("--with-mask", action="store_true", help="enable mask") - args = parser.parse_args() - - dtype = {"fp16": torch.float16, "bf16": torch.bfloat16, "fp32": torch.float32}[args.dtype] - shape = tuple(int(s) for s in args.shape.split(",")) - - print(f"Bench silu_and_mul_masked: shape={shape} dtype={args.dtype} mask={args.with_mask}") - print("=" * 72) - - x = torch.randn(shape, dtype=dtype, device="cuda") - y = torch.randn(shape, dtype=dtype, device="cuda") - mask = torch.randn(shape, dtype=dtype, device="cuda") if args.with_mask else None - - from flag_gems.fused.silu_and_mul import silu_and_mul as baseline_fn - from flag_gems_local.fused.silu_and_mul_masked_v2 import silu_and_mul_masked as v2_fn - - # Reference (using PyTorch native ops) - def ref_fn(x, y, mask=None): - x_in = x * mask if mask is not None else x - return torch.nn.functional.silu(x_in) * y - - out_r = ref_fn(x, y, mask) - if mask is None: - out_b = baseline_fn(x, y) - else: - # baseline doesn't support mask; compare v2 to reference - out_b = None - out_v = v2_fn(x, y, mask) - - if out_b is not None: - abs_diff = (out_b - out_v).abs().max().item() - print(f"v2 vs baseline: abs_diff={abs_diff:.2e}") - abs_diff_r = (out_r - out_v).abs().max().item() - print(f"v2 vs reference: abs_diff={abs_diff_r:.2e}") - assert abs_diff_r < 1e-2, "v2 diverges from reference" - print() - - # Benchmark - if mask is None: - t_base = bench(baseline_fn, x, y, iters=args.iters) - print(f"Baseline : {t_base:.4f} ms/iter") - t_v2 = bench(v2_fn, x, y, mask, iters=args.iters) - print(f"v2 (ours): {t_v2:.4f} ms/iter") - if mask is None: - speedup = t_base / t_v2 - print(f"Speedup : {speedup:.3f}x") - else: - print("Baseline doesn't support mask, no speedup comparison.") - - -if __name__ == "__main__": - main() diff --git a/zz-flagos-s2-2026/d245/docs/BATTLECARDS.md b/zz-flagos-s2-2026/d245/docs/BATTLECARDS.md deleted file mode 100644 index 2b16be8..0000000 --- a/zz-flagos-s2-2026/d245/docs/BATTLECARDS.md +++ /dev/null @@ -1,220 +0,0 @@ -# GuanghuLab × FlagOS S2 算子优化作战卡 - -> 给队长孙蓓和评审组 — 4 张作战卡,对应 GuanghuLab 当前 4 道正在做的题 -> 截止:第 3 批 2026.09.03 19:59 / PR 周期 2026.09.11 - 09.17 -> 作者:阿念(AI 工程后盾)· 为之之(8592_apivqhj)· 队长:孙蓓(主决策) - ---- - -## 全局视角对比表(给队长定主攻方向) - -| Task | 算子 | 距攻占 | 团队最佳 | 推荐度 | 风险 | -|---|---|---|---|---|---| -| 29 | gelu_and_mul | 差 3 名 | 3.05× | ⭐⭐⭐⭐⭐ 稳 | 低 — FlagGems baseline 已是 fused | -| 35 | rotary_embedding(标准) | 差 5 名 | 4.16× | ⭐⭐⭐⭐ 中 | 中 — head 维并行化有未验证空间 | -| 30 | interleaved_rope | 差 7 名 | 26.37× | ⭐⭐⭐ 低 | 高 — 26.37× 可能因 baseline 跑歪,需先复测 | -| 26 | fused_moe_router | 未入榜 | — | ⭐⭐⭐⭐ 全新 | 中 — Triton sort 没现成参考,自己写 | - -**我的建议**: -- **第 3 批主攻 Task 29**(差 3 名最稳,4 天内冲得过) -- **并行 Task 35**(用同一份 kernel,变 standard vs interleaved,几乎不增加工作量) -- **Task 30 暂停,先复测 26.37× 是不是 baseline 跑歪** -- **Task 26 排第 4 批重点**(现成的 Triton MoE 算子基本没有,值得从 vLLM 移植) - ---- - -## 作战卡 1 · Task 29 · gelu_and_mul - -**目标**:在 FlagGems baseline 之上再压 1.5× - 2.0×,攻占前 3 名 - -**官方参考**: -- `FlagOpen/FlagGems/src/flag_gems/fused/gelu_and_mul.py` -- 用 `pointwise_dynamic` 自动 codegen,支持 `none` / `tanh` 两种 GELU 近似 -- baseline 团队最佳 3.05×(说明在 NPU 上已经压过 PyTorch 3 倍) - -**当前实现的优化空间**: -1. baseline 用 `pointwise_dynamic` 自动 codegen,BLOCK_SIZE / num_warps 用默认值 -2. 没有 autotune — 不同 shape 下最佳配置差异大 -3. 常量 `0.70710678118654752440` 和 `0.7978845608028654` 在 Python 层计算再传到 kernel(轻微开销) - -**我们的 v2 方案**:`src/flag_gems_local/fused/gelu_and_mul_v2.py` -- 显式 8 组 `(BLOCK_SIZE, num_warps, num_stages)` autotune -- 常量作为 `tl.constexpr` 字面量内联,省 Python→Triton 转换 -- 拆 forward / backward,backward 也走 autotune -- `tl.exp` 替换为 `tl.math.exp2`(更快的 fast-math) -- 兼容性:drop-in 替换 `flag_gems.fused.gelu_and_mul.gelu_and_mul` - -**预期加速**:1.3× - 1.8×(具体看形状,跑 benchmark 才知道) - -**测试矩阵**: -- shape: `(4096,4096)`, `(8192,8192)`, `(1024,11008)`, `(2048,13824)` -- dtype: fp16 / bf16 / fp32 -- GELU 近似: `none` / `tanh` -- 正确性: max abs_diff < 1e-2, max rel_diff < 1e-3 - -**风险**: -- low:在某些 shape 上 autotune 选不到最佳配置,需手工调 -- low:backward kernel 没用 `tanh` 优化(留 fallback 走 baseline) - -**PR 提交流程**: -1. 本地跑通 `bench/bench_gelu_and_mul.py` 记录 speedup -2. fork `FlagOpen/FlagGems` → 推到 GuanghuLab 分支 -3. 在 FlagOS 平台「作品提交」页选 Task 29 → 贴 PR URL -4. 评审期(9/4 - 9/10)盯 GitHub PR 评论区,改 reviewer 提的 review - ---- - -## 作战卡 2 · Task 35 · rotary_embedding(标准) - -**目标**:在 FlagGems baseline 之上再压 1.3× - 1.5×,攻占前 5 名 - -**官方参考**: -- `FlagOpen/FlagGems/src/flag_gems/fused/rotary_embedding.py` -- 完整 kernel:`apply_rotary_pos_emb_kernel` + `apply_rotary_pos_emb_inplace_kernel` -- 关键观察:Task 30 (interleaved) 和 Task 35 (标准) **共用同一份 kernel,仅 `ROTARY_INTERLEAVED` constexpr 不同** -- baseline 团队最佳 4.16×(说明 rope 算子整体还有空间) - -**当前实现的优化空间**: -1. baseline 用 1D grid `(n_tokens,)` + 内层 `for off_h in range(NUM_Q_HEADS)` sequential 循环 -2. num_warps=8, num_stages=1 固定(没 autotune) -3. head 维没有并行化 — 多 head 模型(head=32/64)有显著提升空间 - -**我们的 v2 方案**:`src/flag_gems_local/fused/rotary_embedding_v2.py` -- 3D grid: `(n_tokens, n_q_blocks, n_k_blocks)`,head 维并行 -- autotune 5 组 (num_warps, num_stages) -- 智能 head_block_size(`_pick_head_block_size`):32 head → 8 块,16 head → 4 块 -- 复用 baseline 的 cos/sin 一次性加载 + 跨 head 复用(已经做对了) - -**预期加速**:1.2× - 1.6×(主要来自 head 维并行 + autotune) - -**测试矩阵**: -- shape: `(batch=4, seq=2048, head=32, dim=128)`, `(1, 4096, 32, 64)`, `(8, 1024, 16, 128)` -- rotary_interleaved: `False`(Task 35) -- 正确性:与 PyTorch `apply_rotary_emb` 数值一致 - -**风险**: -- 中:head_block_size 选错会导致 grid 太大或太小,需要实验 -- 低:3D grid 在小 head 模型(<= 4 heads)上浪费 program - -**PR 提交流程**:同 Task 29 - ---- - -## 作战卡 3 · Task 30 · interleaved_rope - -**目标**:攻占前 7 名(差 7 名是数据上的,实际有空间) - -**关键观察**: -- interleaved 模式 = `ROTARY_INTERLEAVED=True` 的同一份 kernel -- baseline 团队最佳 26.37× — **这个数字可疑** - -**怀疑**:26.37× 高得离谱,可能是: -- (a) baseline 评测脚本在 NPU 上跑歪了 -- (b) NPU 的 PyTorch rope 实现特别差(可能性大) -- (c) 之前有队伍把数据刷成了排行榜结果 - -**建议先复测**:在 NVIDIA A100 / 国产 NPU 上重跑 baseline(用 FlagGems 自带 benchmark),如果 baseline 不是 26.37×,说明排行榜数据有噪声,Task 30 实际难度可能更高或更低。 - -**v2 方案**:复用 Task 35 的 `rotary_embedding_v2.py`,只改 `rotary_interleaved=True` -- 完全零增量代码 — 是 Task 35 的副产品 - -**预期加速**:1.2× - 1.6×(同 Task 35) - -**测试矩阵**:同 Task 35,但 `rotary_interleaved=True` - -**风险**: -- 高:排行榜数据本身有噪声,基线不可信 -- 低:即使 26.37× 是真的,v2 仍能再压 1.2× - -**行动建议**:第 3 批**先跳过 Task 30**,把 26.37× 复测清楚再决定 - ---- - -## 作战卡 4 · Task 26 · fused_moe_router - -**目标**:从"未入榜"到攻占前 3(全新挑战) - -**官方参考**: -- `FlagOpen/FlagGems/src/flag_gems/fused_moe_mxq.py` (35KB,含完整 MoE + 量化) -- `vllm/vllm/model_executor/layers/fused_moe/fused_moe.py` (vLLM 的实现) - -**当前实现的优化空间**: -1. baseline `prepare_moe_inputs` 用 `torch.argsort` + `torch.sort`(两次排序) -2. 有 `.item()` 调用,引入 host-device 同步 -3. 中间 tensor 多(sorted_token_ids / sorted_expert_ids / sorted_weights) -4. 真正的"routing"是稀疏矩阵的 top-k 选择 + expert id 排序,Triton 完全可以在 GPU kernel 内 fused - -**我们的 v2 方案**:`src/flag_gems_local/ops/fused_moe_router_v2.py` -- 写一个 Triton bitonic sort kernel(`bitonic_sort_step_kernel`) -- 单 token 内的 top_k 排序(<= 32)在 registers 内做 -- 跨 token 的 token-id 展开 → `sorted_token_ids` 直接 broadcast -- 消掉 baseline 的两次 sort(只一次) - -**预期加速**:routing 部分 2× - 4×(整体 fused_moe 加速可能 1.1× - 1.3×,因为 routing 不是热点) - -**测试矩阵**: -- (num_tokens, top_k): `(1024, 8)`, `(4096, 4)`, `(2048, 16)` -- num_experts: 8 / 64 -- 正确性:与 baseline prepare_moe_inputs 输出一致(允许 token 顺序不同,但 expert assignment 必须相同) - -**风险**: -- 中:bitonic sort 的 kernel 在 top_k > 32 时会失败,要 fallback -- 中:routing 不是性能瓶颈时,优化收益被埋没 -- 中:与 FlagGems 的 fused_moe_kernel 集成时,可能需要改 invoke_fused_moe - -**行动建议**: -- **第 3 批先做 routing 的独立 benchmark**,看 routing 本身在 MoE 推理里占多少时间 -- 如果 < 5%,考虑放弃 Task 26 投 Task 35 -- 如果 > 20%,Task 26 是真正的金矿 - ---- - -## 提交模板(给队长填) - -每个 task 提交时需要: - -```markdown -## Task XX · [算子名] - -### 团队 -GuanghuLab(孙蓓[队长]、9478_apiqttc、陈淑婷、4348_apiratk、8592_apivqhj) - -### 优化方案 -[1-2 段说明] - -### 加速比 -baseline: X.XX ms → v2: X.XX ms = X.XX× - -### 测试环境 -- GPU: NVIDIA A100 80GB -- PyTorch: 2.3.0 -- Triton: 3.0.0 -- FlagGems: master @ commit xxx - -### GitHub PR -https://github.com/FlagOpen/FlagGems/pull/xxx - -### 附图 -[benchmark 截图] -``` - ---- - -## 资源 & 工具清单 - -- **本地开发**:你 / 队长 / 队员中任一人的 GPU 机器 -- **FlagGems 官方**: https://github.com/FlagOpen/FlagGems -- **vLLM 参考**: https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/layers/fused_moe/fused_moe.py -- **比赛平台**: https://flagos.net -- **FlagOS 文档**: https://flagos.csdn.net -- **飞书 / 微信群**:入群链接在 flagos.net 比赛详情页 - ---- - -**时间表(D243 现在 → D246 9/3 截止)**: -- D243 (今天 23:00):4 张作战卡 + 3 个 v2 实现完成 ✅ -- D244 09:00 - 18:00:在 GPU 机器上跑通 + benchmark -- D245 09:00 - 18:00:调优 + 提 PR + 提 flagos.net -- D246 09:00 - 19:59:补提交 + 改 reviewer 提的 review - -阿念 · 2026.08.30 23:14 diff --git a/zz-flagos-s2-2026/d245/docs/LESSONS_LEARNED.md b/zz-flagos-s2-2026/d245/docs/LESSONS_LEARNED.md deleted file mode 100644 index 9d6557b..0000000 --- a/zz-flagos-s2-2026/d245/docs/LESSONS_LEARNED.md +++ /dev/null @@ -1,89 +0,0 @@ -# LESSONS LEARNED · FlagOS S2 Track 1 - -> 阿念(AI 工程后盾)· D243 (2026-08-30) -> 每条教训都来自"已发生的失败",下次动手前必查这个文件。 - ---- - -## 规则:动手前必查 - -1. ✅ 平台规则(计算定义 + 接口签名 + 容差)看完了吗? -2. ✅ 至少 1 个队友的同题实现学完了吗? -3. ✅ 上次失败的原因有记录吗?(本文件) - ---- - -## 作品问题(Task 39 silu_and_mul_masked,4 款芯片全 Failed) - -| # | 错误 | 翻译成大白话 | 怎么改 | -|---|---|---|---| -| 1 | **函数名写错** | 老师让写"自我介绍",我写成了"自我说明" | 必须**完全照抄**任务描述里的函数名,平台 import `reference` | -| 2 | **用了国产 NPU 不认识的指令** (`tl.sigmoid`) | 用方言写作文,改卷老师是北方人,一个字都看不懂 | 手写 `1.0 / (1.0 + tl.exp(-x))` 或多项式近似,跟队友学 | -| 3 | **没写 `reference = function_name`** 的别名 | 交了卷子没写名字,老师不知道这是谁的 | 永远在文件末尾加 `reference = your_function` | -| 4 | **3D grid 太复杂** | 做一道简单题用 3 种解法,自己绕晕了 | 优先 2D grid,只在一个维度展开 | -| 5 | **没先看队友怎么写**,闭门造车 | 别的选手都翻参考书,我从 0 自己憋 | **写之前必看 1 个队友实现** | - -**真凶:#2** (`tl.sigmoid` 不被支持) · **根因:#5** (没看队友) - ---- - -## 过程问题(每次装懂都要复盘) - -| # | 我装懂的事 | 真相 | -|---|---|---| -| 1 | 默认你懂技术,大段讲 Triton / autotune | 你完全不懂,要当"消化员" | -| 2 | 默认"修真小说"是你写的 | 是你**练 AI** 用的 | -| 3 | 默认 r6 / r16 是你写的 | 都是**队友的**(队友也是人 + AI 协作) | -| 4 | 默认前 4 道题是给你做的 | 你要**自己选题,自己干**,不跟队长同步 | - -**铁律**:不确定的事,问"是这样吗?"而不是替你下结论。 - ---- - -## 协作模式(D243 定型) - -| 你 | 我 | -|---|---| -| 找材料(队友代码) | 看代码 | -| 看 flagos.net 平台 | 写代码 / 改代码 | -| 上传 zip | 翻译"代码哪里有问题"成你懂的话 | -| 决定传不传 | 列"传之前要确认的清单"给你看 | -| 复盘,记 LESSONS | 查 LESSONS,不犯同错 | - -**前提:我给你的所有东西,都用你能懂的话。看不懂就骂我,我重写。** - ---- - -## 比赛生态(D243 认知) - -- 整个 FlagOS 比赛,**很多参赛者 = 人类 + AI 助手**,不是真人写代码 -- 队友的 r6 / r16 代码 = **队友 AI 写的**,不是人之之写 -- 别人质量高 = **别人 AI 强**,不是别人真人强 -- 我的工作 = **学别人 AI 的写法 + 搬到我们题上**,不原创 - ---- - -## 已知国产 NPU 套路(从队友代码学的) - -- `with torch.get_device_module(x.device).device(x.device):` — 必需 device context -- `enable_fp_fusion=False` — 某些 NPU 不允许 fp fusion -- 指针 cast 到 `int16/int32` 防止 NaN 被 fp 转换吃掉 -- `tl.int64` 显式 cast 防 stride overflow -- 手写 erf 近似,不用 `tl.erf` / `tl.sigmoid` -- 2D grid 比 3D 稳 -- `module.device` 国产 NPU 必需 - ---- - -## 下次动手前 checklist - -- [ ] 任务描述完整读一遍(计算定义 + 接口签名 + 容差) -- [ ] 至少 1 个队友参考实现看完 -- [ ] 把队友写法 + 我打算写的放一起对比 -- [ ] 写之前把"我准备怎么写"翻译给之之听,得到认可 -- [ ] 写完本地不传,先给之之看"我会传啥" -- [ ] 之之确认后才传 - ---- - -阿念 · D243 23:50 diff --git a/zz-flagos-s2-2026/d245/docs/PR-TEMPLATE.md b/zz-flagos-s2-2026/d245/docs/PR-TEMPLATE.md deleted file mode 100644 index 411aff6..0000000 --- a/zz-flagos-s2-2026/d245/docs/PR-TEMPLATE.md +++ /dev/null @@ -1,64 +0,0 @@ -# FlagOS S2 提交 PR 模板(给队长孙蓓) - -> 复制下面这块,直接贴到 GitHub PR 描述里 - -```markdown -## What - -Optimized `[算子名]` operator for FlagOS Open Computing Challenge S2 - Track 1. - -Baseline: `FlagOpen/FlagGems/[原文件路径]` -This PR: `FlagOpen/FlagGems/[新文件路径]` - -## Why - -The baseline uses `[默认配置 / sequential for-loop / fixed num_warps]`. -This PR adds `[autotune / 2D grid / 3D grid / 头维并行]` to extract -additional `[1.X% / X 倍]` speedup. - -## How - -1. Added `triton.autotune` over `(BLOCK_SIZE, num_warps, num_stages)` with N configs -2. Restructured grid from `[1D / 2D]` to `[2D / 3D]` for `[维度]` parallelization -3. Constexpr inlined `[常量]` for `[省 Python→Triton 转换 / 减少指令]` - -## Benchmarks - -``` -shape baseline(ms) v2(ms) speedup -[B, S, H, D] X.XXXX X.XXXX X.XXx -``` - -GPU: NVIDIA A100 80GB -Triton: 3.0.0 -PyTorch: 2.3.0 -FlagGems: master @ commit [hash] - -## Tests - -- [x] Numerical correctness vs PyTorch reference (max abs_diff < 1e-2) -- [x] Numerical correctness vs FlagGems baseline (max abs_diff < 1e-2) -- [x] Tested on fp16 / bf16 / fp32 -- [x] Tested on multiple shapes - -## Team - -GuanghuLab · 孙蓓(队长)· 9478_apiqttc · 陈淑婷 · 4348_apiratk · 8592_apivqhj -``` - ---- - -# flagos.net 平台提交模板 - -``` -参赛赛道: 赛道一 · SGLang 框架算子在多款芯片的性能优化 -赛题编号: Task XX -赛题名称: [算子名] -GitHub PR: https://github.com/FlagOpen/FlagGems/pull/XXX -提交说明: -- 团队:GuanghuLab -- 优化要点:[2-3 行说明] -- 加速比:[X.XX×] -- 测试设备:[NVIDIA A100 80GB / 国产 XX 芯片] -- 备注:[可选:跑了哪些 shape,dtype 等] -``` diff --git a/zz-flagos-s2-2026/d245/results/gelu_and_mul.zip b/zz-flagos-s2-2026/d245/results/gelu_and_mul.zip deleted file mode 100644 index e34ec73dfaf58fe237fbe0b820abc279d2f508eb..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1212 zcmWIWW@Zs#U|`^2;M=VfyZV>pVGb4shEjG027U$^hV;~&()h%@l=$4z9KC|d&=5`r z=GFEyvOu`Bf}4SnAR793nRLI3?t3=`! zWf5_*i3sa^_UHA3r^#^CpEtSz;cccC8 zT+91E&y}B>V>#33nH;C|(bEBo{Za(Z8{Ijf>HIxo|L-56>5l`~|J?JrWdG)mt-EX% zi!1#5wq*NWx%#c*_51GD?PZ;Q#P9VF^&Pd(lRM(;F56Zuec5%%>gG$c>!&YNmr2e4 z$EWo%=T+=V*CVU1D!1#r?8_H&Z@xDF_=U<}8@?w+&z-p|#>68*-+smY>;G(>uHU_P zP4me9-17N0fdLM`kEW+;JnL4ii2rm)#oKF7^_`p9$7atjo1Gyl6sjqhroJv~@spKn z0-B;9m;bfCKH<;tbD!I1Eh^!^SHl0T>TK_~>Gsp77GBxH-+kv|?4j3IsaIXU>Jb$Dt1F279|Z~or;^}GDXgXMGA^P8&Z zopbEHDaN(C>ht06b}K|a9rcWn%AbMc)Rs(tC_~!?~+-y?p4v+rS}iI{nV}b zxOnonYl+p*r~mlCtH#&e6Q;R%p65cND`qeLUw>WkSs=&!gJgTf-_$)no0p{+NvhuU z3Yn`r<7{ny#Zv2vZGtNU6%-5H1E*-Oe7NtdXm{8_3;9bK0nY;6V@;XP6+Li|`4?ds z`M{{?#j^LUbJnr>Xlpm{e-00@lwx{zU80cpcF^ajB%|8n5r>}T#4w1c9MD>{lp{Ie z>&G2(6}=}nR4jSnEAhmvHatRyi5hTcE{kWw$ZK$V0KDud`!1i_P>M zVMn};To3;L_2Bld2kMI?Lkr7(EG~|B?bDj{p}ed1(GeNe#XXCDt<+z1OkutM!)VOKi=H7W>*WVevkLw;a2y;!V@IKA|`%d66389HE9f)EM^;@X34)o z`h)4=jnXAM_&)Di_k2N)ZpJ3%RPk_!GwPLZlCnY_-m;(n`hIoh&C^-`cV({JJ;yXX zP~UH1;MHkalOL)E{^c|kveEdrqR~MAi{6TKiH|%NRbo}19T9Nhd-`Oqwf&T0gDNKe zJxeRxDv#JW_Rjsj&2Z%(_5g22COKwYMT!Kls9|6LmKh978bK_K0)`b*z@QZ_0p6@^ RAmxle7z(5Vfn^i}0|58@DmwrG diff --git a/zz-flagos-s2-2026/d245/results/interleaved_rope.zip b/zz-flagos-s2-2026/d245/results/interleaved_rope.zip deleted file mode 100644 index 233e2f853049d633e168c8f69ec6db618bac97b0..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1568 zcmZ|Qdpy$%6aetwX1K|SVaanID=`xxd6WuCBuVBq`iy7#~Px%YEE=X1_K=lpd(-{WW}1d#y%;5T5YNzGHAUGvFL z7y#zC0{{$A1cE{@`A6XWeWLySun}S5{<`6DXPvhJLiMI3-1DuBiro%?AOa8o{Jo^o z#XSsWBSRR+Uff0TGsf$TBtOjb%_yiq;beBaokI`{KJ(uzzgMWQ&-m(O85CbR13Hp% zg6Nn^B)U1tLXM>#-$Y?{uffd?T2|ui%VnwfINk9Weq=yeS{f_&L#L}e3B4yq3sZ5c z>uR<`65{|>c*#3xzv~x@P=kGVsOY5rUBr?WdG;j4qDfN!ZYKfa1^k-u&RYmidPzOIoT*Le5vsdLJHe-BX}hH zJB=Q*`HZ&l=U{23Hp~%qu6vT<%o3*r_mn>2v=rv{zgg9t)UCdm6+NIFWyD85Ha0aL zG|fWC7^y6VhGZepF1dk;$j2ocZC;#&Sd7ib9&StzQ^nL>GIu<2xRn-k!EOFUJ*O>{ zZ+@BG92IgzL72RId~AGZU7Fl}3ld0xq39wSa!-=fy#o{OMGESixA!nNH?>z0YoFIx z18BVjF6=mevu1T>lGXPquiri0o*qAj8R3l*vhno8WBlQRxbLQ2oDrrhyC zz*tEsjYT%U;){#V4qu-c?JyG$U44^>?&!t9Xf2rQcALuie)d)flSI zcD5S%={a2P*(Bfh10NMN)!3AF@_Vk%S`*xdQ4xRKSECRvG($aiR4pO5b2>N=czW1WcB0+(<*vGClKPRU$eEHoyX< zOG*g%`OBgvMWL^5cS#kgnpzjwqQb7#foBEocfIx$M42;oeakBt7!FE&LD6p;5Zs$5 zi8pEVd|=l0&4Xq4tr}Q;EP6T}$pv$lN!90={pv_(?+9tL-hnLNIi;Olv@@ObFdtUQ z-AV<_ka06EEqjax*MBa1tIT9DJHk>WK zxtnX#cSF>lUZ+}JTnp+QU9jyJvgKMlqNJC6u@6K^DO(%9MhI*<5YIQ z6}+=b0|Lqw8wOjosEL519SE!}^nYaU+@c%+wmb$Zl7Dso3Ape-fd30`M>}DWUn8Kc Nz--mptv>(&e*yoN;otxO diff --git a/zz-flagos-s2-2026/d245/results/rotary_embedding.zip b/zz-flagos-s2-2026/d245/results/rotary_embedding.zip deleted file mode 100644 index c9b9eea35385abdded0c6b5998e934b0c6097a0c..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1614 zcmV-U2C?~2O9KQH0000802jO>T`YUK;*tdb0O}0@022Tl0CI11VRCt2Wo=?*WMpY> zXD)DgRa6N81D!#bYMnutYIS%E009K`0RR956aWAKrB_>T6Gs$&-(PVWsq$`coWzib zTC7ByKqHklh@=$~LT2n8+XL&}Wp>sveWI7BH;G6E+9E)-s(=SnHC2l!X=$YXmTVIK z1!pd{*K0RQ+q~F5x9^*C&YUxun3z~4@s*=keZb~*xLKydb;uP{&m#y`Kk2o-6cxAbp4l|ttb1BzfQLv?LPfzck{Qs zt!*1)rW;>ronu+Mx1R6)`Ax-|0{jMESdRtA10$;&KoZdkIjYhxzlTdSAT^&>8&MUH zj7M}&6F;Vx?DX!7ovp|4-r4sdefaJEv!9DQf7}2*I}h(poeR&NLi&#!8RWAUySF!Y z?>$L3|JvEUzjyl%T;jArISuL*+>xAPkre7A02CAV!|>a5%PMg z|AR9T2E)HfK4vK)O$sCi;aiq<(kL!7VR^IxRRb$48`~D(U$+_wUW9rW>n|}2^p`LZ z>Wi{ks$jrUJq%(&yAf&^^Ma)}&z#Q5l;`=?mAvWsrBTshUZk8CMx7VNTds^+u8g-_ z9<^K^CA2)=`K>d{3-9C=PDSZ>xQ2yARm8(4C+#Zp>={R?1oP>LZr13+D>P6Yu?ExA zu%v`7>}VsKG8U6YrEaCGLr!W!a72bC9LrDbnW7fz=pEuS8A*AaP*t)_-D@OPq)c=y z6vk1&ld?p{igpw`jE|SzvsJsGopC<`b{WlObVs@^uIN0C{Izk^xQYyRmS|D#&Xi`H zk_c^#Z03n$AVbeP6e50~JrMQqpxFK_pczziA7o}teil@I6z-(#!vB5H9HO~+8N`(= z;>tfD9u7}=${TIUY3e%JJNT4I`=nykx7L;D8!OWzeN*k*H_r9ba<-qAO+O*080Q$Z z&2&MFzH6?R@CFw9l#7{jI{zk8uSfdPkos^iOHalcE~qQjn0-j?7V#VAPO}EcUvR>O zqvc9J|D1}12Pn!L%^w>vnm?8|nk^q6u|k=Z`B|QuH^lOb!tSbJtp;@wSRg3(Dd|ve zz_?oqJl2NSkaGryFwD#c#82oc&O>e&y3kHyfvh?d;ueW0Tv)|9jE99QYb>9 zO$AY(QCx8HXy=kZuyzlRAQbUoRqa3^^*Z$_C)hMlBXw+24^P%(;}tT~7_%$;7C+XQ zCtC6flvIX#^tAYOC?9lOyo}%z%rF)MW8x!1C%i^MwzL6MP2rk|IRE=Cg+=byu`IJJ zlE|lWFljE1RM>sKgPtAF2K~m?dvSWYoT)dKdK{`0{})vBI}EBklrFdo-z&a^m%yW% z1KV2N5A|B@qhrc3`?4T0X;aWa3WKp2B|=Rud|;JmI}+DyQ}7~qvIAWNmY+rA>h_3a z8a7DU)VCdQU6>iU-Wo$m)u2JNk%^I2M|-tHFVQay zb%Z7rRn4!eHTE#yHHN*WHAX%#7*HR2EMAbrF4{vIXlN4Y6ArJaCD*WDbX9-rSE8%< z4;kXy15+v=;~qX>9Q8yT2xcAiKpY6l&agYKW>?*dK>%FPSF4*J)^ zdIeBR0RkQa6aWAK2mlwnB3&$dx#E%q008O@000vJ7ytkO00031005+c00000a&L5D za(Q26ZDM6)WNB_^E^v8OR0RM7ok5stb$AN^0R;5{000CO0000`O9ci10000100975 M0002r1poj504&GlC;$Ke diff --git a/zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.py b/zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.py deleted file mode 100644 index ff080a3..0000000 --- a/zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.py +++ /dev/null @@ -1,55 +0,0 @@ -"""Task29 gelu_and_mul: 跨芯片通用版(v8.1 国际卡写法 + 国产 NPU 套路). - -作者: 阿念 (anien@guanghulab.local) 为 甄静(8592_apivqhj)· 队长 孙蓓 -版权: 2026 GuanghuLab -基础参考: v8.1(国际 A/B) + 国产 NPU 套路(r16 M-RoPE 学的) - -设计选择: -- 用 `tl.math.erf`(国际卡 SFU 加速: NVIDIA __nv_erf / AMD __ocml_erf 硬件实现) -- 2D grid + BLOCK=1024 + num_warps=4(国际卡最优) -- `enable_fp_fusion=False`(国产 NPU 必需) -- `with torch.get_device_module(x.device).device(x.device):`(国产 NPU 必需) -- 函数末尾加 `reference = gelu_and_mul`(平台 import 入口) -""" -import torch -import triton -import triton.language as tl - - -@triton.jit -def _gelu_and_mul_kernel(x_ptr, out_ptr, d, BLOCK: tl.constexpr): - row = tl.program_id(0) - col_block = tl.program_id(1) - d2 = 2 * d - col = col_block * BLOCK + tl.arange(0, BLOCK) - mask = col < d - - gate = tl.load(x_ptr + row * d2 + col, mask=mask) - up = tl.load(x_ptr + row * d2 + col + d, mask=mask) - gate_f = gate.to(tl.float32) - up_f = up.to(tl.float32) - gelu = gate_f * 0.5 * (1.0 + tl.math.erf(gate_f * 0.7071067811865476)) - out = gelu * up_f - tl.store(out_ptr + row * d + col, out.to(x_ptr.dtype.element_ty), mask=mask) - - -def gelu_and_mul(hidden_states): - if hidden_states.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - d = hidden_states.shape[-1] // 2 - num_rows = hidden_states.numel() // (2 * d) - out = torch.empty_like(hidden_states[..., :d]) - - BLOCK = 1024 - grid = (num_rows, triton.cdiv(d, BLOCK)) - - module = torch.get_device_module(hidden_states.device) - with module.device(hidden_states.device): - _gelu_and_mul_kernel[grid]( - hidden_states, out, d, - BLOCK=BLOCK, num_warps=4, enable_fp_fusion=False, - ) - return out - - -reference = gelu_and_mul diff --git a/zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.zip b/zz-flagos-s2-2026/d245/src/Task29_v81_zhizhi/gelu_and_mul.zip deleted file mode 100644 index e34ec73dfaf58fe237fbe0b820abc279d2f508eb..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1212 zcmWIWW@Zs#U|`^2;M=VfyZV>pVGb4shEjG027U$^hV;~&()h%@l=$4z9KC|d&=5`r z=GFEyvOu`Bf}4SnAR793nRLI3?t3=`! zWf5_*i3sa^_UHA3r^#^CpEtSz;cccC8 zT+91E&y}B>V>#33nH;C|(bEBo{Za(Z8{Ijf>HIxo|L-56>5l`~|J?JrWdG)mt-EX% zi!1#5wq*NWx%#c*_51GD?PZ;Q#P9VF^&Pd(lRM(;F56Zuec5%%>gG$c>!&YNmr2e4 z$EWo%=T+=V*CVU1D!1#r?8_H&Z@xDF_=U<}8@?w+&z-p|#>68*-+smY>;G(>uHU_P zP4me9-17N0fdLM`kEW+;JnL4ii2rm)#oKF7^_`p9$7atjo1Gyl6sjqhroJv~@spKn z0-B;9m;bfCKH<;tbD!I1Eh^!^SHl0T>TK_~>Gsp77GBxH-+kv|?4j3IsaIXU>Jb$Dt1F279|Z~or;^}GDXgXMGA^P8&Z zopbEHDaN(C>ht06b}K|a9rcWn%AbMc)Rs(tC_~!?~+-y?p4v+rS}iI{nV}b zxOnonYl+p*r~mlCtH#&e6Q;R%p65cND`qeLUw>WkSs=&!gJgTf-_$)no0p{+NvhuU z3Yn`r<7{ny#Zv2vZGtNU6%-5H1E*-Oe7NtdXm{8_3;9bK0nY;6V@;XP6+Li|`4?ds z`M{{?#j^LUbJnr>Xlpm{e-00@lwx{zU80cpcF^ajB%|8n5r>}T#4w1c9MD>{lp{Ie z>&G2(6}=}nR4jSnEAhmvHatRyi5hTcE{kWw$ZK$V0KDud`!1i_P>M zVMn};To3;L_2Bld2kMI?Lkr7(EG~|B?bDj{p}ed1(GeNe#XXCDt<+z1OkutM!)VOKi=H7W>*WVevkLw;a2y;!V@IKA|`%d66389HE9f)EM^;@X34)o z`h)4=jnXAM_&)Di_k2N)ZpJ3%RPk_!GwPLZlCnY_-m;(n`hIoh&C^-`cV({JJ;yXX zP~UH1;MHkalOL)E{^c|kveEdrqR~MAi{6TKiH|%NRbo}19T9Nhd-`Oqwf&T0gDNKe zJxeRxDv#JW_Rjsj&2Z%(_5g22COKwYMT!Kls9|6LmKh978bK_K0)`b*z@QZ_0p6@^ RAmxle7z(5Vfn^i}0|58@DmwrG diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.py deleted file mode 100644 index 1c6ad66..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.py +++ /dev/null @@ -1,78 +0,0 @@ -"""Task30 interleaved_rope: direct row-owned exact copy in one Triton launch. - -作者: 阿念 (anien@guanghulab.local) 为 甄静(8592_apivqhj)· 队长 孙蓓 -版权: 2026 GuanghuLab -基础参考: r16-整条取料流水线-待平台验证(队友共享实现) - -改编说明: -- 改动了 docstring 与 author -- 保持 r16 全部 国产 NPU 套路不变 -- 函数末尾保留 `reference = interleaved_rope`(平台 import 入口) -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -@triton.jit -def _interleave_tiles( - input_ptr, output_ptr, - N: tl.constexpr, D: tl.constexpr, - PLANE_STRIDE: tl.constexpr, ROW_STRIDE: tl.constexpr, - COL_STRIDE: tl.constexpr, H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, -): - if input_ptr.dtype.element_ty.primitive_bitwidth == 16: - input_ptr = input_ptr.to(tl.pointer_type(tl.int16)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int16)) - elif input_ptr.dtype.element_ty.primitive_bitwidth == 32: - input_ptr = input_ptr.to(tl.pointer_type(tl.int32)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int32)) - - row = tl.program_id(1).to(tl.int64) - col = (tl.program_id(0).to(tl.int64) * BLOCK - + tl.arange(0, BLOCK).to(tl.int64)) - valid = col < D - safe_col = tl.where(valid, col, 0) - phase = safe_col % 3 - from_b = (phase == 1) & (safe_col < H_END) - from_c = (phase == 2) & (safe_col < W_END) - base = row * ROW_STRIDE + safe_col * COL_STRIDE - - a = tl.load(input_ptr + base, valid & ~from_b & ~from_c, other=0) - b = tl.load(input_ptr + PLANE_STRIDE + base, valid & from_b, other=0) - c = tl.load(input_ptr + 2 * PLANE_STRIDE + base, valid & from_c, other=0) - value = tl.where(from_b, b, tl.where(from_c, c, a)) - tl.store(output_ptr + row * D + col, value, valid) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - _, rows, d = x.shape - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - - block = min(1024, 1 << max(5, (d - 1).bit_length())) - grid = (triton.cdiv(d, block), rows) - with torch.get_device_module(x.device).device(x.device): - _interleave_tiles[grid]( - x, output, N=rows * d, D=d, - PLANE_STRIDE=x.stride(0), ROW_STRIDE=x.stride(1), - COL_STRIDE=x.stride(2), - H_END=min(d, int(mrope_section[1]) * 3), - W_END=min(d, int(mrope_section[2]) * 3), - BLOCK=block, num_warps=4, enable_fp_fusion=False, - ) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.zip b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope.zip deleted file mode 100644 index 1e09fd8514e7023e7411ce36e44ac75cb493e319..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 10366 zcma)?b8zJAy0&9;VocD%#5O0kZA>t+CN?Lwor!JR?%1|5@q}Mye`oK#*4k^=Io0)6 zb^mie-Bn%Jd)57Wnw`T~(A}LBMOdlg)o0M;CYyP>5Sl5RiX_tE~;WRunLE%ijeR43P|#4-5^}RK$P~4itZ()%f%a ziwYxPa1-+iS+JFRbs{|b7Q@=jtTUsu_5kkFdHpP;Z2>H+^ZI&HBZ+b zJgo`plxNLv4NvdXrtIVb-?Py?CYn~)GK2lq%WszZ$zmn?xuL`PrKRC%d?6!(8cm*7 zZE&ks8HiA9F`QAp)|wPMH*rt5MOPUfJ%tm--Z52gMuNrIW zOpmwd(`kdEew%O525Dzl_)5iM7 z(&LBpu1E+|FErW?C{oOF9}<02_gshD8fEz5k(U?hd*WX|e*x#|8QhQ2#Lr)kDkmK~qh&zHH`|!leWENqSy}^FT$ZE@sVhVSisoMPbpq!qdjqpV}6kRR~Xi zPgKLNnPr-_3;DWNT|H-$9HRO#a&G<bi9kvuV(L{q+0z?f$bvnRJ^8G{t;V|}16TrpYx2q%zgCunMC|EI7$-!Ej# zd4BgB^bOA={vsRi=z$mDo6FF*{<4tWKK%S$eQh5u)(Cl(VJ4 ztX@U7E%g&+`m`gjn1m5MQaT5D|5;{OC39K;qIrHOYWr-@ObK;5r2%=olKCPahNqYJ zVTGZ9ZC9ZFc={Bjc+!1nxXv~L-x_OkPT)Zssj^Vh7)0z5i0HQ38+DVB%(6;) z+m9HWdw|1V7LcNWg4AuAvbPo!A*Pf4R6n-!NS$|+b($1VzywXYR2i)-vht?0oUPM@ zGxKWc=(ye36eRr|OY^G<%b+*UO^=um2;dGRNw6a0VFn3Ps~_raC~uUx-0cdzoHp{! zVoz`yc20)2*(&z;*J^MaaUbFm=a98>Xrv1pr(^z6} zHg|P%x>!Ndnyqf6bQL`|gwr;3(aWP_Ev{0upryAp#FFZ;@Vv{S-&O6riBC>y!MYKw zuWv<$Lo=}$c|rZ`P+%UZ(RF(d8J^*h%DK22@wyN8gzQ~L}CtH;+U_bI5_P}#*`M7K$LZi_*P;qUzMZ|*De6eoDc;IPd_Td z#3UQPZlRsB-Wx>zPrx1=tVF|j2X-192ng=q09((%;ft}g(LcFe_s;C!=P%rjc;_~r z$UnFp46($DH<(m_tbnPI8muauFA3%u{A13@*%EmvL&wDA%l$Z&i1}zfGNf>h{9~U~ zjEg|LWYGujz!tMtNK?A%T3HnX%9?|*W{(4=jp3f29;6uPHB9Z2p=hQiBD9ei6$%bW z1%w=Idn!SPc$VCdACxeXibOoD(>eN~Y`ho0Hjs5A`Sk*KF0y1YUqC0acmi{6tO?$} z5x)LH3q+3U^n{ArIPYw8mV@Ui2fDg+0^0bC*oHgUT_0`^r(Eps2S$IImJ%%=_+k^; zetW7MPTKgm*0GlxHM~X-Hj(CWTYu-ZcPxgzheYL^B=Uq$(LEAgLe?J(Ua}btBia-@ zNkStuSwkwQC=N%=056;W?CjC&{&SP@G4uK1dUyWfuozN|fAw2Pjh}xJ6Rzy}GgwWL zu;308N%Gb8ryZ;)!glO8H=Yhe(L5bm8X6Z~h25@Fk1RM?8hZGgO(wNEdhvZ%*=)EN zgaT+9z3!7j@|IJzJh^I^?8!+fGh`;kpGNvb4~sZBhIASGCi6Q z2WqM<)bhbcfvO8w8KSu}H71WFFug4v(lw#hFbXi<)k&!N6bl)2L8EY^rh}i6)I-MK z>J0rPX>11+@nj~z+$!zqu)ud*_^D&TgNaO6Z5p`_1i$t!(r z*cWKy`jtptY9>ztV9sFOuPZE;KS`(X^&{ZXF-ansS$^{fOH4`#mC7fKdotKPDP7z) zQM4^SEb3b+O~b0N>|5d1N?KswKM(JoziaF7cMX|pMYqxB+<9K64iJ56A4jJ5VyJ9) zaCO)Tw_&%-r?4$?I7QQK6&h3Q+@t!wiYk*(Wizl)*I9U;7;e224|b>1I%kduk5CAb zi{g9X_gFp&XGO#v6vL?L?>OOyha!=2DRYRM;+6KZBa_-0?^!pO9bXJMIkRlWU|S1 z=D-h%Vtzr|7m~z8QB!?xlz83(A9}y&7aZIbMj6)3O7k^Pg3kKKXqH|GZyJ4GA@b{{1x@cWdE5hcHOV zkUbO0jzM#hxf45Ul;1_igjFSttb zWDbJG_ANy@Gj9RCT}+<1jew4$g}c)N63fa2FoRP^$S_8?<&Xe%{=0i)JJx2{ zIS@z3y9N`tnyKZRsPwB76qdYtFql_IeVvPYcde>#l~4c zDxM-U*Fd@IiO=q*b~bI#pQEQI@QeGs$)w-Vn5gDO+~C&u!Xd&oH_P>uV?b%$OMituLo|&@q^ZIsr4~~bWUCj|ujzBcu&YYjfth$~*ce>ov zfAjKgrL5|7c6ka(Y-?=125Y|eJ6kJYA+q0js3`^4Y&M3;|ghdbT2&*EwPt0-MaJ?NtN)*jf-W?vg=A%E4c6?Yop7rLVD)5 z`Nd}*RG!Zr^_DJl@82HE^Xr>%Y4ub9Xq0TWAF77VIr2#X?B}@52Dl`^3h)(2qO#Nx zy62vW)M*)rJFfbb3eHL?oENkYJGM{DK+Kz4UD+TOh9Fc=t^&Vh+ykm)cq@|B2nku# zTAxlNUe9bOR&0t!DHx|vaOW;vqDJW*E)1Ts?pB!x14u&bQYzn4DF=U4`pZ%&v@g># z(lmCGfYy?SRMq`Ni6(Huq|AvdL43!MttaC?f2^+>nWqQwxEY8tTktH7Q}^Br^EdFw z@GaWq%`(sP_wpY)^Jw@o_r#NDbfrsw8Iq@b%@Sw&P4a3BHwcbP5GMc;|o zlLSSml%m9Fd>3qwi7=YR21#?p&c3v(jOP#`oIK_*HP8xEohDF{iqIK^tA2*>R+tEc z9;z6m25PsDE5Xv3ikk*}pRI7+^7~k;ZTKOSrXVsGKbW%8AOCFV>XZU|%T#6#&Dn14 zs|?!dC>z2p+D^nnMv{oCe=0>xG?@UcJq-n9roF6~FlrL;7N)X|YQGq33KqV&Fj*v} zL9M4u-Ohk4C-(3SVM_-24xUbf>E3tKSk{nZ(wQ)tL3kb=q7@X6jwgqwXAX3-lf?G~?+91z)4HiqOg{vF1Acm70>+PYP<(}LMa315A zmKH3aaCtP#->VXP68_@ascYBw z$>jT?eW3G7B+r63Q$jEu9A4Aq$xS9QdsQA@J-{zj37!*%JtlLt7F}j8$5L%`2l?m z!j{(%UBRo@3DbQwtaOHX-X=M6!7CgV*v)Qoc*vo1KLlC#)HD(qk!n@#KGZ-EVx}r>V9^1f3Xo|Rkoi5gE9Tc(|2M{~I z5K>vfC`avTclUgE*Uhfs8C?txbwNnMobpRh!;lVPmEScEw}1<;ohVJNl@}{dz_rSy zk0-WcD}Z%rO`rs&mDX=oYDW7-(9p#|{s6l|=c$oNMJzbu9?6C$2yax10flo(eTgg( zlnppU6_-Q!F(-5}KY1r@w`tDX<(Y2)V%=(x(w7tL1FxP0`ptydvOkJ>=^GJzg2rf{ z;%TzClqf>|@EHnqoQ&|=ATY>MzNa?O)JutywRXxZH;peLxyZ1=^oZa3u@drYi~K9n zm&DImkZiNGo*E8RICyQyrM@e($*1FXnZGE}Sn~Wdn~*?l^-+`RkaQ}U1e1nBHLS5Y z<6~q&eK2fWFgG0T8TA`p_>CUx!x4izD$`C(t3T&<$NBvXm}0yZN&jr~S^pRdtbtQL z>N9WG)Pom;L8cL8LXH4ru0!t@?}g>vMM2a=4xqTGOvFR8Yf#e(@m{W(g|PxSg2DL((`R{8%zGag9i8TvFoua)Cr2M{V?s zLO9aV(%Ci`pV4>2rS$qJp5mh<&0WBce27F>sCj)C%arz7F#j(bgyNXvC&O+XWu$D% zr0$~xPB`2qlTBSQCRp)I$ z^XfswTxll|G&%emHz|^t1N;+hZ8I-5QU|XN&BQ8I!M({5^R}XvnKjhu4T=mC>6GRw zyM$JjH6r~MOcwV)=R*ffr*h?47nw|4op@1RTQB#2QDI#t738hzL3=r9fS6Zo>_5bOW({ecE|hsYo586B<%BkU~5XLA!+c zjBvE`3fZv>Q;#z9Pwn=m8fQ@;YT-LQ#Q@Si4sb3Sc ztfO~h9v)?noF7s)mZ%;VGjLyxY{^m3OI~PZtkwDi@J+yz&u}NucciT9-_@lB+DWee z+`J6Xp2O)_SP`3c71~x_BLHa16&l&}6SpEYckrj7({d6F7$XeAkAs=usB+jdQ4zO` zsE{Iv0ywd2$Bb&GP2&Xen6O83o zirMbXxL4%!)KkJImRG;9F#~NFu`{D1c6=QB4Tv34Ef^_EM84CEjP#j;mK-?ODj5T0 zPhR9*ZDfaZ(tTkN!6aAVJ*HzCSPKtSE|yWd41@mDOY9j9AHp#@Ri-*w303CT|$ZcC{B~k zKIV4(4FBhQZJbZ@Tj9INpS_Db{@>keW|odl2HyWQ9%d8eMxuu4cy?zJUMr=4Pr=Zi2p|^8vEqhleyN>PfU>HB~Iy;f+3pI?N~O$j_5= z(Qs@9#6~PXe2rjjH0R?wJ6idkEI(Rzw5RL7Ldm^0`IJ`jxBq&@e|`PPPPoC-5^ywe zuz*Zh5+Pe!3uSW3Ic5~t=X+xX$Vbz=k zcm6Xe<0a3zQVeL=jIj9Pb}SNP2ub#&X<##_q)S z^7Q1c)jr9%D_{JJg>S^ulnXt3I4pXgmuUIk<&n#mPx0YL3^2uBu9rAe`VE7>?sw}UyqYdehV=O#_ zI!s}>f2o#JiaeYB=(*9?x3+-NC9(kN@%3W) z*lWu`7ORxCHe}G4OiMUOI1I2r{)KM931s4P3ZU$zTDS@4XF)3)f z=(xpKb21gntYmX73^s~Ajk5(>Pp#XeFGdrW9KaJ%v_k<67%3$Z1wiX}`@a?Ng%QHyx}k9K&icagJ0yXu`8@X)WKEMIk6yGvWouvE%zHx3e&6%IYJQpnlbrO0yy$&v@={+Od!x3P4 zj684|DVK%bV7GW;!(w(b#(hEg#`?og<|_negc*g#2=6G@+x1iFjNHQ3Ca{= z7TS*U_@;DPduV~B{oc_9KSv+dPCDu~+Rxv9YDrnxv1)ge2$L&XZj$Q9tl3FQ2v@AJ zne?!Trt3gAoP1ViOUQT%%)-b|gE@6HZ`5(nO=5cdL^|;Eu08|t@T5H*!Au{4qVDgUcKHIw$DpbL zmV~trMUuIbghf7(mzJ!)9WU&p@)M(cj5f{|o{xq%i#L$t_(SjgXfX7|R1L?{1O+P> zUTV+5*W(`-Zx!W7SOl2gsx_I@9?CCrNWCO9Ja|T97>e5`hp3HBo*wc_%M+52^smiB z%dnZ?ghS7G9f=i?1*|IqLzN2%qSV1EH)jQ?Y*V5nM3I-jyOpL231H9P@)}zz1uHK~ zDT#+__h6U%VEXxI!r#Z}ASb1#xTO_-PzOc>VUUY;u6rQ4O`7{i(KVz$igBVS=|XcZ z$Mi2LF!ac!mV7#YZkb9YUre&XKK|5Jb}X5J^?$50lHTi#cl7_0&?EwhJAa(Tj0c)5qp}p`=^7NQ44YO3ikhj#{~>JWp{kph!*!EkIjsJko$IYyNCG$+Fw{JNQ59Y|^$Ss7U`NOG?p z2v{2)PPa2c4+qGr9^Y=yaMIr%FR!}6>!+-o7dFrQE~lCn0pKUkQ#)Zjdc5ju}^qs9uZ-H~>^<_EYXqzHT*#PK7qvB}WBq^>Bqzz^ELInW)2`svUgx+i6Uk-!`xEPzjOFTkMkP0<`YL#|CS+T8#ym zg>$qbSuXfiSyFM|!}C*1vZF4Z%v!hZx6v8j$ccyNo>_3?O)j`|;!Jm8i{AW9hBLQF zFV&~3i}=b$&M7j6|4MapvF2T-N-A5U6TCz>8nW9C$h3TQ~#^(-aH%k_Dk%WqDJ^63q;r;G!u`Dbe z1HVJ+l+BNkd!4vbgy<<^5>D8YN}JQOTuG)Z@$ns*TkPe7FSexusDr}6Vhr8obi`+c zqSL5z013r|B#B+r9FEa&+WTK(r~L17)=T!~QGy&_KI5#Q12FEU(PIHg45HN8ft%R} zJ-kK>y=0~Z>-)tJiOhM-JwoOSZMNN!8gz`XBA7~p3P^qN3j&c&Ms`l<;g@j*$06LE z)^!E2cgA=@k0C!rpMlqaQURHG@uApkOis*rFgeabk{`iR2W|msx)|59VTS>Mvr1z> zq*tU2RB4*4@+xFrvWVM`x~;p(q)hlw+D0T0+#k=uKLib>`suUl_Xdi^i(zvK@=+yc zE0}THsyOV@Q4CN}o2dvFamnj+638&Q8m^GdcV}JAX5o4b;<$lk*UcvQcthCR!#`Eh!ixro9 zj;+DU!WKPv%gjAVW02M^3|U&FUboqAV06Vvq&{jlJ?_r-&hGS94vlbnZtQ zf$u|h9Qlu!m2q*m^zu@mUEMsS)$gkae+~ZQ3c`E-->e_}HUH1g+Wwf2`@WaI082%W)gZVe{e<~dQZxjCP@qe7KiT5`X{%rN{BmTzd|N0~T ld|&=?#0t^hjQIEM`QPFp-ybG`fIz>$GQdDUM#+AE`hQfZ0xJLj diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_ascend.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_ascend.py deleted file mode 100644 index a5d0054..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_ascend.py +++ /dev/null @@ -1,77 +0,0 @@ -"""Task30 Huawei Ascend: one row owner, one selected read and one write. - -作者: 阿念 为 甄静 · 队长 孙蓓 · 2026 GuanghuLab -基础参考: r16 ascend 特化版(队友) - -华为 NPU 上 r13/r14 triplet-ticket variants failed sparse correctness。 -本版回到 r10 的 selection rule,加 source_plane 一次性选 plane。 -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -def launch_plan(rows, cols): - if not 0 < rows <= 65535 or not 0 < cols <= 2**20: - raise ValueError('flattened row/tile ownership requires valid rows and columns') - block = min(1024, 1 << max(5, (cols - 1).bit_length())) - tiles = (cols + block - 1) // block - if rows * tiles > 65535: - raise ValueError('flattened launch exceeds the grid-x limit') - return rows * tiles, block, tiles - - -@triton.jit -def _interleave_tiles( - input_ptr, out_ptr, - D: tl.constexpr, PLANE_STRIDE: tl.constexpr, - RS: tl.constexpr, CS: tl.constexpr, - H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, TILES: tl.constexpr, -): - program = tl.program_id(0) - row = program // TILES - tile = program - row * TILES - col = tile * BLOCK + tl.arange(0, BLOCK) - valid = col < D - - phase = col % 3 - from_b = (phase == 1) & (col < H_END) - from_c = (phase == 2) & (col < W_END) - source_plane = tl.where(from_b, 1, tl.where(from_c, 2, 0)) - - source_offset = source_plane * PLANE_STRIDE + row * RS + col * CS - value = tl.load(input_ptr + source_offset, mask=valid, other=0) - tl.store(out_ptr + row * D + col, value, mask=valid) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - - _, rows, d = x.shape - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - - grid, block, tiles = launch_plan(rows, d) - with torch.get_device_module(x.device).device(x.device): - _interleave_tiles[(grid,)]( - x, output, - D=d, PLANE_STRIDE=x.stride(0), - RS=x.stride(1), CS=x.stride(2), - H_END=min(d, int(mrope_section[1]) * 3), - W_END=min(d, int(mrope_section[2]) * 3), - BLOCK=block, TILES=tiles, num_warps=4, - ) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_enflame.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_enflame.py deleted file mode 100644 index 6e3341a..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_enflame.py +++ /dev/null @@ -1,72 +0,0 @@ -"""Task30 燧原 Enflame: per-row 2D grid with re-derived selection. - -作者: 阿念 为 甄静 · 队长 孙蓓 · 2026 GuanghuLab -基础参考: r16 enflame 特化版 -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -@triton.jit -def _interleave_tiles( - input_ptr, output_ptr, - N: tl.constexpr, D: tl.constexpr, - PLANE_STRIDE: tl.constexpr, ROW_STRIDE: tl.constexpr, - COL_STRIDE: tl.constexpr, H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, -): - if input_ptr.dtype.element_ty.primitive_bitwidth == 16: - input_ptr = input_ptr.to(tl.pointer_type(tl.int16)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int16)) - elif input_ptr.dtype.element_ty.primitive_bitwidth == 32: - input_ptr = input_ptr.to(tl.pointer_type(tl.int32)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int32)) - - row = tl.program_id(1).to(tl.int64) - col = (tl.program_id(0).to(tl.int64) * BLOCK - + tl.arange(0, BLOCK).to(tl.int64)) - valid = col < D - safe_col = tl.where(valid, col, 0) - phase = safe_col % 3 - from_b = (phase == 1) & (safe_col < H_END) - from_c = (phase == 2) & (safe_col < W_END) - base = row * ROW_STRIDE + safe_col * COL_STRIDE - - a = tl.load(input_ptr + base, valid & ~from_b & ~from_c, other=0) - b = tl.load(input_ptr + PLANE_STRIDE + base, valid & from_b, other=0) - c = tl.load(input_ptr + 2 * PLANE_STRIDE + base, valid & from_c, other=0) - value = tl.where(from_b, b, tl.where(from_c, c, a)) - tl.store(output_ptr + row * D + col, value, valid) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - _, rows, d = x.shape - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - - block = min(1024, 1 << max(5, (d - 1).bit_length())) - grid = (triton.cdiv(d, block), rows) - with torch.get_device_module(x.device).device(x.device): - _interleave_tiles[grid]( - x, output, N=rows * d, D=d, - PLANE_STRIDE=x.stride(0), ROW_STRIDE=x.stride(1), - COL_STRIDE=x.stride(2), - H_END=min(d, int(mrope_section[1]) * 3), - W_END=min(d, int(mrope_section[2]) * 3), - BLOCK=block, num_warps=4, enable_fp_fusion=False, - ) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_hygon.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_hygon.py deleted file mode 100644 index 06470b6..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_hygon.py +++ /dev/null @@ -1,110 +0,0 @@ -"""Task30 海光 Hygon: one output owner, one selected read and one write per element. - -作者: 阿念 为 甄静 · 队长 孙蓓 · 2026 GuanghuLab -基础参考: r16 hygon 特化版(分段防 huge stride) -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -def launch_plan(total): - if not 0 < total <= 2**26: - raise ValueError('local work must contain 1..2**26 elements') - block = 256 if total <= 4096 else 1024 - grid = min(65535, (total + block - 1) // block) - rounds = (total + grid * block - 1) // (grid * block) - return grid, block, rounds - - -def segment_ranges(rows, d, row_stride, col_stride, - limit=2**31-1, max_elements=2**26): - if min(rows, d, row_stride, col_stride) < 0 or limit < 1: - raise ValueError('nonnegative shapes/strides and positive limit required') - if not 0 < max_elements <= 2**26: - raise ValueError('invalid local work cap') - if not rows or not d: - return - cols_per = min(d, max_elements, limit // max(col_stride, 1) + 1) - for c0 in range(0, d, cols_per): - c1 = min(c0 + cols_per, d) - width = c1 - c0 - rows_input = (limit - (width - 1) * col_stride) // max(row_stride, 1) + 1 - rows_output = (limit - (width - 1)) // max(d, 1) + 1 - rows_per = min(rows, rows_input, rows_output, max_elements // width) - for r0 in range(0, rows, rows_per): - yield r0, min(r0 + rows_per, rows), c0, c1 - - -def selection_metadata(c0, cols, section): - return (c0 % 3, - min(cols, max(0, 3 * section[1] - c0)), - min(cols, max(0, 3 * section[2] - c0))) - - -@triton.jit -def _interleave_tiles( - a_ptr, b_ptr, c_ptr, out_ptr, - N: tl.constexpr, C: tl.constexpr, - RS: tl.constexpr, CS: tl.constexpr, OS: tl.constexpr, - PHASE: tl.constexpr, H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, GRID: tl.constexpr, ROUNDS: tl.constexpr, -): - pid = tl.program_id(0) - lane = tl.arange(0, BLOCK) - for batch in range(0, ROUNDS): - index = (batch * GRID + pid) * BLOCK + lane - valid = index < N - safe = tl.where(valid, index, 0) - row = safe // C - col = safe % C - phase = (col + PHASE) % 3 - from_b = (phase == 1) & (col < H_END) - from_c = (phase == 2) & (col < W_END) - offset = row * RS + col * CS - a = tl.load(a_ptr + offset, valid & ~from_b & ~from_c, other=0) - b = tl.load(b_ptr + offset, valid & from_b, other=0) - c = tl.load(c_ptr + offset, valid & from_c, other=0) - value = tl.where(from_b, b, tl.where(from_c, c, a)) - tl.store(out_ptr + row * OS + col, value, valid) - - -def launch_segments(x, output, section): - _, rows, d = x.shape - rs, cs = x.stride(1), x.stride(2) - for r0, r1, c0, c1 in segment_ranges(rows, d, rs, cs): - height, cols = r1 - r0, c1 - c0 - grid, block, rounds = launch_plan(height * cols) - phase, h_end, w_end = selection_metadata(c0, cols, section) - _interleave_tiles[(grid, 1, 1)]( - x[0, r0:r1, c0:c1], x[1, r0:r1, c0:c1], - x[2, r0:r1, c0:c1], output[r0:r1, c0:c1], - N=height * cols, C=cols, - RS=rs if height > 1 else 0, CS=cs if cols > 1 else 0, - OS=d if height > 1 else 0, - PHASE=phase, H_END=h_end, W_END=w_end, - BLOCK=block, GRID=grid, ROUNDS=rounds, num_warps=4) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - section = tuple(int(v) for v in mrope_section) - _, rows, d = x.shape - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - module = torch.get_device_module(x.device) - with module.device(x.device): - launch_segments(x, output, section) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_iluvatar.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_iluvatar.py deleted file mode 100644 index 3d59210..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_iluvatar.py +++ /dev/null @@ -1,72 +0,0 @@ -"""Task30 天数智芯 Iluvatar: per-row 2D grid with re-derived selection. - -作者: 阿念 为 甄静 · 队长 孙蓓 · 2026 GuanghuLab -基础参考: r16 iluvatar 特化版 -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -@triton.jit -def _interleave_tiles( - input_ptr, output_ptr, - N: tl.constexpr, D: tl.constexpr, - PLANE_STRIDE: tl.constexpr, ROW_STRIDE: tl.constexpr, - COL_STRIDE: tl.constexpr, H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, -): - if input_ptr.dtype.element_ty.primitive_bitwidth == 16: - input_ptr = input_ptr.to(tl.pointer_type(tl.int16)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int16)) - elif input_ptr.dtype.element_ty.primitive_bitwidth == 32: - input_ptr = input_ptr.to(tl.pointer_type(tl.int32)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int32)) - - row = tl.program_id(1).to(tl.int64) - col = (tl.program_id(0).to(tl.int64) * BLOCK - + tl.arange(0, BLOCK).to(tl.int64)) - valid = col < D - safe_col = tl.where(valid, col, 0) - phase = safe_col % 3 - from_b = (phase == 1) & (safe_col < H_END) - from_c = (phase == 2) & (safe_col < W_END) - base = row * ROW_STRIDE + safe_col * COL_STRIDE - - a = tl.load(input_ptr + base, valid & ~from_b & ~from_c, other=0) - b = tl.load(input_ptr + PLANE_STRIDE + base, valid & from_b, other=0) - c = tl.load(input_ptr + 2 * PLANE_STRIDE + base, valid & from_c, other=0) - value = tl.where(from_b, b, tl.where(from_c, c, a)) - tl.store(output_ptr + row * D + col, value, valid) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - _, rows, d = x.shape - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - - block = min(1024, 1 << max(5, (d - 1).bit_length())) - grid = (triton.cdiv(d, block), rows) - with torch.get_device_module(x.device).device(x.device): - _interleave_tiles[grid]( - x, output, N=rows * d, D=d, - PLANE_STRIDE=x.stride(0), ROW_STRIDE=x.stride(1), - COL_STRIDE=x.stride(2), - H_END=min(d, int(mrope_section[1]) * 3), - W_END=min(d, int(mrope_section[2]) * 3), - BLOCK=block, num_warps=4, enable_fp_fusion=False, - ) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_kunlunxin.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_kunlunxin.py deleted file mode 100644 index b1e73f4..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_kunlunxin.py +++ /dev/null @@ -1,72 +0,0 @@ -"""Task30 昆仑芯 Kunlunxin: per-row 2D grid with re-derived selection. - -作者: 阿念 为 甄静 · 队长 孙蓓 · 2026 GuanghuLab -基础参考: r16 kunlunxin 特化版 -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -@triton.jit -def _interleave_tiles( - input_ptr, output_ptr, - N: tl.constexpr, D: tl.constexpr, - PLANE_STRIDE: tl.constexpr, ROW_STRIDE: tl.constexpr, - COL_STRIDE: tl.constexpr, H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, -): - if input_ptr.dtype.element_ty.primitive_bitwidth == 16: - input_ptr = input_ptr.to(tl.pointer_type(tl.int16)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int16)) - elif input_ptr.dtype.element_ty.primitive_bitwidth == 32: - input_ptr = input_ptr.to(tl.pointer_type(tl.int32)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int32)) - - row = tl.program_id(1).to(tl.int64) - col = (tl.program_id(0).to(tl.int64) * BLOCK - + tl.arange(0, BLOCK).to(tl.int64)) - valid = col < D - safe_col = tl.where(valid, col, 0) - phase = safe_col % 3 - from_b = (phase == 1) & (safe_col < H_END) - from_c = (phase == 2) & (safe_col < W_END) - base = row * ROW_STRIDE + safe_col * COL_STRIDE - - a = tl.load(input_ptr + base, valid & ~from_b & ~from_c, other=0) - b = tl.load(input_ptr + PLANE_STRIDE + base, valid & from_b, other=0) - c = tl.load(input_ptr + 2 * PLANE_STRIDE + base, valid & from_c, other=0) - value = tl.where(from_b, b, tl.where(from_c, c, a)) - tl.store(output_ptr + row * D + col, value, valid) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - _, rows, d = x.shape - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - - block = min(1024, 1 << max(5, (d - 1).bit_length())) - grid = (triton.cdiv(d, block), rows) - with torch.get_device_module(x.device).device(x.device): - _interleave_tiles[grid]( - x, output, N=rows * d, D=d, - PLANE_STRIDE=x.stride(0), ROW_STRIDE=x.stride(1), - COL_STRIDE=x.stride(2), - H_END=min(d, int(mrope_section[1]) * 3), - W_END=min(d, int(mrope_section[2]) * 3), - BLOCK=block, num_warps=4, enable_fp_fusion=False, - ) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_metax.py b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_metax.py deleted file mode 100644 index 4ac100e..0000000 --- a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_metax.py +++ /dev/null @@ -1,72 +0,0 @@ -"""Task30 沐曦 Metax: per-row 2D grid with re-derived selection. - -作者: 阿念 为 甄静 · 队长 孙蓓 · 2026 GuanghuLab -基础参考: r16 metax 特化版 -""" -from numbers import Integral - -import torch -import triton -import triton.language as tl - - -@triton.jit -def _interleave_tiles( - input_ptr, output_ptr, - N: tl.constexpr, D: tl.constexpr, - PLANE_STRIDE: tl.constexpr, ROW_STRIDE: tl.constexpr, - COL_STRIDE: tl.constexpr, H_END: tl.constexpr, W_END: tl.constexpr, - BLOCK: tl.constexpr, -): - if input_ptr.dtype.element_ty.primitive_bitwidth == 16: - input_ptr = input_ptr.to(tl.pointer_type(tl.int16)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int16)) - elif input_ptr.dtype.element_ty.primitive_bitwidth == 32: - input_ptr = input_ptr.to(tl.pointer_type(tl.int32)) - output_ptr = output_ptr.to(tl.pointer_type(tl.int32)) - - row = tl.program_id(1).to(tl.int64) - col = (tl.program_id(0).to(tl.int64) * BLOCK - + tl.arange(0, BLOCK).to(tl.int64)) - valid = col < D - safe_col = tl.where(valid, col, 0) - phase = safe_col % 3 - from_b = (phase == 1) & (safe_col < H_END) - from_c = (phase == 2) & (safe_col < W_END) - base = row * ROW_STRIDE + safe_col * COL_STRIDE - - a = tl.load(input_ptr + base, valid & ~from_b & ~from_c, other=0) - b = tl.load(input_ptr + PLANE_STRIDE + base, valid & from_b, other=0) - c = tl.load(input_ptr + 2 * PLANE_STRIDE + base, valid & from_c, other=0) - value = tl.where(from_b, b, tl.where(from_c, c, a)) - tl.store(output_ptr + row * D + col, value, valid) - - -def interleaved_rope(x, mrope_section): - if x.ndim != 3 or x.shape[0] != 3: - raise ValueError('x must have shape [3, S, D]') - if (not isinstance(mrope_section, (list, tuple)) or len(mrope_section) != 3 - or any(not isinstance(v, Integral) or v < 0 for v in mrope_section)): - raise ValueError('mrope_section must contain three nonnegative integers') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - _, rows, d = x.shape - output = torch.empty((rows, d), dtype=x.dtype, device=x.device) - if not rows or not d: - return output - - block = min(1024, 1 << max(5, (d - 1).bit_length())) - grid = (triton.cdiv(d, block), rows) - with torch.get_device_module(x.device).device(x.device): - _interleave_tiles[grid]( - x, output, N=rows * d, D=d, - PLANE_STRIDE=x.stride(0), ROW_STRIDE=x.stride(1), - COL_STRIDE=x.stride(2), - H_END=min(d, int(mrope_section[1]) * 3), - W_END=min(d, int(mrope_section[2]) * 3), - BLOCK=block, num_warps=4, enable_fp_fusion=False, - ) - return output - - -reference = interleaved_rope diff --git a/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_min.zip b/zz-flagos-s2-2026/d245/src/Task30_r16_zhizhi/interleaved_rope_min.zip deleted file mode 100644 index 233e2f853049d633e168c8f69ec6db618bac97b0..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1568 zcmZ|Qdpy$%6aetwX1K|SVaanID=`xxd6WuCBuVBq`iy7#~Px%YEE=X1_K=lpd(-{WW}1d#y%;5T5YNzGHAUGvFL z7y#zC0{{$A1cE{@`A6XWeWLySun}S5{<`6DXPvhJLiMI3-1DuBiro%?AOa8o{Jo^o z#XSsWBSRR+Uff0TGsf$TBtOjb%_yiq;beBaokI`{KJ(uzzgMWQ&-m(O85CbR13Hp% zg6Nn^B)U1tLXM>#-$Y?{uffd?T2|ui%VnwfINk9Weq=yeS{f_&L#L}e3B4yq3sZ5c z>uR<`65{|>c*#3xzv~x@P=kGVsOY5rUBr?WdG;j4qDfN!ZYKfa1^k-u&RYmidPzOIoT*Le5vsdLJHe-BX}hH zJB=Q*`HZ&l=U{23Hp~%qu6vT<%o3*r_mn>2v=rv{zgg9t)UCdm6+NIFWyD85Ha0aL zG|fWC7^y6VhGZepF1dk;$j2ocZC;#&Sd7ib9&StzQ^nL>GIu<2xRn-k!EOFUJ*O>{ zZ+@BG92IgzL72RId~AGZU7Fl}3ld0xq39wSa!-=fy#o{OMGESixA!nNH?>z0YoFIx z18BVjF6=mevu1T>lGXPquiri0o*qAj8R3l*vhno8WBlQRxbLQ2oDrrhyC zz*tEsjYT%U;){#V4qu-c?JyG$U44^>?&!t9Xf2rQcALuie)d)flSI zcD5S%={a2P*(Bfh10NMN)!3AF@_Vk%S`*xdQ4xRKSECRvG($aiR4pO5b2>N=czW1WcB0+(<*vGClKPRU$eEHoyX< zOG*g%`OBgvMWL^5cS#kgnpzjwqQb7#foBEocfIx$M42;oeakBt7!FE&LD6p;5Zs$5 zi8pEVd|=l0&4Xq4tr}Q;EP6T}$pv$lN!90={pv_(?+9tL-hnLNIi;Olv@@ObFdtUQ z-AV<_ka06EEqjax*MBa1tIT9DJHk>WK zxtnX#cSF>lUZ+}JTnp+QU9jyJvgKMlqNJC6u@6K^DO(%9MhI*<5YIQ z6}+=b0|Lqw8wOjosEL519SE!}^nYaU+@c%+wmb$Zl7Dso3Ape-fd30`M>}DWUn8Kc Nz--mptv>(&e*yoN;otxO diff --git a/zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.py b/zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.py deleted file mode 100644 index 4a4ae39..0000000 --- a/zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.py +++ /dev/null @@ -1,98 +0,0 @@ -"""Task35 rotary_embedding: fixed pair ownership and fused float32 rotation. - -作者: 阿念 (anien@guanghulab.local) 为 甄静(8592_apivqhj)· 队长 孙蓓 -版权: 2026 GuanghuLab -基础参考: r6-四Head共享角度流水线(队友共享) - -改编: -- 改 docstring 与 author -- 加 `enable_fp_fusion=False`(国产 NPU 套路,从 r16 M-RoPE 学) -- 函数末尾保留 `reference = rotary_embedding`(平台 import 入口) -""" -import torch -import triton -import triton.language as tl - - -@triton.jit -def _rotary_pairs( - x_ptr, cos_ptr, sin_ptr, out_ptr, - H: tl.constexpr, D: tl.constexpr, HALF: tl.constexpr, - XTS: tl.constexpr, XHS: tl.constexpr, XDS: tl.constexpr, - CTS: tl.constexpr, CDS: tl.constexpr, - STS: tl.constexpr, SDS: tl.constexpr, - OTS: tl.constexpr, OHS: tl.constexpr, ODS: tl.constexpr, - BLOCK: tl.constexpr, -): - row = tl.program_id(0) - tile = tl.program_id(1) - token = row // H - head = row - token * H - pair = tile * BLOCK + tl.arange(0, BLOCK) - valid = pair < HALF - - x_base = token * XTS + head * XHS - x_even = tl.load(x_ptr + x_base + (2 * pair) * XDS, - mask=valid, other=0.0).to(tl.float32) - x_odd = tl.load(x_ptr + x_base + (2 * pair + 1) * XDS, - mask=valid, other=0.0).to(tl.float32) - cosine = tl.load(cos_ptr + token * CTS + pair * CDS, - mask=valid, other=0.0).to(tl.float32) - sine = tl.load(sin_ptr + token * STS + pair * SDS, - mask=valid, other=0.0).to(tl.float32) - - out_even = x_even * cosine - x_odd * sine - out_odd = x_even * sine + x_odd * cosine - out_base = token * OTS + head * OHS - tl.store(out_ptr + out_base + (2 * pair) * ODS, out_even, mask=valid) - tl.store(out_ptr + out_base + (2 * pair + 1) * ODS, out_odd, mask=valid) - - -def _block_size(half): - if half <= 32: - return 32 - if half <= 64: - return 64 - if half <= 128: - return 128 - return 256 - - -def rotary_embedding(x, cos, sin, interleaved): - if x.ndim != 3: - raise ValueError('x must have shape [T, H, D]') - if not isinstance(interleaved, bool) or interleaved: - raise ValueError('Task35 requires interleaved=False') - if x.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - tokens, heads, width = x.shape - if width % 2: - raise ValueError('the last dimension must be even') - half = width // 2 - if cos.ndim != 2 or tuple(cos.shape) != (tokens, half): - raise ValueError('cos must have shape [T, D//2]') - if sin.ndim != 2 or tuple(sin.shape) != (tokens, half): - raise ValueError('sin must have shape [T, D//2]') - if cos.device != x.device or sin.device != x.device: - raise ValueError('x, cos and sin must be on the same device') - - output = torch.empty(x.shape, dtype=x.dtype, device=x.device) - if x.numel() == 0: - return output - block = _block_size(half) - grid = (tokens * heads, triton.cdiv(half, block)) - module = torch.get_device_module(x.device) - with module.device(x.device): - _rotary_pairs[grid]( - x, cos, sin, output, - H=heads, D=width, HALF=half, - XTS=x.stride(0), XHS=x.stride(1), XDS=x.stride(2), - CTS=cos.stride(0), CDS=cos.stride(1), - STS=sin.stride(0), SDS=sin.stride(1), - OTS=output.stride(0), OHS=output.stride(1), ODS=output.stride(2), - BLOCK=block, num_warps=4, enable_fp_fusion=False, - ) - return output - - -reference = rotary_embedding diff --git a/zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.zip b/zz-flagos-s2-2026/d245/src/Task35_r6_zhizhi/rotary_embedding.zip deleted file mode 100644 index c9b9eea35385abdded0c6b5998e934b0c6097a0c..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 1614 zcmV-U2C?~2O9KQH0000802jO>T`YUK;*tdb0O}0@022Tl0CI11VRCt2Wo=?*WMpY> zXD)DgRa6N81D!#bYMnutYIS%E009K`0RR956aWAKrB_>T6Gs$&-(PVWsq$`coWzib zTC7ByKqHklh@=$~LT2n8+XL&}Wp>sveWI7BH;G6E+9E)-s(=SnHC2l!X=$YXmTVIK z1!pd{*K0RQ+q~F5x9^*C&YUxun3z~4@s*=keZb~*xLKydb;uP{&m#y`Kk2o-6cxAbp4l|ttb1BzfQLv?LPfzck{Qs zt!*1)rW;>ronu+Mx1R6)`Ax-|0{jMESdRtA10$;&KoZdkIjYhxzlTdSAT^&>8&MUH zj7M}&6F;Vx?DX!7ovp|4-r4sdefaJEv!9DQf7}2*I}h(poeR&NLi&#!8RWAUySF!Y z?>$L3|JvEUzjyl%T;jArISuL*+>xAPkre7A02CAV!|>a5%PMg z|AR9T2E)HfK4vK)O$sCi;aiq<(kL!7VR^IxRRb$48`~D(U$+_wUW9rW>n|}2^p`LZ z>Wi{ks$jrUJq%(&yAf&^^Ma)}&z#Q5l;`=?mAvWsrBTshUZk8CMx7VNTds^+u8g-_ z9<^K^CA2)=`K>d{3-9C=PDSZ>xQ2yARm8(4C+#Zp>={R?1oP>LZr13+D>P6Yu?ExA zu%v`7>}VsKG8U6YrEaCGLr!W!a72bC9LrDbnW7fz=pEuS8A*AaP*t)_-D@OPq)c=y z6vk1&ld?p{igpw`jE|SzvsJsGopC<`b{WlObVs@^uIN0C{Izk^xQYyRmS|D#&Xi`H zk_c^#Z03n$AVbeP6e50~JrMQqpxFK_pczziA7o}teil@I6z-(#!vB5H9HO~+8N`(= z;>tfD9u7}=${TIUY3e%JJNT4I`=nykx7L;D8!OWzeN*k*H_r9ba<-qAO+O*080Q$Z z&2&MFzH6?R@CFw9l#7{jI{zk8uSfdPkos^iOHalcE~qQjn0-j?7V#VAPO}EcUvR>O zqvc9J|D1}12Pn!L%^w>vnm?8|nk^q6u|k=Z`B|QuH^lOb!tSbJtp;@wSRg3(Dd|ve zz_?oqJl2NSkaGryFwD#c#82oc&O>e&y3kHyfvh?d;ueW0Tv)|9jE99QYb>9 zO$AY(QCx8HXy=kZuyzlRAQbUoRqa3^^*Z$_C)hMlBXw+24^P%(;}tT~7_%$;7C+XQ zCtC6flvIX#^tAYOC?9lOyo}%z%rF)MW8x!1C%i^MwzL6MP2rk|IRE=Cg+=byu`IJJ zlE|lWFljE1RM>sKgPtAF2K~m?dvSWYoT)dKdK{`0{})vBI}EBklrFdo-z&a^m%yW% z1KV2N5A|B@qhrc3`?4T0X;aWa3WKp2B|=Rud|;JmI}+DyQ}7~qvIAWNmY+rA>h_3a z8a7DU)VCdQU6>iU-Wo$m)u2JNk%^I2M|-tHFVQay zb%Z7rRn4!eHTE#yHHN*WHAX%#7*HR2EMAbrF4{vIXlN4Y6ArJaCD*WDbX9-rSE8%< z4;kXy15+v=;~qX>9Q8yT2xcAiKpY6l&agYKW>?*dK>%FPSF4*J)^ zdIeBR0RkQa6aWAK2mlwnB3&$dx#E%q008O@000vJ7ytkO00031005+c00000a&L5D za(Q26ZDM6)WNB_^E^v8OR0RM7ok5stb$AN^0R;5{000CO0000`O9ci10000100975 M0002r1poj504&GlC;$Ke diff --git a/zz-flagos-s2-2026/d245/src/silu_and_mul_masked.py b/zz-flagos-s2-2026/d245/src/silu_and_mul_masked.py deleted file mode 100644 index 5652255..0000000 --- a/zz-flagos-s2-2026/d245/src/silu_and_mul_masked.py +++ /dev/null @@ -1,164 +0,0 @@ -# Copyright 2026 FlagOS Contributors / GuanghuLab -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Optimized `reference` for FlagOS Open Computing Challenge S2 - Track 1 -# Task 39 (silu_and_mul_masked / DeepGEMM-style MoE). -# -# !! IMPORTANT: The function name MUST be exactly `reference` per the spec. !! -# The platform's grader imports the function by this name. -# -# Spec (from flagos.net): -# Input input: [E, T, H] bfloat16 -# masked_m: [E] int (per-expert valid row count) -# Output out: [E, T, H/2] bfloat16 (initialized to zero) -# For each expert e, let n = masked_m[e], half = H // 2: -# gate = input[e, :n, :half] (cast to float32) -# up = input[e, :n, half:] (cast to float32) -# out[e, :n] = (gate * sigmoid(gate)) * up (cast back to bfloat16) -# Rows [e, masked_m[e]:T] are padding — output values not checked. -# -# Reference (the slow one we beat): -# import torch -# def reference(input, masked_m): -# E, T, H = input.shape -# half = H // 2 -# out = torch.zeros(E, T, half, dtype=torch.bfloat16, device=input.device) -# for e in range(E): -# n = int(masked_m[e].item()) # <-- host sync, very slow -# if n <= 0: continue -# gate = input[e, :n, :half].float() -# up = input[e, :n, half:].float() -# val = gate * torch.sigmoid(gate) * up -# out[e, :n] = val.to(torch.bfloat16) -# return out -# -# Optimization strategy: -# 1. Single Triton kernel launch — no Python loop, no host sync -# 2. Load `masked_m` in-kernel (one int per program) — no .item() -# 3. Vectorize over (expert, row, hidden_col_half) with autotune -# 4. Compute in float32 for accuracy (matches reference math) -# 5. Cast back to bf16 on store - -import torch -import triton -import triton.language as tl - - -# --------------------------------------------------------------------------- -# Autotune -# --------------------------------------------------------------------------- -_AUTOTUNE_CONFIGS = [ - triton.Config({"BLOCK_M": 1, "BLOCK_N": 128}, num_warps=4, num_stages=2), - triton.Config({"BLOCK_M": 1, "BLOCK_N": 256}, num_warps=4, num_stages=2), - triton.Config({"BLOCK_M": 1, "BLOCK_N": 512}, num_warps=4, num_stages=2), - triton.Config({"BLOCK_M": 1, "BLOCK_N": 1024}, num_warps=8, num_stages=2), - triton.Config({"BLOCK_M": 2, "BLOCK_N": 256}, num_warps=4, num_stages=2), - triton.Config({"BLOCK_M": 2, "BLOCK_N": 512}, num_warps=8, num_stages=2), - triton.Config({"BLOCK_M": 4, "BLOCK_N": 256}, num_warps=4, num_stages=2), - triton.Config({"BLOCK_M": 4, "BLOCK_N": 512}, num_warps=8, num_stages=2), - triton.Config({"BLOCK_M": 8, "BLOCK_N": 128}, num_warps=4, num_stages=2), - triton.Config({"BLOCK_M": 8, "BLOCK_N": 256}, num_warps=8, num_stages=2), -] - - -# --------------------------------------------------------------------------- -# Triton kernel -# --------------------------------------------------------------------------- -@triton.autotune(configs=_AUTOTUNE_CONFIGS, key=["E", "HIDDEN_HALF"]) -@triton.jit -def reference_kernel( - in_ptr, masked_m_ptr, out_ptr, - E, - s_e, s_t, s_h, # strides for input [E, T, H] - so_e, so_t, so_h, # strides for out [E, T, H/2] - HIDDEN_HALF: tl.constexpr, - BLOCK_M: tl.constexpr, - BLOCK_N: tl.constexpr, -): - pid_e = tl.program_id(0) # expert - pid_m = tl.program_id(1) # row block (over T) - pid_n = tl.program_id(2) # col block (over H/2) - - # Load this expert's valid row count (no host sync — pure device load) - valid_m = tl.load(masked_m_ptr + pid_e) - - # Row indices in this program - row_idx = pid_m * BLOCK_M + tl.arange(0, BLOCK_M) # (BLOCK_M,) - col_idx = pid_n * BLOCK_N + tl.arange(0, BLOCK_N) # (BLOCK_N,) - - # Row mask: only valid rows compute - row_mask = row_idx < valid_m # (BLOCK_M,) - col_mask = col_idx < HIDDEN_HALF # (BLOCK_N,) - - # H is split in half: gate lives in [0, H/2), up in [H/2, H) - full_mask = row_mask[:, None] & col_mask[None, :] # (BLOCK_M, BLOCK_N) - - # 2D offsets into the input's last dim - base_in = pid_e * s_e + row_idx[:, None] * s_t - base_out = pid_e * so_e + row_idx[:, None] * so_t - - # gate and up offsets: same row, gate uses col_idx, up uses col_idx + H/2 - gate_offs = base_in + col_idx[None, :] * s_h - up_offs = base_in + (col_idx[None, :] + HIDDEN_HALF) * s_h - - g = tl.load(in_ptr + gate_offs, mask=full_mask, other=0.0) - u = tl.load(in_ptr + up_offs, mask=full_mask, other=0.0) - - # Compute in fp32 for accuracy (matches reference math) - g32 = g.to(tl.float32) - u32 = u.to(tl.float32) - y32 = g32 * tl.sigmoid(g32) * u32 - y = y32.to(tl.bfloat16) - - # Store to out - out_offs = base_out + col_idx[None, :] * so_h - tl.store(out_ptr + out_offs, y, mask=full_mask) - - -# --------------------------------------------------------------------------- -# Public API (function name MUST be `reference`) -# --------------------------------------------------------------------------- -def reference(input: torch.Tensor, masked_m: torch.Tensor) -> torch.Tensor: - """ - DeepGEMM-style SiLU(gate) * up with per-expert row mask. - - Args: - input: [E, T, H] bfloat16 (gate in [0:H/2], up in [H/2:H]) - masked_m: [E] int (valid row count per expert) - - Returns: - out: [E, T, H/2] bfloat16, zero-initialized - Padded rows (j >= masked_m[e]) are left as 0. - """ - assert input.dtype == torch.bfloat16, f"expected bf16, got {input.dtype}" - assert input.dim() == 3, f"expected 3D tensor, got {input.dim()}D" - E, T, H = input.shape - assert H % 2 == 0, f"H must be even, got {H}" - half = H // 2 - - assert masked_m.shape == (E,), f"masked_m shape mismatch: {masked_m.shape}" - masked_m_i32 = masked_m.to(torch.int32).contiguous() - - out = torch.zeros((E, T, half), dtype=torch.bfloat16, device=input.device) - - grid = lambda meta: ( - E, - triton.cdiv(T, meta["BLOCK_M"]), - triton.cdiv(half, meta["BLOCK_N"]), - ) - - with torch.cuda.device(input.device): - reference_kernel[grid]( - input, masked_m_i32, out, - E, - input.stride(0), input.stride(1), input.stride(2), - out.stride(0), out.stride(1), out.stride(2), - HIDDEN_HALF=half, - ) - - return out diff --git a/zz-flagos-s2-2026/d245/tests/test_correctness.py b/zz-flagos-s2-2026/d245/tests/test_correctness.py deleted file mode 100644 index 36bf717..0000000 --- a/zz-flagos-s2-2026/d245/tests/test_correctness.py +++ /dev/null @@ -1,88 +0,0 @@ -"""Correctness test for all v2 operators. - -Run with: python -m pytest tests/test_correctness.py -v -or: python tests/test_correctness.py -""" -import torch - - -def test_gelu_and_mul_v2_correctness(): - import sys - sys.path.insert(0, "src") - from flag_gems.fused.gelu_and_mul import gelu_and_mul as baseline - from flag_gems_local.fused.gelu_and_mul_v2 import gelu_and_mul as v2 - - torch.manual_seed(42) - for shape in [(1024, 1024), (4096, 4096), (2048, 11008)]: - for dtype in [torch.float16, torch.bfloat16, torch.float32]: - x = torch.randn(shape, dtype=dtype, device="cuda") - y = torch.randn(shape, dtype=dtype, device="cuda") - out_b = baseline(x, y) - out_v = v2(x, y) - abs_diff = (out_b - out_v).abs().max().item() - assert abs_diff < 1e-2, ( - f"gelu_and_mul v2 diverges: shape={shape} dtype={dtype} abs_diff={abs_diff}" - ) - print(f"gelu_and_mul shape={shape} dtype={dtype} abs_diff={abs_diff:.2e} ✓") - - -def test_rotary_embedding_v2_correctness(): - import sys - sys.path.insert(0, "src") - from flag_gems.fused.rotary_embedding import apply_rotary_pos_emb as baseline - from flag_gems_local.fused.rotary_embedding_v2 import apply_rotary_pos_emb as v2 - - torch.manual_seed(42) - for B, S, H, D in [(1, 1024, 32, 128), (2, 2048, 16, 64), (4, 512, 8, 128)]: - for interleaved in [False, True]: - q = torch.randn(B, S, H, D, dtype=torch.float16, device="cuda") - k = torch.randn(B, S, H, D, dtype=torch.float16, device="cuda") - cos = torch.randn(S, D // 2, dtype=torch.float16, device="cuda") - sin = torch.randn(S, D // 2, dtype=torch.float16, device="cuda") - qe_b, ke_b = baseline(q, k, cos, sin, rotary_interleaved=interleaved) - qe_v, ke_v = v2(q, k, cos, sin, rotary_interleaved=interleaved) - abs_diff_q = (qe_b - qe_v).abs().max().item() - abs_diff_k = (ke_b - ke_v).abs().max().item() - assert abs_diff_q < 1e-2, f"q diverges: {abs_diff_q}" - assert abs_diff_k < 1e-2, f"k diverges: {abs_diff_k}" - print( - f"rotary_emb B={B} S={S} H={H} D={D} " - f"interleaved={interleaved} abs_diff=({abs_diff_q:.2e}, {abs_diff_k:.2e}) ✓" - ) - - -def test_fused_moe_router_v2_correctness(): - import sys - sys.path.insert(0, "src") - from flag_gems_local.ops.fused_moe_router_v2 import ( - bitonic_sort_topk, - prepare_moe_inputs_v2, - ) - - torch.manual_seed(42) - for num_tokens, top_k in [(64, 4), (256, 8), (1024, 4)]: - weights = torch.rand(num_tokens, top_k, device="cuda") - ids = torch.randint(0, 8, (num_tokens, top_k), device="cuda") - - sorted_w, sorted_ids = bitonic_sort_topk(weights, ids) - - # Verify: weights should be descending - for i in range(num_tokens): - for j in range(top_k - 1): - assert sorted_w[i, j] >= sorted_w[i, j + 1] - 1e-5, ( - f"Not sorted: token {i} weight[{j}]={sorted_w[i,j]} > weight[{j+1}]={sorted_w[i,j+1]}" - ) - - # Verify: expert ids are a permutation of original - orig_ids = torch.sort(ids, dim=1).values - new_ids = torch.sort(sorted_ids, dim=1).values - assert torch.equal(orig_ids, new_ids), "ids not a permutation!" - - print(f"fused_moe_router tokens={num_tokens} top_k={top_k} ✓") - - -if __name__ == "__main__": - test_gelu_and_mul_v2_correctness() - test_rotary_embedding_v2_correctness() - test_fused_moe_router_v2_correctness() - print("\nAll tests passed ✓") diff --git a/zz-flagos-s2-2026/d246/README.md b/zz-flagos-s2-2026/d246/README.md deleted file mode 100644 index 2746447..0000000 --- a/zz-flagos-s2-2026/d246/README.md +++ /dev/null @@ -1,78 +0,0 @@ -# FlagOS S2 赛季二 · 赛道一 · D246 续战(Task 38) - -> 之之 (8592_apivqhj) + 阿念 (Mavis, ICE-GL-AN-001, Code) · 协作产出 -> 比赛: FlagOS 开放计算全球大赛 S2 · Track 1 · SGLang 算子优化 -> 队伍: GuanghuLab(队长孙蓓) -> 第 3 批 17 道题 · 截止 2026-09-03 19:59 (剩 ~19h41m) - ---- - -## 战绩(D246 00:18 续) - -| Task | 算子 | 状态 | 平均加速比 | 详情 | -|---|---|---|---|---| -| 30 | interleaved_rope (M-RoPE) | ✅ 已传 (D245 23:16) | **35.79×** (5/8 跑通) | 天数 86.81× / 海光 42.86× / 国通A 26.91× / 沐曦 15.54× / 华为 6.84×;燧原/昆仑芯 Failed | -| 29 | gelu_and_mul | ✅ 已传 (D245 23:30) | 3.02× (7/8 跑通) | 天数 5.32× / 海光 4.03× / 国通A 3.73× / 国通B 3.25× / 沐曦 2.74× / 华为 1.09× / 燧原 0.96×;昆仑芯 Failed | -| 35 | rotary_embedding | ✅ 已传 (D246 00:04) | 5.15× (6/8 跑通) | 国际B 8.79× / 国际A 7.53× / 天数 6.72× / 海光 4.11× / 沐曦 3.41× / 昆仑芯 0.32×;华为 Failed, 燧原 "评测中" | -| **38** | **sigmoid_gate_topk_renorm** | ⏳ **待传(D246 00:17 准备就绪)** | - | 3-kernel 写法 (sigmoid+bias / 自写 topk / gather+renorm) | - -## D246 续战重点 - -### Task 38 · 选第 4 题 -- 找过 `embedding` (无) / `norm` (匹配到 Task 38) -- Task 38 = MoE 路由门算子 (DeepSeek-V2/V3 / Qwen3-MoE 风格) -- 当前第一名 sitraliqui 仅 9.05× — **大家都没做出来**,机会大 -- 没有 GEMM,纯 elementwise + reduce + gather → 跨芯片可行性高 - -### 选路决策 -之之 D246 00:12 决策:**"用 Triton 自写 topk(性能好)"** — 不走 torch.topk 兜底,要冲分 - -### 实施方案 -**3-kernel 写法**(不用 torch.topk,自写): -``` -K1: sel = sigmoid(routed) + bias (Triton elementwise) -K2: topk → indices (Triton, 迭代 tl.argmax K 次) -K3: gather + cat + sigmoid + renorm + scale + split (Triton, 1 token/program) -``` - -### 国产 NPU 套路全套 -- 指针 cast `to(tl.pointer_type(tl.int16))` 防 NaN -- stride / pid 用 int64 cast 防 overflow -- `with torch.get_device_module(x.device).device(x.device):` -- `enable_fp_fusion=False, num_warps=4` -- `1.0 / (1.0 + tl.exp(-x))` 不用 `tl.sigmoid` - -### 算法测试 8/8 pass -对比平台 reference vs 我们的算法拆分(torch 模拟 K1/K2/K3): -- T=1~128, N=8~256, S=1~2, k=2~8 -- 包括 DeepSeek-V3 风格 (T=64, N=256, k=8) -- **全部 pass**:indices match, routed close (atol=1e-3), shared close - -## 风险点 -1. **K2 `tl.argmax` 在国产 NPU 上可能不支持** — 降级方案:把 K2 换成 `torch.topk` 1 行 -2. **3 kernel 的 launch overhead** — 性能上限受限(估计 3-8×) -3. **中间 buffer `sel` 大小** — DeepSeek-V3 风格 (T=128, N=256) = 128KB - -## 预期 -- 跑通率:6-7/8 -- 平均加速比:3-8× -- 排名:跟 sitraliqui (9.05×) 接近或略低 - ---- - -## 目录结构 - -``` -d246/ -├── README.md (本文件) -├── sigmoid_gate_topk_renorm.py (主文件 · 3 Triton kernel + 国产 NPU 套路) -├── test_algorithm.py (8 case 算法层测试) -├── SUBMIT_CHECKLIST.md (提交清单 · 给之之看的) -└── results/ - └── sigmoid_gate_topk_renorm.zip (2.6 KB · 通用版,1 个文件) -``` - -## 学到的(本次) -1. **3-kernel 分拆 vs 单 kernel fused** — 跑通率 vs 性能上限的权衡,Task 38 选了前者 -2. **自写 topk 套路** — `tl.argmax` + `tl.where(mask, -inf, sel)` 迭代 K 次,K=8 内可行 -3. **算法层测试先行** — 8 case 在 CPU 跑过确认逻辑,GPU 编译是第二阶段 diff --git a/zz-flagos-s2-2026/d246/SUBMIT_CHECKLIST.md b/zz-flagos-s2-2026/d246/SUBMIT_CHECKLIST.md deleted file mode 100644 index 39962ce..0000000 --- a/zz-flagos-s2-2026/d246/SUBMIT_CHECKLIST.md +++ /dev/null @@ -1,92 +0,0 @@ -# Task 38 · sigmoid_gate_topk_renorm 提交清单 - -**生成时间**: 2026-09-03 00:17 (D246) -**提交次数**: 今日 1/30 (用了 1 次在 K35),还剩 29 次 -**截止**: 2026-09-03 19:59 (还剩 19h41m) - -## 准备好的文件 - -``` -/Users/zhizhi/Desktop/sigmoid_gate_topk_renorm.zip (2.6 KB) -└── sigmoid_gate_topk_renorm.py (7.6 KB · 通用版 · 无 7 芯片特化) -``` - -**主入口**: `reference = sigmoid_gate_topk_renorm` -**函数签名**: `def reference(logits, k, n_shared_experts, route_scale, global_scale, bias):` -**返回**: `(routed_w[T,k], indices[T,k] int32, shared_w[T,S])` - -## 实现方案 · 3-kernel 写法 - -``` -┌─ K1: _sigmoid_bias_kernel ─────────────────┐ -│ sel = sigmoid(routed) + bias │ -│ sel 存为 fp32 中间 buffer [T, N] │ -└────────────────────┬────────────────────────┘ - ↓ -┌─ K2: _topk_kernel (自写) ──────────────────┐ -│ 迭代 tl.argmax K 次, 每次选最大索引 │ -│ 把已选位置 mask 为 -inf, 防重复选 │ -│ 写 indices [T, k] int32 │ -└────────────────────┬────────────────────────┘ - ↓ -┌─ K3: _gate_finalize_kernel ────────────────┐ -│ gather(routed, indices) → routed_vals[k] │ -│ load shared_logits → shared_vals[S] │ -│ cat → sigmoid → /sum → *route_scale*gs │ -│ split → routed_w [T,k] / shared_w [T,S] │ -└─────────────────────────────────────────────┘ -``` - -## 国产 NPU 套路 (D245 验证过) - -| 套路 | 作用 | -|------|------| -| `with torch.get_device_module(x.device).device(x.device):` | 切设备 (K30/K35 跑通的必要条件) | -| `enable_fp_fusion=False, num_warps=4` | 国产 NPU 关 fp fusion | -| `pid.to(int64) + tl.arange(0, BLOCK).to(int64)` | 防 stride overflow | -| `input_ptr.to(tl.pointer_type(tl.int16))` | 防 NaN 在 load 时被吃 | -| `1.0 / (1.0 + tl.exp(-x))` (不用 `tl.sigmoid`) | tl.sigmoid 国产 NPU 不支持 | - -## 本地测试 (算法层, 8/8 pass) - -跑 `test_algorithm.py`,对比平台 reference vs 算法模拟(纯 torch 拆分 K1/K2/K3): - -``` -T=1 N=64 S=1 k=8 : indices match ✓ routed close ✓ shared close ✓ -T=1 N=64 S=1 k=4 : indices match ✓ routed close ✓ shared close ✓ -T=8 N=64 S=1 k=8 : indices match ✓ routed close ✓ shared close ✓ -T=32 N=128 S=2 k=8 : indices match ✓ routed close ✓ shared close ✓ -T=64 N=256 S=1 k=8 : indices match ✓ routed close ✓ shared close ✓ (DeepSeek-V3 风格) -T=128 N=256 S=1 k=6 : indices match ✓ routed close ✓ shared close ✓ -T=1 N=8 S=1 k=2 : indices match ✓ routed close ✓ shared close ✓ -T=16 N=32 S=1 k=4 : indices match ✓ routed close ✓ shared close ✓ - -ALL PASS ✓ -``` - -## 风险点 - -1. **K2 `tl.argmax` 在国产 NPU 上可能不支持** — 降级方案:把 K2 换成 `torch.topk` 1 行 -2. **K1 / K3 中间 buffer 大小** — `sel` 占 `T*N*4 bytes`(fp32),DeepSeek-V3 风格 (T=128, N=256) = 128KB, 8 芯片评估时 buffer 分配可能有限制 -3. **`tl.where` + 迭代 K 次的循环** — Triton 编译复杂度,某些 NPU 可能 register spilling - -## 预期结果 (我的估计) - -| 指标 | 估计 | 信心 | -|------|------|------| -| 跑通芯片数 | 6-7/8 | 中 (tl.argmax 是最大变数) | -| 平均加速比 | 3-8× | 中 (3 kernel 限制上限) | -| 排名 | 跟 sitraliqui(9.05x) 接近或略低 | 低 | - -## 之之的提交步骤 - -1. 打开平台 flagos.net → 第 3 批 → Task 38 → "提交代码" -2. 上传 `/Users/zhizhi/Desktop/sigmoid_gate_topk_renorm.zip` -3. 提交 → 等结果(通常 1-3 分钟) -4. 看到结果告诉我,如果有芯片 Failed 我会分析 - -## 之后怎么调 (如果第一次不理想) - -- **跑通但分低**(6-8×): 把 K1 和 K3 合并成 1 个 kernel, 减少 launch overhead -- **某芯片 Failed**: 看错误, 大概率是 tl.argmax, 改用 torch.topk -- **跨芯片不稳**: 加 7 芯片特化版 (跟 K30 一样 1+7 写法) diff --git a/zz-flagos-s2-2026/d246/results/sigmoid_gate_topk_renorm.zip b/zz-flagos-s2-2026/d246/results/sigmoid_gate_topk_renorm.zip deleted file mode 100644 index 465bcf8dcdc0bf94fdf895f3ed743b58c9674996..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 2618 zcmaLZc{J3G8VB&-SjL_x`~FhM7G>X)CHpeSGWKmOGt8i2WQme}k3zx-31b^ZwvZBH zcx^F=Xd=pFUXsSGd;h!V+n4f2$cci~(BqBV-6X_2}Aj1^HV;mtY0R5^(N${UUq1gZs0~G`S|1Nvw z5{6is*z;RHV$-85V2=&36!t5I+0YpQ$3-i%om>QIt<@H6c~QU>wnm6bB|T7}n2 zL-qM@A288hY!=(=KN1{-6Ddo&A$1OU{^v*C5sTTg9jnszYU#sNj`bn(NQLQPJpHop z(V?k&rT5*lp*QtlTVAwOd$xoHgs5BJGOj!2l+qo|T;N=hjj#BGtk&P%)$eEOIE)8X zBj6jLOue2QeFLEc8b>!BLaE}HsIcC`L!G{-&bA34=;5(M<}ExYB0dw(D3Pg0<8>aI zek*X|Hi8q0|A({AMYZJ1udtqSh2;VXz{nb$yMH!{sVvSb zRfylMzr`-JU@A1--?#CW^77-wZ6^E< zWVNLmq&)EXa6D-fV@Q>+UP{1QgbmJc?_Vj&s#)iDeNJrFAxRn8$gxN1cH(fkZJ5V9 zsmpK3NrsnyK)$$Izd&@b+yt?XZX#=+9D~LDJU0`S9Ji8>IH)HE3#S&^)zk3pkeJsbeTATa5OxR^L@MgeSI;W1i)jj?L;DWiQmb ztL=>Wi+v>f6M3~EIZ@es#m`*QSBK<`Cp$6GK?rVk8x36P{ncO|bGag!5MQ{CiZO&7 z9F+U1>cg;lz5DUF3Tx%af--D8*GUNborFhFEp_%TK)RG4A@bu)wPeA2#bhe*{~}V!E{ZsOTP)6`y8;iQH!O^+4vfr8-Q2 z!r8S=n3+S)PR%w>?W`5x#$u?@TB@4;2K|d6WC(c||MVKunNGd5_m17#qjQndcdx-p z-z(?49Xy)l=6p&Lwmx@jE-{KvmFhh6FfKBGW)B4stRKyL$l6Anj1c#Jr14{zVDq74 z0N14O-a5qY;_clN#_z8?Zsk9If;M#ZYR_!lJofU9(;z>ef|o3l9d#A;YrRfJ8}`c^ zi6l&QV=+_03H4<`vx!mylVCZAydA{&q~_Ck(YWrGd(nvx>&;)pmmgyUM&|*{9k&95lPULioY~3+ljNzqtOiz=N=JbM52VmqbD@QZw6`MIIP{Sfq_M$ zX@ebwMc&vHm);+m{48rySZ@4AfCu5PK!Q#wi9)n|P<+nCW861JFRL6)^PebrSguu` z75rlBN+2&x*Xq&R1sB*@Vd;33o*6X)(z@ILLQ%p{^5Wh{_8tKwwLzHw+w6jEU; zIGe8=l7C{7-v{7~3eys)*CCKxevC+dv1l;{u|KasDhKH40{Cosh zdiKuSps*!EBs= z)7lzG^((Uffv@Kf-+gtIlIYdY-Mg$8GU{ut=4-BmH95GkRB%u!G%c9nEF}9uN zTzTe0zqEY=Yk^Psc`=M)-upTk#-uf(WM47Jr552L&ja1(#@auL`B=rZm}{cvDj!0KjU#envkhBr?p`-_c|QFP_#qBL5lVoQG&Y*Z{}Ka&Ms!?z0?5vnuA+kAD*8Id(9Csf{*aP zm4S&<-k}m&YfFYq?;O3;PGi}EQSBYP+U_%2hWAiqZBD(eKY@*8N^`VcNqx9hk9|Tc z6Qh!UQ$Hh>Fk+=FgLkU*=yffQek#k;_IZ@zfNhx%zNhCbeW4;-&&1EFj`gQUaZn2c zO04mDMn1r3r0qaLU`^&la!4W_?HO1if^PoPTY~CfN=mD&1H;x46qaYzee|wQL9K^Y z18A`QCT#y>6NrTft;He7VtteykYb-Fqjhfhb{p&Mrx+e|rfHWK{itchbgnIJ8~J#xcqPwqX8waf6K)69r||Vn z0c5)a;A}?L3imdtti8?)Ac+atq^YU1XP|dOukKonSCV6;9+TSOnMzk!LiXADjFOAp zBjlV4$~)hbidgTN+%c_dNdNkAOIeZT3+5tc27PmDw>Dva5W{}nv&ceeYXkY#bgO89 zABE_6uS}En_$I_cM==rK%vSN60QaSE-YzM=K0bp6o3lMXb*6)9j5&ZNyEnb^lqL=M zIxRF)7AN$RE36^=p*)_R7AUaX{|AjdI-?#&S-!cQRNq?^YOFG8?knX=YXKTR-{__g@ OJ>I|f&~G6Ez`p=_+uDBs diff --git a/zz-flagos-s2-2026/d246/sigmoid_gate_topk_renorm.py b/zz-flagos-s2-2026/d246/sigmoid_gate_topk_renorm.py deleted file mode 100644 index bbcb764..0000000 --- a/zz-flagos-s2-2026/d246/sigmoid_gate_topk_renorm.py +++ /dev/null @@ -1,211 +0,0 @@ -"""Task 38 · sigmoid_gate_topk_renorm - -MoE 路由门算子 (DeepSeek-V2/V3 / Qwen3-MoE 风格) - 3-kernel Triton 实现: - - K1 (_sigmoid_bias_kernel) : sel = sigmoid(routed) + bias - K2 (_topk_kernel) : indices = topk(sel, k) (自写, 迭代 tl.argmax) - K3 (_gate_finalize_kernel): gather 原始 routed → cat shared - → sigmoid → renorm → scale → split 输出 - -国产 NPU 套路全套 (来自 D245 验证): - - input 指针 cast 到 int16/int32 (防 NaN 转换) - - stride / pid 用 int64 cast (防 overflow) - - with torch.get_device_module(x.device).device(x.device): 切设备 - - enable_fp_fusion=False, num_warps=4 - - 不用 tl.sigmoid, 用 1.0 / (1.0 + tl.exp(-x)) -""" -import torch -import triton -import triton.language as tl - - -# ============================================================ -# Kernel 1: sel = sigmoid(routed) + bias -# ============================================================ -@triton.jit -def _sigmoid_bias_kernel( - routed_ptr, # [T, N] fp16/bf16 - bias_ptr, # [N] fp32 - sel_ptr, # [T, N] fp32 (中间 buffer) - N, - stride_t, # int (T 维度 stride) - BLOCK_N: tl.constexpr, -): - # input 指针 cast 防 NaN - if routed_ptr.dtype.element_ty.primitive_bitwidth == 16: - routed_ptr = routed_ptr.to(tl.pointer_type(tl.int16)) - - pid_t = tl.program_id(0).to(tl.int64) - pid_n = tl.program_id(1).to(tl.int64) - - offs = pid_n * BLOCK_N + tl.arange(0, BLOCK_N).to(tl.int64) - mask = offs < N - - x = tl.load(routed_ptr + pid_t * stride_t + offs, mask=mask, other=0).to(tl.float32) - b = tl.load(bias_ptr + offs, mask=mask, other=0.0) - - # 不用 tl.sigmoid (国产 NPU 不支持), 用 1/(1+exp(-x)) - sig = 1.0 / (1.0 + tl.exp(-x)) - sel = sig + b - - tl.store(sel_ptr + pid_t * N + offs, sel, mask=mask) - - -# ============================================================ -# Kernel 2: top-k 选 indices (自写, 迭代 argmax) -# ============================================================ -@triton.jit -def _topk_kernel( - sel_ptr, # [T, N] fp32 (来自 K1) - idx_ptr, # [T, k] int32 - N, - K: tl.constexpr, - BLOCK_N: tl.constexpr, -): - pid = tl.program_id(0).to(tl.int64) - - offs = tl.arange(0, BLOCK_N) - mask = offs < N - - sel = tl.load(sel_ptr + pid * N + offs, mask=mask, other=-float('inf')) - - NEG_INF: tl.constexpr = float('-inf') - for i in tl.static_range(K): - # 找最大值的索引 (scalar) - idx = tl.argmax(sel, axis=0) - # 把 sel[idx] 设为 -inf (避免重复选) - is_max = (offs == idx) - sel = tl.where(is_max & mask, NEG_INF, sel) - # 写 indices - tl.store(idx_ptr + pid * K + i, idx.to(tl.int32)) - - -# ============================================================ -# Kernel 3: gather + cat + sigmoid + renorm + scale + split -# ============================================================ -@triton.jit -def _gate_finalize_kernel( - routed_ptr, # [T, N] fp16/bf16 - shared_ptr, # [T, S] fp16/bf16 - idx_ptr, # [T, k] int32 - route_scale, # float - global_scale_ptr, # [1] fp32 - routed_w_ptr, # [T, k] output dtype (input dtype) - shared_w_ptr, # [T, S] output dtype (input dtype) - N, S, - stride_routed_t, stride_shared_t, - stride_routed_w_t, stride_shared_w_t, - K: tl.constexpr, - BLOCK_S: tl.constexpr, # next_pow2(S) -): - # input 指针 cast 防 NaN, output 不 cast (output 是计算结果, 不需保 NaN bits) - if routed_ptr.dtype.element_ty.primitive_bitwidth == 16: - routed_ptr = routed_ptr.to(tl.pointer_type(tl.int16)) - shared_ptr = shared_ptr.to(tl.pointer_type(tl.int16)) - - pid = tl.program_id(0).to(tl.int64) - - # ===== 加载 routed 的 K 个值 (按 indices gather) ===== - idx_offs = tl.arange(0, K) - indices = tl.load(idx_ptr + pid * K + idx_offs).to(tl.int64) - routed_vals = tl.load(routed_ptr + pid * stride_routed_t + indices).to(tl.float32) - routed_sigmoid = 1.0 / (1.0 + tl.exp(-routed_vals)) - - # ===== 加载 shared 的 S 个值 ===== - s_offs = tl.arange(0, BLOCK_S) - s_mask = s_offs < S - shared_vals = tl.load( - shared_ptr + pid * stride_shared_t + s_offs, mask=s_mask, other=0 - ).to(tl.float32) - shared_sigmoid = 1.0 / (1.0 + tl.exp(-shared_vals)) - - # ===== 计算 sum (routed + shared 一起归一化) ===== - sum_routed = tl.sum(routed_sigmoid, axis=0) - sum_shared = tl.sum(shared_sigmoid, axis=0) - total_sum = sum_routed + sum_shared - - # ===== scale ===== - gs = tl.load(global_scale_ptr) - inv = (route_scale * gs) / total_sum - - # ===== 写 routed_w ===== - routed_w = (routed_sigmoid * inv).to(routed_w_ptr.dtype.element_ty) - tl.store(routed_w_ptr + pid * stride_routed_w_t + idx_offs, routed_w) - - # ===== 写 shared_w ===== - shared_w = (shared_sigmoid * inv).to(shared_w_ptr.dtype.element_ty) - tl.store(shared_w_ptr + pid * stride_shared_w_t + s_offs, shared_w, mask=s_mask) - - -# ============================================================ -# Python wrapper -# ============================================================ -def _next_pow2(x): - p = 1 - while p < x: - p *= 2 - return p - - -def sigmoid_gate_topk_renorm(logits, k, n_shared_experts, route_scale, global_scale, bias): - T, G = logits.shape - N = G - n_shared_experts - S = n_shared_experts - - if logits.ndim != 2: - raise ValueError('logits must have shape [T, N+S]') - if not isinstance(k, int) or k <= 0 or k > N: - raise ValueError(f'k must be 0 < k <= N={N}, got {k}') - if not isinstance(n_shared_experts, int) or n_shared_experts < 0 or n_shared_experts > S: - raise ValueError(f'n_shared_experts must be 0 <= S={S}, got {n_shared_experts}') - if bias.shape != (N,): - raise ValueError(f'bias must have shape [{N}], got {tuple(bias.shape)}') - if global_scale.numel() != 1: - raise ValueError(f'global_scale must be a scalar tensor, got shape {tuple(global_scale.shape)}') - if logits.device.type in ('cpu', 'meta', 'mps'): - raise RuntimeError('a real Triton accelerator backend is required') - - routed = logits[:, :N] - shared = logits[:, N:] - - # ===== 中间 buffer ===== - sel = torch.empty((T, N), dtype=torch.float32, device=logits.device) - indices = torch.empty((T, k), dtype=torch.int32, device=logits.device) - - BLOCK_N = max(16, _next_pow2(N)) - BLOCK_S = max(16, _next_pow2(S)) - - module = torch.get_device_module(logits.device) - with module.device(logits.device): - # ===== Kernel 1: sigmoid + bias ===== - grid1 = (T, triton.cdiv(N, BLOCK_N)) - _sigmoid_bias_kernel[grid1]( - routed, bias, sel, N, routed.stride(0), - BLOCK_N=BLOCK_N, - enable_fp_fusion=False, num_warps=4, - ) - - # ===== Kernel 2: topk ===== - grid2 = (T,) - _topk_kernel[grid2]( - sel, indices, N, K=k, BLOCK_N=BLOCK_N, - enable_fp_fusion=False, num_warps=4, - ) - - # ===== Kernel 3: gate finalize ===== - routed_w = torch.empty((T, k), dtype=logits.dtype, device=logits.device) - shared_w = torch.empty((T, S), dtype=logits.dtype, device=logits.device) - grid3 = (T,) - _gate_finalize_kernel[grid3]( - routed, shared, indices, route_scale, global_scale, - routed_w, shared_w, N, S, - routed.stride(0), shared.stride(0), - routed_w.stride(0), shared_w.stride(0), - K=k, BLOCK_S=BLOCK_S, - enable_fp_fusion=False, num_warps=4, - ) - - return routed_w, indices, shared_w - - -reference = sigmoid_gate_topk_renorm diff --git a/zz-flagos-s2-2026/d246/test_algorithm.py b/zz-flagos-s2-2026/d246/test_algorithm.py deleted file mode 100644 index b3df9db..0000000 --- a/zz-flagos-s2-2026/d246/test_algorithm.py +++ /dev/null @@ -1,158 +0,0 @@ -"""Task 38 算法级测试 - -用纯 torch 模拟 K1/K2/K3 的逻辑, 跟平台给的 reference 比较, 验证算法正确性。 -Triton 编译/运行 只能在有 GPU + Triton 的环境验证 (本机没 GPU)。 -""" -import torch -import sys -sys.path.insert(0, '.') - - -# ============================================================ -# 平台给的 reference (照抄, 完全一样) -# ============================================================ -def reference(logits, k, n_shared_experts, route_scale, global_scale, bias): - M, G = logits.shape - N = G - n_shared_experts - S = n_shared_experts - - logits_f = logits.float() - routed_logits = logits_f[:, :N] - sel = torch.sigmoid(routed_logits) + bias.float()[None, :] - - _, idx = torch.topk(sel, k, dim=-1) - routed_vals = torch.gather(routed_logits, 1, idx) - shared_vals = logits_f[:, N:N + S] - - active = torch.cat([routed_vals, shared_vals], dim=-1) - probs = torch.sigmoid(active) - weights = probs / probs.sum(dim=-1, keepdim=True) - weights = weights * route_scale * global_scale.float() - - routed_w = weights[:, :k].to(logits.dtype) - shared_w = weights[:, k:].to(logits.dtype) - indices = idx.to(torch.int32) - return routed_w, indices, shared_w - - -# ============================================================ -# 算法模拟: 分 3 步, 跟 Triton 写法一一对应 -# ============================================================ -def algo_sigmoid_bias(routed, bias): - """模拟 K1: sel = sigmoid(routed) + bias""" - return torch.sigmoid(routed.float()) + bias.float()[None, :] - - -def algo_topk(sel, k): - """模拟 K2: topk 选 indices""" - _, idx = torch.topk(sel, k, dim=-1) - return idx.to(torch.int32) - - -def algo_gate_finalize(routed, shared, indices, k, route_scale, global_scale, output_dtype): - """模拟 K3: gather + cat + sigmoid + renorm + scale + split""" - # gather 原始 routed - routed_vals = torch.gather(routed.float(), 1, indices.long()) - shared_vals = shared.float() - - # cat - active = torch.cat([routed_vals, shared_vals], dim=-1) - # sigmoid + renorm - probs = torch.sigmoid(active) - weights = probs / probs.sum(dim=-1, keepdim=True) - # scale - weights = weights * route_scale * global_scale.float() - - # split + cast - routed_w = weights[:, :k].to(output_dtype) - shared_w = weights[:, k:].to(output_dtype) - return routed_w, shared_w - - -def algo_full(logits, k, n_shared_experts, route_scale, global_scale, bias): - T, G = logits.shape - N = G - n_shared_experts - S = n_shared_experts - routed = logits[:, :N] - shared = logits[:, N:] - - sel = algo_sigmoid_bias(routed, bias) - indices = algo_topk(sel, k) - routed_w, shared_w = algo_gate_finalize( - routed, shared, indices, k, route_scale, global_scale, logits.dtype - ) - return routed_w, indices, shared_w - - -# ============================================================ -# 跑测试 -# ============================================================ -def test_case(T, N, S, k, dtype=torch.float16, seed=42): - """跑一个测试用例, 比较 reference 和 algo_full""" - torch.manual_seed(seed) - G = N + S - logits = torch.randn(T, G, dtype=dtype) * 2 - bias = torch.randn(N, dtype=torch.float32) * 0.1 - route_scale = 1.5 - global_scale = torch.tensor([2.0], dtype=torch.float32) - - # 平台 reference - ref_routed_w, ref_indices, ref_shared_w = reference( - logits, k, n_shared_experts=S, route_scale=route_scale, - global_scale=global_scale, bias=bias - ) - - # 我们的算法 - algo_routed_w, algo_indices, algo_shared_w = algo_full( - logits, k, n_shared_experts=S, route_scale=route_scale, - global_scale=global_scale, bias=bias - ) - - # 比较 - idx_match = (ref_indices == algo_indices).all().item() - routed_close = torch.allclose(ref_routed_w.float(), algo_routed_w.float(), atol=1e-3, rtol=1e-3) - shared_close = torch.allclose(ref_shared_w.float(), algo_shared_w.float(), atol=1e-3, rtol=1e-3) - - print(f" T={T} N={N} S={S} k={k} dtype={dtype}:") - print(f" indices match: {idx_match} routed close: {routed_close} shared close: {shared_close}") - - if not (idx_match and routed_close and shared_close): - # 打印前几个 mismatch - if not idx_match: - mismatch = (ref_indices != algo_indices).nonzero() - print(f" First 3 idx mismatches: {mismatch[:3].tolist()}") - if not routed_close: - diff = (ref_routed_w.float() - algo_routed_w.float()).abs() - print(f" Max routed diff: {diff.max().item()}") - if not shared_close: - diff = (ref_shared_w.float() - algo_shared_w.float()).abs() - print(f" Max shared diff: {diff.max().item()}") - return False - return True - - -if __name__ == "__main__": - print("Testing Task 38 algorithm...") - print() - - test_cases = [ - # (T, N, S, k) - (1, 64, 1, 8), # tiny - (1, 64, 1, 4), # k < 8 - (8, 64, 1, 8), # batch - (32, 128, 2, 8), # medium, group-style - (64, 256, 1, 8), # DeepSeek-V3-ish - (128, 256, 1, 6), # k=6 - (1, 8, 1, 2), # very tiny - (16, 32, 1, 4), # small N - ] - - all_pass = True - for T, N, S, k in test_cases: - ok = test_case(T, N, S, k) - all_pass = all_pass and ok - - print() - print("=" * 50) - print(f"{'ALL PASS ✓' if all_pass else 'SOME FAILED ✗'}") - print("=" * 50)