Skip to content

[BACKEND] Triton 3.6:重建 GPU CommonIR 并直接转换为 TTGIR - #92

Open
macqueen09 wants to merge 1 commit into
triton_v3.6.xfrom
commonir-triton36-rebuild-20260908
Open

macqueen09 wants to merge 1 commit into
triton_v3.6.xfrom
commonir-triton36-rebuild-20260908

Conversation

@macqueen09

@macqueen09 macqueen09 commented Sep 8, 2026 •

Copy link
Copy Markdown
Collaborator

修改目的

修复 Triton 3.6 GPU CommonIR 转换链路中的依赖方向和语义问题:独立的 CommonIRToTTGIR conversion 直接生成 Triton/TTGIR 原生操作,去除 conversion 对 TLE dialect 的 include、link 和 pass 依赖,也不再通过创建 TLE op 完成普通 buffer copy。

例如 tle.gpu.copy(global_pointer_tensor, shared_buffer, shape) 在 CommonIR 开启时先生成 tile.copy,再转换为原生 ttg.async_copy_global_to_local / commit / wait;shared→global 使用 ttg.local_load + tt.store。保留原有 layout 和显式转换流程,并对尚未支持的语义给出诊断。

本 PR 从 triton_v3.6.x 的精确 base 3f863acd700bda89514cf3dce6c57020d593d8f3 创建,包含一个提交 8e1d33333a391912066c7e0d9388b08b564a5afe,共 33 个代码/测试文件。关联 #88 的两条审查意见:

转换流程与代码组织

flowchart LR
  A[TLE GPU Python API] --> B[TileIR / CommonIR]
  B --> C[独立 CommonIRToTTGIR conversion]
  C --> D[原生 TTGIR copy / memdesc]
  D --> E[既有 GPU lowering / LLVM / PTX]
  F[既有 pipe / TMA / source TLE op] --> E
Loading
部分 主要修改 目的
构建与依赖 FLAGTREE_COMMON_IR 显式构建开关;条件加入 FLIR、conversion、dialect/pass 注册和 lit feature ON/OFF 在编译时决定,OFF 无需 FLIR 源码
Python 开关 新增 _common_ir.py,导入时缓存 C++ 暴露的开关值;删除旧 Python 运行时 TileIR 开关 Python 前端与已编译后端使用同一个开关
GPU DSL core.py、semantic.py、types.py 构造真实 TileIR alloc/copy/subview/to_tensor/store_tensor,并携带布局信息 保持显式 CommonIR 流程和 GPU 布局
FLIR builder triton_tle.cc 适配冻结 FLIR;通过 builtin cast 连接既有 memdesc 消费者 保留现有 pipe、local_ptr 等接口的类型衔接
独立 conversion 新增 third_party/nvidia/include/CommonIRToTTGIR 与 lib/CommonIRToTTGIR conversion 无 TLE include/link、无 TLE op 创建
NVIDIA pipeline inliner 后、canonicalizer/CSE 前执行 conversion,SM80+ 显式开启 async-copy 在冻结 FLIR 的部分 memory effect 建模不足影响 CSE 前完成 lowering
TTIR→TTGIR 衔接 在 CommonIR 宏下增加原生 async-copy 的动态合法性及类型转换规则 使新产生的 global pointer tensor 获得后续 TTGIR 所需 layout
pipe / TMA pipe 保留原同步协议和 memdesc 操作接口;descriptor/TMA copy 保留既有路径 不将已有复杂同步错误替换为简单 barrier
测试 独立 IR 正反例、真实 GPU copy roundtrip、DSL 错误诊断、更新 native matmul/flash/primitives dump 验证转换、诊断和端到端执行

FLIR 固定为 kateyijian/flir@a884fd3c608db2203f08ec5a281327f5dc0153b0,未修改其源码;未修改 TleOps.td 或社区 TMA op 定义。原 #88 中与功能无关的 test/CommonIR/gpgpu/matmul_ampere_gluon.py 格式修改未移入。

原生 copy 与同步

  • global→shared:支持的 GPU 上直接生成 TTGIR async-copy/commit/wait;同步路径生成 tt.load + ttg.local_store。CLI pass 的 enable-async-copy 默认关闭,NVIDIA 编译流水线按 capability >= 80 启用。
  • shared→global:ttg.local_load + tt.store。
  • 只合并同一 block 内、不同已知 allocation 的独立 copy 等待;只跨越不含 region 且无内存副作用的计算。遇到读写、控制流、相同 allocation 或未知 alias 即停止合并,保留 copy 对后续消费者已完成的语义。
  • conversion 不添加 tle.local_ptr_async_store 标记。源程序中原有的 TLE op 继续交给原有 TLE passes。
  • tle.gpu_layout 临时属性用于恢复布局,不构成 TLE dialect 链接依赖;冻结 FLIR 所需的 builtin.unrealized_conversion_cast 在转换后清除。

支持边界与错误诊断

  • 普通 copy 限于 full-buffer、shape/dtype 一致的 global pointer tensor 与 shared buffer;不支持 partial copy、transpose/NZ 或通用 engine 参数。
  • 同 rank subview 支持静态、边界内、单位 stride;降 rank 支持 leading stage 索引和零 trailing offset。
  • to_tensor/store_tensor 检查完整 buffer shape 和元素类型;普通 copy 的 offsets 不再被静默忽略。
  • 不支持的 tmem 在构造布局前报错;未实现的 tile engine/event 同步明确报错,不替换为 CTA barrier。
  • descriptor/TMA copy 仍走既有路径,不能将本 PR 描述为“所有 copy overload 均经 TileIR”。
  • 较旧 GPU 的同步 lowering 有 IR 测试;此次实际 GPU 执行验证在 H800 上完成,不代表所有架构均已实测。

验证环境与源码版本

项目 实际使用值
FlagTree base 3f863acd700bda89514cf3dce6c57020d593d8f3
本 PR 8e1d33333a391912066c7e0d9388b08b564a5afe
旧 #88 HEAD,仅核心补充对照 59d3f12ef29742263109bc846e29abffd2b334d8
Python aikg-t34a-src 的 Python 3.11.14
实际导入的 Triton 3.6.0;base/new/OFF 分别验证 CommonIR false/true/false
Torch / CUDA runtime 2.9.1+cu128 / 12.8
GPU NVIDIA H800;精度 GPU 2,主性能配对 GPU 5,旧 HEAD 补充配对 GPU 6
LLVM f6ded0be897e2878612dd903f7e8bb85448269e5
CMake 3.31.10
FlagAttention 90aa7aecb350a38b41f19c24be863c92dd48757f
FlagGems 38f083068f1e18b27665b24bdab39bfc7ea0dc83
FlagGems-vllm f142968ae0138e5b84fc05a4537bd61b72c8aeed

原 conda 环境直接导入曾为 Triton 3.4.0;测试通过隔离运行目录显式加载实际 3.6.0,并逐进程断言版本/开关。算子源码的 Python 文件与固定下载归档逐个核对一致,未修改算子实现或放宽精度容差。

本轮依赖补测环境

在优先环境 aikg-t34a-src 的 Python 3.11.14 上建立完全隔离的工作目录 venv,原 Conda 环境不变。正式组合为 Torch 2.11.0+cu129、CUDA runtime 12.9、vLLM 0.20.0+cu129、FLA/fla-core 0.5.2、NumPy 2.3.5、pytest 9.0.1;pip check 为 No broken requirements found。算子源码通过 PYTHONPATH 加载,本轮使用的 NumPy 2.3.5 与 FlagGems pyproject 原固定值 1.26.4 不同,受影响原测试的实测结果按下表记录。

实际加载的 Triton 仍为 3.6.0,候选 CommonIR C++ 开关为 true,Base 为 false。NVIDIA H800 / SM90、驱动 550.163.01;正式补测使用空闲 GPU 0。候选 wheel 另通过新环境下 native matmul、native flash GPU 精度检查。代码、FLIR、算子及断言不变,33 个编译器文件和 3,679 个算子 Python 文件哈希复核一致。

构建、专项与既有测试

检查 精确 base 本 PR,CommonIR ON 本 PR,CommonIR OFF
native matmul 精度 PASS PASS PASS
native flash 精度 已有编译失败 PASS 同 base 失败
local_ptr / remote / TMA / warp capture 4/4 PASS 4/4 PASS 4/4 PASS
TLE lit 56 PASS / 5 FAIL 同 5 项失败 同 5 项失败
新 CommonIR lit 不适用 5/5 文件 PASS 5 项按宏禁用
新 DSL/GPU semantics 不适用 11/11 PASS 不适用

5 个 CommonIR lit 文件覆盖:alloc/copy/subview/tensor/bridge、async-copy 与后续 layout 转换、函数参数/调用与 warp partition capture、copy 等待合并边界,以及 17 个 negative IR case。11 个 DSL/GPU 用例包含 3 个实际 GPU copy roundtrip 和 8 个错误诊断。

最终 matmul TTGIR 保留 2 个 async-copy / 1 个 wait;flash 保留 3 个 async-copy / 2 个 wait,均无 TileIR buffer 残留。最终 wheel 独立安装并清除 PYTHONPATH 后,再次验证实际 Triton 3.6.0 / CommonIR true,matmul、flash、primitives 全部通过;primitives 脚本用于验证 IR 流程。

格式及静态检查:clang-format 19.1.6、ruff 0.9.1、yapf 0.43.0、mypy 1.15.0 通过。pre-commit 初始化 GitHub hook 仓库时网络阻塞,随后使用同版本本地 hooks 与对应文件过滤完成等价检查;不将原 pre-commit 命令记为成功。以上为本地验证结果,不代表 GitHub CI 已通过。

算子精度与性能表(含 2026-09-08 依赖补测)

本轮补齐依赖后,Base 和候选各完成 40 项精度检查,均为 0 失败、0 跳过:两份 KV-cache 各 10 项、top-k prefill 14 项、FLA 对照 6 项。以下完整 18 行汇总中,17 行有有效 Base 性能配对且均达到无回归门槛,包含 2 行已注明的 fallback;native flash 因 Base 编译失败仍为 UNPAIRED。

P=通过实例数,S=原测试跳过实例数。未受依赖影响的条目沿用原测数据。FA-001 的 P34/S12 是原测 28 pass 与本轮 6 pass 合并;没有将完整 46 项声称为新环境的一次全量运行。FA-001、FG-004 性能沿用原测;两份 KV-cache 精度、性能均为本轮新增。同一行内 Base 和候选始终使用相同环境和 GPU,不跨批次配对。耗时为 CUDA Graph GPU 执行时间。

ID / 算子 Base 精度 候选精度 Base → 候选(μs) 变化 性能结论 路径 证据批次(精度;性能)
GATE-001 native_matmul_gpu PASS PASS 6.0880 → 6.0768 -0.18% PASS CommonIR 原测;原测
GATE-002 native_flash_gpu FAIL(base 已有) PASS — — UNPAIRED CommonIR 原测;原测
FA-001 gated_delta_rule P34 / S12 P34 / S12 35.2688 → 35.1392 -0.37% PASS 既有 TLE/PTX 原测28项+本轮6项,env-cu129;原测
FG-001 w8a8_block_fp8_matmul PASS(3 shape) PASS(3 shape) 2.0096 → 2.0056 -0.20% PASS 既有 TLE/PTX 原测;原测
FG-002 flash_mla_sparse P1 P1 9.1240 → 9.1504 +0.29% PASS CommonIR 原测;原测
FG-003 flash_mla_with_kvcache P10 P10 15.1856 → 15.1896 +0.03% PASS CommonIR 本轮,env-cu129;本轮,env-cu129
FG-004 top_k_per_row_prefill P14 P14 8.3824 → 8.3696 -0.15% PASS CommonIR 本轮,env-cu129;原测
FG-005 top_k_per_row_decode P7 P7 17.8336 → 18.0728 +1.34% PASS CommonIR 原测;原测
FG-006 bin_topk P2 / S36 P2 / S36 4.7944 → 4.8048 +0.22% PASS CommonIR 原测;原测
FG-007 moe_align_block_size P61 P61 11.7840 → 11.7072 -0.65% PASS FALLBACK 原测;原测
FG-008 fft PASS(1 shape) PASS(1 shape) 6.0848 → 6.0416 -0.71% PASS CommonIR 原测;原测
FGV-001 flash_mla P4 P4 35.7776 → 36.0328 +0.71% PASS CommonIR 原测;原测
FGV-002 flash_mla_sparse P1 P1 9.1536 → 9.0976 -0.61% PASS CommonIR 原测;原测
FGV-003 flash_mla_with_kvcache P10 P10 15.1840 → 15.1680 -0.11% PASS CommonIR 本轮,env-cu129;本轮,env-cu129
FGV-004 chunk_gated_delta_rule PASS(1 shape) PASS(1 shape) 28.3200 → 28.3008 -0.07% PASS CommonIR 原测;原测
FGV-005 chunk_kda P8 P8 12.1312 → 12.1424 +0.09% PASS CommonIR 原测;原测
FGV-006 moe_align_block_size P48 P48 11.7408 → 11.6696 -0.61% PASS FALLBACK 原测;原测
FGV-007 bin_topk P2 / S36 P2 / S36 4.7712 → 4.7840 +0.27% PASS CommonIR 原测;原测

交替计时与样本选择

  1. 每项开始前检查目标 GPU 利用率为 0%;A/B 使用同一张 GPU,各自在常驻进程完成 JIT/autotune 和预热,精度另行验证。
  2. 统一用预先 capture 的 CUDA Graph 测 warm-cache GPU 执行时间,排除 Python 提交间隙。每个 graph 含 20 次 workload 调用;每轮记录 30 次 replay 的 CUDA event 耗时,除以 20 后取中位数。
  3. 固定执行 3 组 ABBA,共 12 个样本。6 对相邻 A/B 或 B/A 中,仅选择 monotonic 测量中点距离最小的一对。表格每行只展示这一对;不按耗时相近程度或性能优劣选择。
  4. 回归门槛为同时变慢 >5% 且 >0.5 µs。保留全部时间戳、分位数、配对索引、GPU 状态及原始 IR;报告脚本独立重算并验证配对选择。
  5. 初轮普通 launch 计时在 PTX 完全一致的 workload 上也有明显波动,因此仅作为诊断资料保存。主表统一来自之后完整的 CUDA Graph 测试,不混合两种计时结果。

旧 HEAD 的核心补充对照如下,仅用于说明相对既有可运行 CommonIR 版本的性能,不替代主表 base:

用例 旧 HEAD µs 本 PR µs 变化 结果
GATE-001 6.094 6.127 +0.54% PASS
GATE-002 2.832 2.819 -0.45% PASS

已有失败、阻塞与覆盖限制

native flash 在 base 上的编译失败

错误为 cannot downgrade invalid async copy with non-trivial async token users,指向 tle.gpu.copy(k_tile, k_buf, [BLOCK_N, D])。发生在 TTGIR 的 TritonTleDowngradeInvalidAsyncCopy,kernel 尚未启动,故没有精度比较或有效 base 性能值。

补充 IR 诊断显示:TritonGPUCoalesceAsyncCopy 将 K copy 的源布局从 sizePerThread=[1,4] 调整为 [1,1],按后续检查对应 FP16 的 2 字节连续拷贝,不在该 pass 接受的 4/8/16 字节范围内。降级为同步 load/store 时,K 与 V 共用双 token commit,而 base 的降级逻辑只处理独立单 token 链,因而失败。

本 PR 的 CommonIR 原生 TTGIR 路径保持合法异步拷贝并通过该用例;未修复 base 的通用降级 pass,OFF 仍会复现原错误。旧 HEAD 补充对照通过不能替代缺失的 base 性能配对。

依赖补齐后的 FG-003 / FGV-003、FG-004 和 FA-001

两份 KV-cache 原来因缺少 vLLM 而触发文件级 skip。本轮补齐参考后,两侧均为每文件 10 pass / 0 skip:6 个数值对照实例、4 个参数/metadata 错误检查,测试源码与断言未修改。vLLM 0.15.1 虽可导入,却缺少测试要求的 out 参数;最终使用官方 vLLM 0.20.0+cu129 与匹配的 Torch 2.11.0+cu129。输出指针复用、V32/ MODEL1 / extra KV / dense 均按原测试通过。

FG-004 从 P11/S3 变为完整 P14;FA-001 的 6 个 FLA 对照全部通过,余下 12 个可选外部 benchmark 保留原禁用状态。FG-006、FGV-007 的各 36 个跳过来自源码显式标记的既有 #2352 问题,不能通过安装 vLLM 解决,未取消标记。

两份 KV-cache 的稀疏 TLE 分派在冻结源码中被注释为 TODO,因此新增性能采用原 dense decode 配置:batch=4、seq_q=1、h_q=128、d_qk=576、head_dim_v=512、page_size=64、32 pages、seqlens=[128,256,256,320];预分配输出并复用 scheduler。两项候选前端均有 12 个 tile.alloc、0 个 tile.copy,转换后无残留 tile.buf,覆盖实际 CommonIR buffer 转换路径。

FG-003 的选中测量中点间隔为 11.402585 ms,Base/候选样本索引 7/6(0-based);FGV-003 为 11.392320 ms,索引 4/5。每项均保留完整 12 个 ABBA 样本,表格只采用时间距离最小的一对。两项性能均 PASS。

FG-007 / FGV-006:实际执行 fallback

实际 PTX 为普通 Triton stage kernels,未产生 TileIR alloc/copy。FG-007 的显式 dispatch 诊断显示 base 和本 PR 都因算子库自动调参重复传入 BLOCK_EXPERT,使 TLE cluster/atomic fused 调用失败并回退。因此两项仅表示 fallback 路径精度和性能无回归,不表示 TLE fused 实现通过。

TLE lit 的 5 项既有失败

base、ON 和 OFF 均为相同的 56 PASS / 5 FAIL,无新增失败:

  • Analysis/test-alignment.mlir
  • GPU/test_tle_reduce_data_duplication_remote_if.mlir
  • GPU/test_tle_remote_atomic_splat_ptr_lowering.mlir
  • GPU/test_tle_remote_call_arg_store_lowering.mlir
  • GPU/test_tle_wgmma_pipeline_serialized_warning.mlir

Wheel 与本地复现证据

最终 wheel 对应本 PR 提交 8e1d33333,文件名:

flagtree-0.6.0+commonir.rebuild20260908.git8e1d3333-cp311-cp311-linux_x86_64.whl

SHA256:d703d0962354bc016b7434c4f3cde5d70f0b7da37cb5a3b58ff40964f12010b8。

服务器本地文件:

/supercloud/llm-code/mkl/project/bench/commonIR/results/pr88_rebuild_20260908/wheel/flagtree-0.6.0+commonir.rebuild20260908.git8e1d3333-cp311-cp311-linux_x86_64.whl

以下为服务器上的本地审计资料,不是 GitHub 下载链接,也不随本 PR 提交构建目录或第三方源码快照:

  • plan/27_triton36_pr88_rebuild_code_audit.md:原 GPU CommonIR:适配冻结 FLIR main 并保持显式 TileIR 转换 #88 全部 28 个文件逐项审查。
  • plan/28_triton36_pr88_rebuild_results.md:完整验证报告与复现命令。
  • results/pr88_rebuild_20260908/operator_results.csv、operator_results.json:精度、所选性能配对和 shape。
  • 同结果目录的 metadata/、accuracy/、performance_graph/、candidate_final/、wheel_final_validation/:版本/hash、原始测试、ABBA 样本、IR 和 wheel 验证。
  • diagnostics/native_flash_base/、diagnostics/FG-003/:失败前 IR 与详细 skip/import 诊断。

本地复现入口(在上述 commonIR 工作目录,测试前选择空闲 GPU):

bash plan/pr88_rebuild_build.sh on
/home/aeye/anaconda3/envs/aikg-t34a-src/bin/python plan/pr88_rebuild_validate.py candidate --gpu 2 --primitives --output-name candidate_final
/home/aeye/anaconda3/envs/aikg-t34a-src/bin/python plan/pr88_rebuild_accuracy.py base --gpu 2
/home/aeye/anaconda3/envs/aikg-t34a-src/bin/python plan/pr88_rebuild_accuracy.py candidate --gpu 2
/home/aeye/anaconda3/envs/aikg-t34a-src/bin/python plan/pr88_rebuild_perf.py --gpu 5 --timer graph
/home/aeye/anaconda3/envs/aikg-t34a-src/bin/python plan/pr88_rebuild_perf.py --gpu 6 --timer graph --comparison old_head --cases GATE-001 GATE-002

这些工作目录辅助脚本不包含在该代码提交中。仓库内可直接审查的专项用例位于 test/CommonIR/;CommonIR ON 编译后可运行 native_matmul_gpu.py、native_flash_gpu.py、tle_gpu_primitives.py、pytest test/CommonIR/test_gpu_semantics.py 及上述 5 个 lit 文件。

依赖补测的本地证据

详细说明:plan/29_pr92_dependency_retest_results.md。完整大表、JUnit、原始时刻与 IR:results/pr92_dependency_retest_20260908/。环境入口:results/pr92_dependency_retest_20260908/env-cu129/bin/python。版本锁、pip check、源码及下载 wheel 哈希保存在该目录的 metadata/。原测 results/pr88_rebuild_20260908/ 完整保留。以上是本地验证结果,不代表 GitHub CI 状态。

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant