diff --git a/README.md b/README.md index d152b38..d5b74ff 100644 --- a/README.md +++ b/README.md @@ -1,4 +1,4 @@ -# zl_bench — FlagGems 单算子性能测试工具 +# zl_bench — FlagGems 算子性能测试工具 针对编译器(FlagTree)改动做算子级 A/B 性能对比的 pytest 插件集 + 驱动脚本。在 FlagGems benchmark 体系之上解决三个问题: @@ -6,6 +6,8 @@ 2. **可解释**:自动按 shape 收集每次运行实际使用的 ttgir(带可读的变体命名),供 IR 级 diff; 3. **口径统一**:cudagraph 计时消除 launch 开销,小 kernel 的对比不被 CPU 侧噪声淹没;capture 前按 warmup 预算显式预热吸收 autotune/JIT,首轮即稳态、多轮一致。 +单算子精测流程之上另有批量筛查层(`run_batch.py`,见「批量多算子测试」),可对双仓库上千个算子做例行体检,产物结构与单算子完全同构。 + ## 快速开始 ```bash @@ -54,6 +56,7 @@ softmax: |-----|------|------| | `REPLAY_FROM` | 空(record 模式) | 指向某次历史 run 目录,replay 其 autotune 选择。做 A/B 时 B 侧必须设(见「A/B 对比测试标准流程」) | | `USE_FLAGTUNE` | `0` | `0` 走普通 autotune(跳过搜索、快速验证);`1` 用 FlagTune 扩展调优空间(首跑全量搜索、慢)。A/B 两侧须取同值 | +| `DTYPES` | 空(上游默认 dtype 扫描) | 空格分隔的 dtype 白名单(如 `"bfloat16 float16"`),逐个转为上游 `--dtypes`。算子不支持时上游报 `can't be supported by this op` | | `PARALLEL_WARMUP_GPUS` | 空(串行) | 设为 `N`(≥2)时把 autotune sweep 分片到 N 卡并行,测量仍单卡串行;不设或 `<2` 则完全串行(见「加速:多卡并行 sweep」) | **输出** @@ -161,11 +164,46 @@ WARNING 1 config key(s) got different winners across shards — parallel interfe **冲突数就是这次加速的可信度指标**:0 可放心用;偏多说明 config 选择已被污染,出正式结论前不设该变量重跑一遍。其余行为(`REPLAY_FROM` 已设或可见卡不足 2 张时跳过、分片失败回退串行、轮转分片而非按块切、用子进程而非 xdist 的原因)见插件 docstring。分片绑卡遵守继承来的 `CUDA_VISIBLE_DEVICES`——用它选过卡时,分片只会落在你选的那几张上。 +## 批量多算子测试(run_batch.py) + +单算子流程之上的批量筛查层:按清单逐算子调 `run_pytest.sh`,多卡分片、每卡内部串行,产出逐算子状态表和逐测量行总表。**定位是筛查口径**——跨卡并行测量存在功耗/散热耦合噪声(幅度可到百分之几),发现可疑算子后回单算子串行流程(前几节)确认,不要直接拿批量数字下精细结论。 + +```bash +# 全量筛查(双仓库全部可收集算子,排除带无条件 skip 标记的与 ops/dismiss.txt 里的) +python run_batch.py --all --gpus 0,1,2,3,4,5,6,7 --dtypes bfloat16 --op-timeout 1800 + +# 指定子集:默认 main 仓,vllm 仓加前缀;同名测试函数用 @文件名 消歧 +python run_batch.py --ops softmax vllm:fused_marlin_moe nextafter_@nextafter_ + +# 断点续跑(PASS/SKIP 的复用,FAIL/TIMEOUT 的重跑) +python run_batch.py --all --batch-dir runs/batch_xxx --resume + +# 锁 config 复测:逐算子 REPLAY_FROM //,语义同单算子 REPLAY_FROM +python run_batch.py --ops-file my_ops.txt --replay-root runs/batch_xxx +``` + +**算子资产(`ops/`)**: + +| 文件 | 内容 | 维护方式 | +|-----|------|---------| +| `inventory_main.csv` / `inventory_vllm.csv` | 双仓库全部 benchmark 测试函数(op、op_file、无条件 skip 标记、pytest collect 核实结果、op_name 提取) | 上游更新后重跑 `python ops/gen_inventory.py --verify-collect` | +| `shapes_single.yaml` | 每算子单 shape 的筛查集(537 个精选 op_name 键 + 上游 core_shapes 的类名键底座——上游 shape 回退链是 op_name → MRO 类名 → 基类默认,缺类名键会让 BLAS 族跌到 1 维默认值崩溃) | 手工增补;上游变动后原地刷新:`python ops/gen_inventory.py --migrate-shapes ops/shapes_single.yaml --shapes-out ops/shapes_single.yaml`(未匹配键落 `.unmatched.yaml` 供复核) | +| `shapes_multi.yaml` | 每算子多 shape 的深查集(结构同上) | 同上 | +| `dismiss.txt` | 批量排除清单(`[repo:]op` 每行一个),只放本仓复核确认的失败项并注明原因/日期 | 批量跑出 FAIL 并确认原因后手工添加 | + +无 shape 条目的算子用上游默认 shape,照常可测。`--shape-file ops/shapes_multi.yaml` 切换深查集。 + +**执行语义**:每算子独立子进程(CUDA crash 只废单个算子);`--op-timeout` 对进程组 SIGTERM→SIGKILL,记 `TIMEOUT` 不重试、批次继续;pytest rc=1 与信号杀最多重试 2 次,连续同 rc 视为确定性失败提前止损;`--dtypes` 有两级降级——遇上游 `can't be supported by this op` 即去掉限制重跑(备注 `dtype_fallback`),其余失败且无成功行时也去掉限制最后救一次(备注 `dtype_rescue`,覆盖 torch baseline 对受限 dtype 编译失败的场景)。runtime skipif 与上游 `xfail` 都判 SKIP(上游声明的无信号状态);`.complete` 标记只给 PASS/SKIP,`--resume` 据此复用。 + +**产物**:`///` 与单算子 run 目录结构完全一致(run.log、autotune_records/、ttgir/……),所以任何一个算子都可以事后单独 `REPLAY_FROM` 复测。批量层额外产出 `ops_status.csv`(每算子一行:状态/耗时/行数/备注,边跑边原子更新)与 `summary.csv`(每测量行一行:dtype、latency、speedup、Size Detail,来源是 run.log 的结果表,`batch_summary.py` 也可单独对旧批次重跑)。备注列聚合 `no_cudagraph_rows`/`replay_fallback` 计数,解读口径时先看这列。 + ## 计时口径:cudagraph 的预热、回退与精度 **capture 前的显式 warmup**:`do_bench_cudagraph` 自带的内部预热只有 5 次迭代,对首跑要触发 autotune 编译(尤其含 FlagTune 扩展空间)、libtuner 选择、lazy JIT 的算子远远不够——这些一次性开销若漏进被捕获的 graph 或第一个计时迭代,测出的 latency 会 run-to-run 抖动(M=1 多 kernel 路径最明显,实测首轮可低到稳态的 ~1/3)。`_cudagraph_plugin` 因此在 capture 前显式预热:先跑一次并丢弃(吸收 autotune/JIT 编译),再按调用方传入的 warmup 时间预算(`Config.warm_up`)循环稳态预热,然后才 capture。预热次数按稳态单次耗时换算,并夹在 5–200 次之间——亚毫秒 kernel 的实际预热时长因此低于名义预算(1000ms),实测足够;若换新算子仍见首轮抖动,优先调大 `_warmup_before_capture` 里的次数上限。这样首轮即稳态、多轮一致(实测同一 M=1 shape 两轮 speedup 差 <0.1%)。 -部分算子本身不支持 CUDA graph capture(测量函数内含 host 同步、动态显存分配、不合法的流操作等),这类算子会自动回退到普通 do_bench 计时,`run.log` 中打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` 标记(含失败阶段与具体原因;`phase=warmup` 表示预热阶段就失败了,并非 graph capture 被拒)。 +部分算子本身不支持 CUDA graph capture(测量函数内含 host 同步、动态显存分配、autograd backward、不合法的流操作等),这类算子会自动回退到普通 do_bench 计时,`run.log` 中打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` 标记(含失败阶段与具体原因;`phase=warmup` 表示预热阶段就失败了,并非 graph capture 被拒)。 + +回退路径的两个防御措施(都吃过亏):(1)致命错误判定按错误原文匹配而非宽泛子串——新版 torch 给每个 CUDA 错误都追加 "enable device-side assertions" 提示语,宽松匹配会把所有 capture 失败误判为致命错误、令回退路径整体失效;(2)capture 在一次性 RNG state 替身下执行——`torch.cuda.graph` 会无条件注册默认 CUDA RNG 生成器,capture 中途失败可能让生成器卡在 capturing 态,此后进程内任何 `torch.randn` 都抛 "Offset increment outside graph capture"(毒化显形时机不定,事后修复不可靠),替身隔离让真身永不参与 capture,成败都原样归位。 回退本身不影响 A/B 公平性(两侧同一算子回退行为一致),但**回退口径的测量误差更大**:do_bench 每次迭代都走完整的 Python → launch 路径,kernel 越小,launch 开销和 CPU 侧抖动在数字里占比越高——亚毫秒级 kernel 上两种口径可差 2 倍以上,且行间波动更明显。解读这类算子的结果时: diff --git a/_cudagraph_plugin.py b/_cudagraph_plugin.py index 3b54a08..9c5564b 100644 --- a/_cudagraph_plugin.py +++ b/_cudagraph_plugin.py @@ -8,7 +8,9 @@ fall back to plain do_bench and print a BENCHMARK_DIRECT_NO_CUDAGRAPH marker into the log. Before capturing, fn is explicitly warmed up for the caller's warmup time budget (do_bench_cudagraph's own 5-iter warmup is too short to settle autotune/JIT for the MoE kernels, which makes the first captured -measurement unstable run-to-run). +measurement unstable run-to-run). The capture itself runs under a scratch CUDA +RNG state so that an aborted capture cannot poison the process-wide generator +(rationale: _capture_with_scratch_rng). """ from __future__ import annotations @@ -25,13 +27,19 @@ _ORIGINAL_DO_BENCH = _tt.do_bench def _is_fatal_cuda_error(exc: Exception) -> bool: """True for CUDA context errors (illegal address, OOM, device-side assert): the context is poisoned, so re-raise rather than fall back and record a bogus - latency as success.""" + latency as success. + + Markers must match the error text itself, not torch's generic footer: newer + torch appends "Compile with `TORCH_USE_CUDA_DSA` to enable device-side + assertions." to EVERY AcceleratorError, so a loose "device-side assert" + marker turned all capture-incompatible ops into fatal re-raises and killed + the documented do_bench fallback path entirely.""" msg = f"{type(exc).__name__}: {exc}".lower() markers = ( "illegal memory access", "cudaerrorillegaladdress", "out of memory", - "device-side assert", + "device-side assert triggered", "an illegal instruction", "misaligned address", "uncorrectable ecc", @@ -83,6 +91,33 @@ def _warmup_before_capture(fn, warmup_ms, grad_to_none): torch.cuda.synchronize() +def _capture_with_scratch_rng(call): + """Run a graph capture under a throwaway CUDA RNG state. + + torch.cuda.graph's capture_begin unconditionally registers the default CUDA + RNG generator and sets its capturing flag; when a capture aborts midway the + epilogue may never run, the flag sticks, and every later CUDA RNG call in + the process (the next shape's torch.randn included) raises "Offset + increment outside graph capture". Repairing after the fact is unreliable — + when the poisoning surfaces depends on which stage the capture died in — + so prevent it instead: swap in a scratch state via graphsafe_set_state + before capturing and restore the original state object afterwards, success + or failure. The real generator never takes part in a capture, so poisoning + can only land on the discarded scratch state, and _seed_plugin determinism + is left untouched.""" + import torch + dev = torch.cuda.current_device() + gen = torch.cuda.default_generators[dev] + saved = gen.graphsafe_get_state() + scratch = torch.Generator(device="cuda") + scratch.manual_seed(gen.initial_seed()) + gen.graphsafe_set_state(scratch) + try: + return call() + finally: + gen.graphsafe_set_state(saved) + + def _patched_do_bench(fn, warmup=25, rep=100, grad_to_none=None, quantiles=None, return_mode="mean"): """Drop-in replacement for triton.testing.do_bench using cudagraph capture.""" @@ -100,13 +135,13 @@ def _patched_do_bench(fn, warmup=25, rep=100, grad_to_none=None, # Settle autotune/JIT before capture (rationale: _warmup_before_capture). _warmup_before_capture(fn, warmup, grad_to_none) phase = "capture" - return _tt.do_bench_cudagraph( + return _capture_with_scratch_rng(lambda: _tt.do_bench_cudagraph( fn, rep=rep, grad_to_none=grad_to_none, quantiles=quantiles, return_mode=return_mode, - ) + )) except Exception as exc: # Re-raise genuine device failures; anything else falls back to plain # do_bench, with the failing phase (sync/warmup/capture) in the marker. diff --git a/_device_guard_plugin.py b/_device_guard_plugin.py index f735a1a..bea39a9 100644 --- a/_device_guard_plugin.py +++ b/_device_guard_plugin.py @@ -4,7 +4,7 @@ FlagGems' `runtime.backend.device_finder` falls back to `subprocess.run(nvidia-smi)` **without a timeout**. Under this conda env's forked/inconsistent subprocess (`_posixsubprocess` symbol mismatch), that child can hang indefinitely, leaving `import flag_gems` stuck in `wait4` -(seen as run_pytest "卡住 with no result", GPU 0%). +(seen as run_pytest hanging with no output, GPU at 0%). This plugin runs at import time — before any test module imports flag_gems — detects the vendor via torch (no subprocess), and sets GEMS_VENDOR so diff --git a/batch_summary.py b/batch_summary.py new file mode 100644 index 0000000..3092ecf --- /dev/null +++ b/batch_summary.py @@ -0,0 +1,186 @@ +#!/usr/bin/env python3 +"""聚合一次批量跑(run_batch.py 产物目录)的结果为 summary.csv。 + +数据源是各算子 run 目录里的 run.log(run_pytest.sh 已去 ANSI 色): +- 逐表解析 `Operator: (dtype=..., mode=..., level=...)` 标题 + (兼容 _pretty_report_plugin 与上游原生两种格式); +- 逐行解析 SUCCESS/FAILED 表行:状态 + 前三个数值列固定为 + torch_lat / gems_lat / speedup(tflops/gbps 等追加列收进 extra_metrics), + 余下为 size_detail; +- 汇入 run_batch.py 写的 ops_status.csv(每算子总状态 / 备注)。 + +也可单独使用:python batch_summary.py runs/batch_xxx [-o summary.csv] +""" + +from __future__ import annotations + +import argparse +import csv +import os +import re +import sys +from pathlib import Path +from typing import Dict, List, Optional, Tuple + +# 兼容两种标题: +# pretty: Operator: softmax (dtype=torch.float16, mode=kernel, level=core) +# upstream: Operator: softmax Performance Test (dtype=torch.float16, mode=kernel,level=core) +OPERATOR_HEADER_RE = re.compile( + r"^Operator:\s+(?P\S+)\s+(?:Performance Test\s*)?" + r"\(dtype=(?P[^,]+),\s*mode=(?P[^,]+),\s*level=(?P[^)]+)\)" +) +ROW_RE = re.compile(r"^(?PSUCCESS|FAILED)\s+(?P\S.*)$") +NUM_TOKEN_RE = re.compile( + r"^(?:N/A|nan|-?inf" + r"|-?(?:\d+(?:\.\d*)?|\.\d+)(?:[eE][+-]?\d+)?)$" +) +PYTEST_SKIPPED_RE = re.compile(r"(?P\d+)\s+(?:skipped|xfailed)\b") +DIRECT_NO_CUDAGRAPH_RE = re.compile(r"BENCHMARK_DIRECT_NO_CUDAGRAPH\b") +REPLAY_FALLBACK_RE = re.compile(r"AUTOTUNE_REPLAY_FALLBACK\b") + +SUMMARY_FIELDS = [ + "repo", "op", "fg_op_name", "dtype", "case_id", "row_status", + "torch_lat", "gems_lat", "speedup", "extra_metrics", "size_detail", + "op_status", "notes", +] + + +def parse_run_log(text: str) -> List[Dict[str, str]]: + """把 run.log 里的全部结果表解析为行列表(跨多个 dtype 表)。""" + rows: List[Dict[str, str]] = [] + cur_op, cur_dtype = "", "" + for line in text.splitlines(): + stripped = line.strip() + header = OPERATOR_HEADER_RE.match(stripped) + if header: + cur_op = header.group("name") + cur_dtype = header.group("dtype").strip() + continue + m = ROW_RE.match(stripped) + if not m: + continue + tokens = m.group("rest").split() + metrics: List[str] = [] + while tokens and len(metrics) < 7 and NUM_TOKEN_RE.match(tokens[0]): + metrics.append(tokens.pop(0)) + # 前三个数值列固定:torch_lat, gems_lat, speedup(上游列序承诺不变) + while len(metrics) < 3: + metrics.append("") + rows.append({ + "fg_op_name": cur_op, + "dtype": cur_dtype, + "row_status": m.group("status"), + "torch_lat": metrics[0], + "gems_lat": metrics[1], + "speedup": metrics[2], + "extra_metrics": " ".join(metrics[3:]), + "size_detail": " ".join(tokens), + }) + return rows + + +def pytest_skipped_count(text: str) -> int: + """runtime skipif 与上游 xfail 标记都算"无测量信号",供 SKIP 判定。""" + count = 0 + for m in PYTEST_SKIPPED_RE.finditer(text[-4000:]): + count += int(m.group("n")) + return count + + +def marker_notes(text: str) -> str: + """cudagraph 回退 / replay 兜底标记计数,供备注列(口径解读用)。""" + notes = [] + n = len(DIRECT_NO_CUDAGRAPH_RE.findall(text)) + if n: + notes.append(f"no_cudagraph_rows={n}") + n = len(REPLAY_FALLBACK_RE.findall(text)) + if n: + notes.append(f"replay_fallback={n}") + return ";".join(notes) + + +def classify(rc: int, rows: List[Dict[str, str]], log_text: str) -> str: + """算子级状态:FAIL / SKIP / PASS(TIMEOUT 由 run_batch 在外层判)。""" + if rc != 0: + return "FAIL" + if not rows: + return "SKIP" if pytest_skipped_count(log_text) else "FAIL" + if any(r["row_status"] != "SUCCESS" for r in rows): + return "FAIL" + return "PASS" + + +def _load_ops_status(batch_dir: Path) -> Dict[str, Dict[str, str]]: + path = batch_dir / "ops_status.csv" + if not path.is_file(): + return {} + with path.open(newline="") as f: + return {f"{r['repo']}/{r['run_name']}": r for r in csv.DictReader(f)} + + +def summarize(batch_dir: Path, out_path: Optional[Path] = None) -> Path: + batch_dir = batch_dir.resolve() + out_path = out_path or (batch_dir / "summary.csv") + status_by_dir = _load_ops_status(batch_dir) + + all_rows: List[Dict[str, str]] = [] + for repo_dir in sorted(p for p in batch_dir.iterdir() if p.is_dir()): + repo = repo_dir.name + for op_dir in sorted(p for p in repo_dir.iterdir() if p.is_dir()): + log = op_dir / "run.log" + if not log.is_file(): + continue + text = log.read_text(errors="replace") + status_row = status_by_dir.get(f"{repo}/{op_dir.name}", {}) + op = status_row.get("op") or op_dir.name + table_rows = parse_run_log(text) + if not table_rows: + all_rows.append({ + "repo": repo, "op": op, "fg_op_name": "", "dtype": "", + "case_id": "", "row_status": "", "torch_lat": "", + "gems_lat": "", "speedup": "", "extra_metrics": "", + "size_detail": "", + "op_status": status_row.get("status", ""), + "notes": status_row.get("notes", ""), + }) + continue + case_counter: Dict[Tuple[str, str], int] = {} + for row in table_rows: + key = (row["fg_op_name"], row["dtype"]) + idx = case_counter.get(key, 0) + case_counter[key] = idx + 1 + all_rows.append({ + "repo": repo, "op": op, "case_id": f"case_{idx:03d}", + "op_status": status_row.get("status", ""), + "notes": status_row.get("notes", ""), + **row, + }) + + tmp = out_path.with_name(out_path.name + ".tmp") + with tmp.open("w", newline="") as f: + writer = csv.DictWriter(f, fieldnames=SUMMARY_FIELDS) + writer.writeheader() + writer.writerows(all_rows) + os.replace(tmp, out_path) + return out_path + + +def main() -> None: + p = argparse.ArgumentParser(description=__doc__.splitlines()[0]) + p.add_argument("batch_dir", type=Path, help="run_batch.py 的产物目录") + p.add_argument("-o", "--output", type=Path, default=None, + help="输出 CSV(默认 /summary.csv)") + args = p.parse_args() + if not args.batch_dir.is_dir(): + raise SystemExit(f"目录不存在: {args.batch_dir}") + out = summarize(args.batch_dir, args.output) + with out.open(newline="") as f: + rows = list(csv.DictReader(f)) + n_ok = sum(1 for r in rows if r["row_status"] == "SUCCESS") + n_bad = sum(1 for r in rows if r["row_status"] == "FAILED") + print(f"[summary] {out}: {len(rows)} 行(SUCCESS={n_ok} FAILED={n_bad})", + file=sys.stderr) + + +if __name__ == "__main__": + main() diff --git a/ops/dismiss.txt b/ops/dismiss.txt new file mode 100644 index 0000000..2142d4b --- /dev/null +++ b/ops/dismiss.txt @@ -0,0 +1,120 @@ +# 批量测试排除清单。每行一个算子:`op`(默认 main 仓)或 `repo:op`。 +# run_batch.py --all 会跳过这里列出的算子;带无条件 pytest.mark.skip 的算子由 +# inventory 的 marked_skip 列自动剔除,不必写在这里。 +# 只放"经本仓批量跑复核确认失败"的条目,注明原因与复核日期。 +# +# 2026-08-12 复核记录: +# - 旧回归清单 24 个候选全部重验:14 个新上游已修复(移出),3 个系本仓 +# _cudagraph_plugin 回退失效 bug(已修复,移出),7 个确认为上游/环境问题; +# - 全量筛查(1103 算子)新增确认 67 个:43 个 torch jiterator 环境问题 + +# 24 个上游 bug/环境限制。shape 语义类失败(28 个)已通过 shape yaml 合并 +# 上游 core_shapes 类名键修复,不在此列; +# - test_FLA/ 补齐后新增确认 2 个(硬件限制 + 专用 harness),共 76 条。 + +# ==== 上游 kernel 编译错误 [2026-08-12] +# CompilationError(关 cudagraph 后仍复现) +hstack +scaled_mm_benchmark +scaled_mm_out_benchmark +# tt.load result/ptr type 校验失败 +flash_mla_with_kvcache +# PassManager::run failed(FlagTree triton 编译失败) +flash_mla_sparse_fwd +vllm:flash_mla_sparse_fwd +# capture 失败回退后仍报 kernel 错误 +vllm:flash_mla +vllm:flash_mla_with_kvcache +linalg_lu_factor_ex + +# ==== 上游 benchmark 代码 bug [2026-08-12] +# baseline torch.trace 收到非矩阵输入(expected a matrix) +trace +# set_shapes() 签名与基类调用不符(TypeError) +get_paged_mqa_logits_metadata +# 输入解包 ValueError(expected 3, got 1) +grouped_mm +# 上游 benchmark assert 失败 +get_scheduler_metadata +# is_backward 对未参与计算图的张量求梯度(allow_unused) +embedding_backward +margin_ranking_loss_backward +# baseline 函数签名不符(unexpected keyword 'q' / 缺 constexpr 实参) +fp8_fp4_paged_mqa_logits +vllm:fp8_fp4_mqa_logits +# baseline 返回 None(NoneType has no attribute 'device') +special_gammainc +# normalized_shape=[](上游默认 shape 语义错误) +native_layer_norm +# 上游断言 K>=16 与其默认 attention shape 冲突 +perf_scaled_dot_product_flash_attention_backward +perf_scaled_dot_product_cudnn_attention_backward +# Triton Error [CUDA]: invalid argument(上游默认 shape 下 kernel 启动参数非法) +index_select_backward +pairwise_distance +# test_blas_perf_parallel 专用 harness(shape kind 语义与通用注入冲突;常规 mul 已覆盖) +perf_mul +# 需要 pytest-benchmark fixture(非 FlagGems Benchmark 体系) +vllm:triton_unified_attention_perf +# 自定义输出格式(非 FlagGems Benchmark 表)且预编译超 40min +vllm:perf_chunk_gla + +# ==== 硬件限制 [2026-08-12] +# kernel 共享内存需求 245-335KB 超 H20 上限 227KB,autotune 全候选 OutOfResources +vllm:chunk_gdn2 + +# ==== 环境限制(换环境后应重验)[2026-08-12] +# vllm 0.20.2 无 vllm.utils.deep_gemm.fp8_mqa_logits +fp8_mqa_logits +# FA2 不支持 num_splits > 1(flash-attn 版本) +flash_attn_varlen_opt_func +# magma 显存分配失败(cannot allocate memory on GPU, info=-113) +cholesky_solve +cholesky_solve_out +# cusolver INTERNAL_ERROR(Xgeev) +linalg_eigvals +# ---- torch jiterator/NVRTC 在本环境编译失败(fp32 也复现), +# ---- sinc/bessel/special 族的 torch baseline 整体不可用: +erfc +erfc_ +erfinv +erfinv_ +lcm +lcm_ +lgamma +lgamma_ +mvlgamma +mvlgamma_ +polygamma +polygamma_inplace +polygamma_out +sinc +sinc_ +special_airy_ai +special_airy_ai_out +special_bessel_j0 +special_bessel_j1 +special_chebyshev_polynomial_u +special_chebyshev_polynomial_v +special_chebyshev_polynomial_w +special_chebyshev_polynomial_w_out +special_erfc +special_erfcx +special_erfinv +special_erfinv_out +special_gammaln +special_gammaln_out +special_hermite_polynomial_h +special_i1e +special_legendre_polynomial_p +special_log1p_non_tensor +special_modified_bessel_k0 +special_modified_bessel_k0_out +special_modified_bessel_k1 +special_modified_bessel_k1_out +special_scaled_modified_bessel_k1 +special_scaled_modified_bessel_k1_out +special_shifted_chebyshev_polynomial_u +special_shifted_chebyshev_polynomial_u_ +special_shifted_chebyshev_polynomial_v +special_shifted_chebyshev_polynomial_w +special_sinc diff --git a/ops/gen_inventory.py b/ops/gen_inventory.py new file mode 100644 index 0000000..b2dc9e7 --- /dev/null +++ b/ops/gen_inventory.py @@ -0,0 +1,294 @@ +#!/usr/bin/env python3 +"""重建批量测试的算子清单(inventory),并可校验迁移 shape yaml。 + +对每个 FlagGems 仓库静态扫描 benchmark/ 下的 test_*.py(含 test_ 前缀的 +子目录,如 test_FLA/——op_file 记相对路径去前后缀,run_pytest.sh 的 +`benchmark/test_${OP_FILE}.py` 拼接对其天然成立;models_benchmark/ 等 +非 test_ 前缀路径不在性能批量范围): + +- 每个模块级 test_* 函数出一行 inventory:op(函数名去 test_ 前缀,即 + run_pytest.sh 的 OP)、op_file(文件名去 test_ 前缀/.py 后缀,即 OP_FILE); +- 标注 marked_skip(函数带无条件 pytest.mark.skip,跑了必 SKIP,批量驱动默认剔除); +- 提取 op_name 字符串常量(op_name= 关键字实参,以及 *Benchmark 类 + __init__ 里 super().__init__ 的首个字符串位置实参),用于 shape yaml 键校验; +- --verify-collect 时额外跑一次 pytest --collect-only 核实函数确实可被收集 + (AST 见到 ≠ pytest 收得到,import 失败/条件定义都会导致差异)。 + +产物: + ops/inventory_