TileLang 每日动态报告(2026-11-02)
调研窗口:2026-10-30 07:00 ~ 2026-11-02 07:00(北京时间,72 小时,覆盖周末;上一期截点为 10-30 07:00)。
信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / TileOPs-nightly / 文档站)+ 合入与在审真值检索(合入 35 笔、在审更新 82 笔)+ 新增 issue 全量核对(11 条)+ 主分支逐笔提交核查(tilelang 7、TileOPs 15、TileFoundry 3、tilelang-ascend 5、tilelang-hygon 2、tilelang-mlir-ascend 1、文档站 4)+ 发布页与 tag 核查(6 仓)+ 第三方生态仓核查(FlashQLA、TileKernels、RISC-V 研究链三仓、PTO-ISA/tilelangGym)+ Google News 中英文查询(走代理)+ Hacker News + arXiv。
本期索引
- 今日重点:主线为 TIRX 迁移回退 ThreadSync 的精确线程计数(#3733),并新开跟踪单承认正确性改进一并撤回;TileOPs 单日 13 笔性能合入;昇腾宿主固定开销从约 450 微秒砍到约 106 微秒;72 小时合入 35 笔、在审更新 82 笔、新增 issue 11 条
- 一、核心项目进展
- 1.1 主线:ThreadSync 的精确线程计数回退,为 TIRX 迁移让路(#3733)(11-01)
- 1.2 主线:CUDA 跨 warp 全归约合并为一次共享内存交换,256 线程从 6 道栅栏降到 2 道(#3715)(10-30)
- 1.3 主线:PTO 后端的 T.Parallel 向量化落地第一里程碑,逐元素算子走 VMI 指令(#3512)(10-30)
- 1.4 主线:昇腾 DeepGEMM 的 scale 缓冲改用写计数多缓冲,数据加载可提到条件加载之前(#3629)(10-30)
- 1.5 主线:两笔修复合入,CPU 标量 GEMM 只清目标区域、负索引合法化跳过不可达分支(#3709 / #3712)(10-30)
- 1.6 主线:语言层清掉 proxy 里两个无人使用的 Ref 占位(#3722)(10-31)
- 1.7 新开缺陷与在审:TMEM 静默丢结果、ROCm 符号找不到、TF32 就近偶数舍入(#3721 / #3711 / #3714)(10-30 ~ 10-31)
- 二、多后端适配
- 2.1 昇腾:Cython wrapper 惰性构造 TVM Analyzer,910B3 单次调用从约 450 微秒降到约 106 微秒(#1948)(10-30)
- 2.2 昇腾:RFC 提出把同步实现换成主线核心加独立 Z3 证明服务(#1971)(10-30)
- 2.3 昇腾在审:把 TileLang-Ascend 整体迁到主线多后端架构(#1975)(10-31)
- 2.4 昇腾在审:GLM-5.3 kpool 的 fp8_mqa_logits 示例,双 kernel 加 AIV/Cube 分工(#1966)(10-31)
- 2.5 昇腾:四笔文档与 API 事实核查合入,clamp 在 PTO 上真正生效(#1899 / #1893 / #1620 / #1600)(10-30)
- 2.6 昇腾:日测 2970 用例全部通过,11-01 有一次 workflow 失败(#1977 / #1976)(10-31 ~ 11-02)
- 2.7 海光:MLS 边界与持续 WASP 同步修复合入,gfx948 兼容在审(#31 / #32)(10-30)
- 2.8 mlir-ascend:CI 清掉内部缓存服务引用,Codex 工作流在审(#206 / #208)(10-31)
- 2.9 摩尔线程与沐曦窗口内无推送(11-02)
- 三、生态与采用方
- 3.1 TileOPs:72 小时 15 笔合入,13 笔压在 11-01,几乎每笔都带实测表与清单回归门槛(11-01)
- 3.2 TileOPs:GQA 稠密前向用一条实测规则选 decode 或 prefill 程序,选择遗憾从 1.818 降到 1.017(#2808)(11-01)
- 3.3 TileOPs:分页 MHA decode 扩到共享内存装得下的每个 head dim,GQA 分页前向去掉旧 decode 核(#2807 / #2797)(11-01)
- 3.4 TileOPs:coop2 生产者收成单 warp、新增 192 宽转置块,W4A8 部分一波改成一 CTA 一 tile(#2806 / #2805)(11-01)
- 3.5 TileOPs:编译缓存按构建目标键控,算子层统一检查输出与输入重叠(#2802 / #2801)(11-01)
- 3.6 TileOPs 在修缺陷:Norm 四算子在远离零的行丢精度、约 1e36 溢出成 NaN(#2812)(11-01)
- 3.7 TileFoundry:lowering 把 issuing mesh 与窗口切分收进一处决策,83 个调度夹具字节一致(#280 / #282)(10-31 ~ 11-01)
- 3.8 TileFoundry:站点发布 0.3 版调度文章《Scheduling, Designed for Agents》(#4 / #281)(10-31 ~ 11-01)
- 3.9 采用方:FlashQLA 合入 SM120 的 Dual-CP 前反向,2 张 RTX 5090 上对 32K 至 128K 序列与 Inter-CP 对比(#51)(10-30)
- 3.10 采用方:TileKernels 自 10-14 起无推送(11-02)
- 四、社区、教程与活动
- 4.1 RISC-V 研究链成形:三个仓把 Numba CUDA Python 转 TileLang,再把受限 PrimFunc 落到单 hart RVV C(10-31)
- 4.2 PTO-ISA/tilelangGym 从 PTO-ISA/tilelang 独立成仓,主回归门是昇腾 SimtVF 路径(10-30)
- 4.3 新闻侧:Google News 中英文查询、Hacker News、arXiv 均无窗口内实质命中(11-02)
- 五、趋势观察
- 附、素材与核查说明
今日重点
这个窗口最重的动作是减法。主线把三周前合入的精确线程计数(#3291)整笔回退(#3733),理由写得很直白:那套改动把昂贵的计数查询搬到了满块栅栏检查之前,即使所有线程都会到达栅栏也要跑证明,当前 TVM 的求解预算会被耗尽并返回部分计数,于是本来合法的 remap 与 sequence 辅助 kernel 直接以「Cannot prove a thread rectangle」编译失败。回退在两套 ThreadSync 实现上同时进行,删掉 #3291 引入的测试,补一个「循环对多数线程为空、256 线程仍全部到达栅栏」的归约用例;商守卫、精确参与线程界、非矩形线程集这些正确性改进也一起撤回,新开的 #3732 留作迁移后重新优化的跟踪单。这是明确标注为临时措施的迁移让路:TileLang 正在准备 TIRX 迁移以摆脱当前 TVM fork 依赖,配套的 TVM fork 优化提案已经关闭。
第二条线是昇腾在补地基。#1948 修的是宿主侧一笔与算子规模无关的固定损耗:Cython wrapper 每次调用都无条件构造 TVM Analyzer,实测在 910B3 上约 0.28 毫秒一次,而它只在形状为符号表达式时才用到;改成惰性构造后,vec add 单块从 442 至 467 微秒降到 105.1 微秒,32 块从 441 至 493 微秒降到 108.0 微秒,kpool 压缩从 447 至 450 微秒降到 111.3 微秒,作为下限的裸 ctypes 调用约 12 微秒没变。同期 #1971 的 RFC 给出同步实现的替换路线:移植主线同步核心加独立的 Z3 证明服务,用旧 TIR 适配层接进现有编译流水,保留 TVM、调度、内存规划、AscendC/PTO codegen 与运行 ABI,验收通过后再切默认实现并删除旧算法——与 #1975 把整个 Ascend 仓迁到主线多后端架构是一件事的两个落点。
第三条线在算子库。TileOPs 在 11-01 一天合入 13 笔,几乎全部是性能补丁,且每笔都附 H200 实测表与 manifest 回归门槛。其中 #2808 把 GQA 稠密前向的程序选择收成一条规则,相对最快候选程序的选择遗憾从 1.818 降到 1.017(2 至 128 个 query 位置、586 次调用),并修掉主线一个真实错误:score scale 在乘上 log2(e) 后于 FP32 下消失,输出变成 NaN 或错值。
一、核心项目进展
1.1 主线:ThreadSync 的精确线程计数回退,为 TIRX 迁移让路(#3733)
合于 11-01 14:50。被回退的 #3291 把精确线程计数移到满块栅栏检查之前,并加入矩形认证,代价是即使全部线程都到栅栏也要跑一遍 CountSatisfyingValues;TVM 侧求解预算耗尽后返回部分计数,remap 与 sequence 辅助 kernel 随即以「[ThreadSync] Cannot prove a thread rectangle」失败。回归在主线 cc903dcb9 加未改动的 TVM pin 4211874e9e 上复现,新用例把它缩成「循环对多数线程为空、256 线程仍全部到达栅栏」。回退同时覆盖共享与昇腾两套实现,恢复满块快速路径与旧的部分计数行为;原文明确这是 TIRX 迁移期间的临时缓解,TVM fork 侧的优化提案 tile-ai/tvm#89 已关闭,后续优化转入 #3732。
来源:https://github.com/tile-ai/tilelang/pull/3733
1.2 主线:CUDA 跨 warp 全归约合并为一次共享内存交换,256 线程从 6 道栅栏降到 2 道(#3715)
合于 10-30 20:26。旧实现在 CUDA 归约头文件 reduce.h 里从 threads/2 递归到 scale,凡是 offset 不小于 32 的每一步都做一次完整的共享内存写读交换并配两道栅栏:256 线程块的三步(128、64、32)合计 3 次共享内存往返、6 道栅栏,512 线程块是 4 次、8 道。改法是把所有跨 warp 步骤压成一次交换(每元素一次 st.shared、共 2 道栅栏):先在 warp 内做 scale 小于 32 的那一段,再让每个 warp 内组里不超过 cross_radix 的若干 lane 协作折叠,寄存器压力有上界。标量与批量两条路径都改了。
来源:https://github.com/tile-ai/tilelang/pull/3715
1.3 主线:PTO 后端的 T.Parallel 向量化落地第一里程碑,逐元素算子走 VMI 指令(#3512)
合于 10-30 01:29。目标是让用户在 T.SimdVF(lanes=L) 里写标量并行的逐元素 T.Parallel,由编译器向量化成 PTO VMI 指令,再经 PTOAS 在昇腾 950 上执行。分工写得很清楚:TileLang 侧负责逻辑 lane 布局推导、向量化 pass 流水线、VMI 指令生成与 Fragment 逻辑映射;PTOAS 侧负责物理寄存器切分、掩码具体化、Fragment 缓冲与向量寄存器分配、mem_bar 插入。路线图分五个里程碑,本笔是 M1 核心框架(逐元素计算,覆盖 extent 能否被 lanes 整除两种情况),M2 控制流、M3 reduce、M4 UB 容量规划与寄存器分配、M5 非连续访问与更宽的 PACK 数据类型依次跟进。
来源:https://github.com/tile-ai/tilelang/pull/3512
1.4 主线:昇腾 DeepGEMM 的 scale 缓冲改用写计数多缓冲,数据加载可提到条件加载之前(#3629)
合于 10-30 05:48。DeepGEMM 会在多个 K 块之间复用 scale factor,本笔新增显式的写计数多缓冲模式,让 scale 缓冲的版本按真实写入推进:每个 owner 循环迭代里只在第一个实际执行的写任务处推进一次,后续写入共享该版本,没有写入的迭代沿用旧值,支持多个不相交 owner 与共享时钟。这样数据加载可以排到条件化的 scale 加载之前;同时把 DeepGEMM 的 K 循环摊平,在 L1 与 L0 上都先取数据再取 scale,并补了守卫更新、嵌套更新、保留读、flag 与缓冲复用以及 FP4 的 K 主序与 MN 主序带尾块和环形回绕的用例。该模式要求显式版本数,并假定每个 owner 迭代最多一次更新。
来源:https://github.com/tile-ai/tilelang/pull/3629
1.5 主线:两笔修复合入,CPU 标量 GEMM 只清目标区域、负索引合法化跳过不可达分支(#3709 / #3712)
两笔都在 10-30 合入,属上一期在审条目的落地。#3709 让 CPU 标量 GEMM 只清理目标区域,此前会连带清掉不属于本次输出的范围;#3712 让 LegalizeNegativeIndex 跳过不可达分支,避免对不会执行的代码做负索引改写。这两条与上一期报告的 #3708、#3710 属同一批缺陷清理,其中 #3708 的位宽改 int64、#3710 的关缓存时绕过两层进程内缓存仍在审。
来源:https://github.com/tile-ai/tilelang/pull/3709、https://github.com/tile-ai/tilelang/pull/3712
1.6 主线:语言层清掉 proxy 里两个无人使用的 Ref 占位(#3722)
合于 10-31 17:31。tilelang/language/proxy.py 里同时存在两个 Ref:一个给类型检查用的 tirx.Var 桩,一个空的运行时类,两者在树内都没有消费者;对外的 T.Ref 实际来自 eager.builder,那里的实现才是宏引用参数的本体。本笔删掉两个占位、_T 类型变量以及随之无用的 Generic、TypeVar 导入,保留 T.Ref 与宏引用行为,只移除 tilelang.language.proxy.Ref 这条无人使用的直接导入路径。前端构造测试 179 项通过、6 项跳过,运行时身份检查确认 T.Ref 仍是 eager.builder.Ref。
来源:https://github.com/tile-ai/tilelang/pull/3722
1.7 新开缺陷与在审:TMEM 静默丢结果、ROCm 符号找不到、TF32 就近偶数舍入(#3721 / #3711 / #3714)
#3721 是本窗口最值得警惕的一条:在 sm_100a 上对 TMEM 做 T.gemm,当 threads/32 不能整除 N 时,编译能过、PTX 里却一条 tcgen05.mma 都没有,T.copy(Ct, Cl) 读的是没人写过的 TMEM,乘积被静默丢弃;复现里 threads=256 出一条,threads=384 出零条,作者怀疑 src/cuda/op/gemm.cc 给 TCGEN05 返回 {1, num_warps} 分区时缺了覆盖性检查。#3711 报告 ROCm 7.14 与 10.0 的 PyTorch wheel 下 kernel 报「named symbol not found」。#3714 在审,为 fp32 到 tf32 的转换补上就近偶数舍入(cvt.rn.tf32.f32 与饱和形式),非饱和要求 SM90 以上、饱和要求 SM100 以上,关闭对应的功能请求 #3713。
来源:https://github.com/tile-ai/tilelang/issues/3721、https://github.com/tile-ai/tilelang/issues/3711、https://github.com/tile-ai/tilelang/pull/3714
二、多后端适配(昇腾 / 海光 / 摩尔线程 / 沐曦)
2.1 昇腾:Cython wrapper 惰性构造 TVM Analyzer,910B3 单次调用从约 450 微秒降到约 106 微秒(#1948)
合于 10-30 07:44。CythonKernelWrapper.forward() 此前每次调用都无条件构造一个 TVM Analyzer(),在 910B3 上是约 0.28 毫秒的固定宿主开销,而该对象只在 result 或 workspace 形状为符号表达式时才用到;静态 shape 的 kernel 完全用不上,却同样付这笔钱。改法是把构造延迟到 tir.PrimExpr 分支内并按需复用。实测:vec add 单块 442 至 467 微秒降到 105.1 微秒,32 块 441 至 493 微秒降到 108.0 微秒,kpool 压缩 447 至 450 微秒降到 111.3 微秒,裸 ctypes 调用约 12.5 微秒降到约 12 微秒。新增两例回归,分别验证静态路径不构造、符号路径按需构造。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1948
2.2 昇腾:RFC 提出把同步实现换成主线核心加独立 Z3 证明服务(#1971)
本 RFC 不预设提速结论,只固定比较基线:Ascend 仓在 TileLang 43b0b69 与 TVM c2921fd 上、IR 基础是旧 tvm.tir;主线在 TileLang 1c0f402 与 TVM 4211874e9e 上、IR 基础已是 tvm.tirx 与 tvm.s_tir,对象与注册体系也换成新 FFI 与反射设施。因此迁移不能简化为复制一个 InsertSync 文件。建议路线是移植主线同步核心与独立 Z3 证明服务,通过旧 TIR 适配层接入现有编译流水,保留现有 TVM、调度、内存规划、AscendC/PTO codegen 与运行 ABI,完成支持范围、正确性、性能与安装验收后再切换默认实现并删除旧同步算法。配套 PR 为 #1965。
来源:https://github.com/tile-ai/tilelang-ascend/issues/1971
2.3 昇腾在审:把 TileLang-Ascend 整体迁到主线多后端架构(#1975)
10-31 提出并持续更新,目标是按提案 #1911 把 TileLang-Ascend 迁移到 TileLang 主线的多后端架构上;描述里给出两份差异对照,一份对 ascendc_pto 分支,一份对主线。这条与 2.2 的 RFC 是同一条收敛路径的两个层面:一条改同步算法,一条改仓的骨架。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1975
2.4 昇腾在审:GLM-5.3 kpool 的 fp8_mqa_logits 示例,双 kernel 加 AIV/Cube 分工(#1966)
对照主仓 kpool 目录下的 GLM-5.3 fp8_mqa_logits 示例做昇腾版。算子语义是 GLM-5.3-Flash 稀疏注意力 indexer 的池化历史打分:对每个 query 行,在有效 pool 区间内按分页表从 64-pool 页读 E4M3 key(uint8 承载字节加 fp32 scale),反量化后算 32 个 indexer 头的 MQA 打分,乘 per-pool scale,经 ReLU gate 与按头加权归约为每个 pool 一个 FP32 logit,范围外或无效页输出负无穷。实现是双 kernel:mask kernel 按 64-pool tile 生成逐 lane 的 0 或负无穷掩码以替换宿主预填,main kernel 每行每页一个 block、两个 AIV 各处理 32-pool 半页,Cube 读整页做一次 64×32×128 的 GEMM。头数 32、维度 128、页大小 64。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1966
2.5 昇腾:四笔文档与 API 事实核查合入,clamp 在 PTO 上真正生效(#1899 / #1893 / #1620 / #1600)
四笔都在 10-30 合入,属持续进行的 API 事实核查批次:#1899 核对 T.tile.bitwise_xor 文档(#1681 的重提),#1893 核对 T.tile.bitwise_rshift(#1676 的重提),#1620 让 T.tile.clamp 在 PTO 上真正生效并同步文档,同时补测试与 API 文档,#1600 更新 axpy 的 docstring、加强测试并补 API 文档。同批次另有 #1968、#1967、#1974 在审,分别是 sigmoid 与 round 计数折半、符号形状解析缓存、整块行切分下 tile 尺寸被错误折半的修复。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1620、https://github.com/tile-ai/tilelang-ascend/pull/1899
2.6 昇腾:日测 2970 用例全部通过,11-01 有一次 workflow 失败(#1977 / #1976)
11-02 06:42 的日测记录为 2970 项、通过 2970 项、失败 0 项,通过率 100%;10-31 06:56 的一期同样是 2970 项全过。中间 11-01 05:08 有一期标记为 workflow test job failure,与本窗口内 #1975 的迁移改动、#1971 的 RFC 同期出现,属需要留意的波动而非稳定回归。
来源:https://github.com/tile-ai/tilelang-ascend/issues/1977、https://github.com/tile-ai/tilelang-ascend/issues/1976
2.7 海光:MLS 边界与持续 WASP 同步修复合入,gfx948 兼容在审(#31 / #32)
#31 合于 10-30 10:56,三件事:改进动态 MLS 边界推断与 MatrixStore 的 stride 处理;修 MMAC 与 blockscaled GEMM 的 clear_accum 下沉;在持续 WASP 路径改用 M 分区消费者与独立的 A/B 栅栏,同时保留 slot 与 phase 轮转并对尾部消费者加同步。#32 在审,把 gfx948 当作 gfx946 兼容处理,让既有指令能直接在新型号上走通。前一窗口报道的普通拷贝 LDS swizzle(#30)已在同期落地。
来源:https://github.com/tile-ai/tilelang-hygon/pull/31、https://github.com/tile-ai/tilelang-hygon/pull/32
2.8 mlir-ascend:CI 清掉内部缓存服务引用,Codex 工作流在审(#206 / #208)
#206 合于 10-31 07:24,把 CI 里指向内部缓存服务的内容清掉,属可复现性清理。#208 在审,为 Operator Harness 加首批 Codex 工作流:用迭代提问完成项目初始化、把已有算子仓接入、以及从持久化上下文恢复开发、调优或分析;三个 skill 共用 .workspace/ 下的可编辑 Markdown 记录,任务交接区分「当前候选」与「已验证的最佳候选」,保留失败与回退实验。作者明确本次没有真实 NPU 实测,NPU 执行适配与 npm 分发列为后续工作。
来源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/206、https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
2.9 摩尔线程与沐曦窗口内无推送(11-02)
tilelang-metax 最后推送停在 10-26,tilelang-musa 停在 09-30,两仓在本窗口的 72 小时内都没有提交。这意味着本期的后端动态集中在昇腾、海光与 mlir-ascend 三侧。
来源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa
三、生态与采用方(TileOPs / TileFoundry / TileKernels / FlashQLA)
3.1 TileOPs:72 小时 15 笔合入,13 笔压在 11-01,几乎每笔都带实测表与清单回归门槛(11-01)
v0.1.0 发布后,TileOPs 在 11-01 一天合入 13 笔,加上 10-30 的 #2794 与 10-31 的 #2795,72 小时共 15 笔。改动几乎全部是 perf 类,覆盖注意力、GEMM、逐元素、归约与 MoE,每一笔都给出 H200 上的 busy 与 back-to-back 两组时间、或与 vLLM、FlashInfer 的对照,并声明 manifest 里没有已测行回退。这种「带表合入」的方式说明算子库已经进入按 shape 逐行校准的阶段,而不是补功能。
来源:https://github.com/tile-ai/TileOPs/commits/main
3.2 TileOPs:GQA 稠密前向用一条实测规则选 decode 或 prefill 程序,选择遗憾从 1.818 降到 1.017(#2808)
合于 11-01 13:04。GQADenseFwdOp 与 GQADenseWithLSEFwdOp 改由同一个规则 AttentionCall.dense_decode_program 决定走哪套程序,输入是 query 与 key 长度、batch、头分组、head dim、RoPE 与 SM 数:2 至 64 个 query 位置走 warp 专用 decode,一个 KV 头组内每个位置打包进 64 行块;单个 query 位置时按 32 段切分填满一波、或走 1 序列 1 至 2 个 KV 头的 batch-1 程序。相对最快可用程序的选择遗憾(几何均值与最大值):单位置 1.068 与 1.87 降到 1.018 与 1.69,2 至 128 位置 1.818 与 25.2 降到 1.017 与 1.81。同时修掉主线两个错误:score scale 乘上 log2(e) 后在 FP32 下消失导致 NaN 或错值;D256 多位置调用在 SM90 上因共享内存不足启动失败。
来源:https://github.com/tile-ai/TileOPs/pull/2808
3.3 TileOPs:分页 MHA decode 扩到共享内存装得下的每个 head dim,GQA 分页前向去掉旧 decode 核(#2807 / #2797)
#2807 合于 11-01 12:30,把分页 MHA decode 在「一个块的共享内存装得下」的前提下扩到每个 head dim,此前遇到较大 head dim 会直接不走这条路径。#2797 合于同日 03:54,从 GQAPagedFwdOp 里删掉旧的 paged decode kernel,并把分页 MHA decode 的上限压到能编译出来的 head dim 上,等于把前一版留下的越界能力收回。两条合起来是一次「先扩能力、再削冗余」的收尾。
来源:https://github.com/tile-ai/TileOPs/pull/2807、https://github.com/tile-ai/TileOPs/pull/2797
3.4 TileOPs:coop2 生产者收成单 warp、新增 192 宽转置块,W4A8 部分一波改成一 CTA 一 tile(#2806 / #2805)
#2806 合于 11-01 12:02,把 Hopper GEMM 的 coop2 生产者从多 warp 收成一个加载 warp,并在共享内存允许三级流水时提供 192 宽的均匀转置块;H200 上 wgrad-q-b 从 384.9 微秒降到 382.3,mlp-up fp16 从 317.9 微秒降到 311.9(back-to-back 分别 458.5 到 449.6、369.0 到 367.3),作者声明没有已测 GEMM 行回退,本笔属 #2791 目标的一部分,双卡确认仍未完成。#2805 合于同日 10:36,让填满至少四分之三 SM 的部分一波 W4A8 prefill 栅格改成一 CTA 一 tile 而非 stream-K;llama-8b-up-b33 从 23.52 微秒降到 17.98(vLLM 20.42),b128 从 31.33 降到 23.39(vLLM 24.05),清单外 M 33 至 128 的 44 组 N/K 扫描中改调度形状快 5% 至 36%、中位 18%、无一变慢。
来源:https://github.com/tile-ai/TileOPs/pull/2806、https://github.com/tile-ai/TileOPs/pull/2805
3.5 TileOPs:编译缓存按构建目标键控,算子层统一检查输出与输入重叠(#2802 / #2801)
#2802 合于 11-01 11:31,修的是编译程序缓存跨设备复用的问题:构建器与 Op 条目都按构建目标做键,懒编译与自动调优留在预期设备上,验收为 99 项目标隔离、供应商与分组 GEMM 测试。#2801 合于同日 10:33,把「输出不得与任何其他调用张量共享字节」的检查提到算子层,由 manifest 的 overlaps 字段按算子声明例外,删掉了 MoE unpermute、MoE 分组 GEMM、持续 GEMM 模板等四处 kernel 侧检查;顺带修掉主线一个竞态:IndexedExpertMLPFwdOp 允许 out 与 topk_weights 重叠,而后者在写入时仍被归约读取。MoEPostPermuteFwdOp 带 out 的宿主耗时从 51.48 微秒降到 39.33。
来源:https://github.com/tile-ai/TileOPs/pull/2802、https://github.com/tile-ai/TileOPs/pull/2801
3.6 TileOPs 在修缺陷:Norm 四算子在远离零的行丢精度、约 1e36 溢出成 NaN(#2812)
11-01 新开。LayerNorm、AdaLayerNorm(含 -Zero)、GroupNorm 与 BatchNorm 目前都用原始行的 fp32 直接求和取均值与方差(BatchNorm 用 E[x²] 减均值平方),这正是 #2773 为 FusedAddLayerNorm 修过、#2811 为 InstanceNorm 训练核在修的同一类缺陷。表现是:常数 10000.1 的 fp32 行归一化后输出约等于正负 weight 而不是 bias,10000 加随机数的行在 fp32 容差下对不上 float64 参考;接近 1e36 的行求和溢出,fp32 与 bf16 下全部输出 NaN。受影响路径覆盖 fragment、共享内存走查、一行一 warp、分块与流式等多种实现。
来源:https://github.com/tile-ai/TileOPs/issues/2812
3.7 TileFoundry:lowering 把 issuing mesh 与窗口切分收进一处决策,83 个调度夹具字节一致(#280 / #282)
#280 合于 10-31 00:47,把「每次调用发射什么」的决策收进一处;#282 合于 11-01 11:47,继续把「由哪个 mesh 发射」与「窗口怎么切」也收到一处:analysis/emission.py 新增 copy_mesh、issue_owner、issue_mesh,copy_mesh 与 issue_owner 共用一个助手,两者唯一的差别(Partial 怎么读)只留在一处;三处重复的 Slice 切分改由 _cut 助手统一。验收方式是字节级对照:83 个调度夹具里 75 个 finalize 输出与主线完全一致,其余 8 个在主线上同样失败;docs/spec/passes.md 第 7.1 节把这套 mesh 决策写进规范。
来源:https://github.com/tile-ai/TileFoundry/pull/282、https://github.com/tile-ai/TileFoundry/pull/280
3.8 TileFoundry:站点发布 0.3 版调度文章《Scheduling, Designed for Agents》(#4 / #281)
站点仓在 10-30 合入 #4,发布 0.3 版文章《Scheduling, Designed for Agents》,同时补了文章与配图的目录结构、公式的 MathJax 支持、图注与表格样式以及每个图表锚点,首页 Blog 卡片改为打开博客索引,索引列出两篇文章;upstream-ref 未变,站点仍渲染 v0.0.3。仓库侧 #281 于 10-31 在 README 里通告这篇文章。对 TileFoundry 这条线来说,这是把「面向 agent 的调度」从代码与发布说明搬到公开叙述的一步。
来源:https://github.com/tile-ai/TileFoundry.github.io/pull/4、https://github.com/tile-ai/TileFoundry/pull/281
3.9 采用方:FlashQLA 合入 SM120 的 Dual-CP 前反向,2 张 RTX 5090 上对 32K 至 128K 序列与 Inter-CP 对比(#51)
10-30 03:02 合入。改动为 SM120 增加 Dual-CP 的前向与反向:前向复用已有的 state-scan kernel 做卡级状态聚合,有序 M 乘积按 32 列一块计算以控制在共享内存内;反向用转置后的 M 逆序传播状态,复用前向缓存下来的 M 矩阵;同时修掉反向 kernel 里一处共享内存竞态,办法是让第 7 阶段与 dV 存储 warp 同步。性能在 2 张 RTX 5090 上以 BF16、chunk 32、序列长 32K 至 128K 对比 Inter-CP,每例 3 次中位数、每次 25 次预热加 100 次计时。
来源:https://github.com/QwenLM/FlashQLA/pull/51
3.10 采用方:TileKernels 自 10-14 起无推送(11-02)
深度求索的 TileKernels 最后推送停在 10-14,本窗口无变化。采用方这一侧本期的有效动态全部来自通义 FlashQLA。
来源:https://github.com/deepseek-ai/TileKernels
四、社区、教程与活动
4.1 RISC-V 研究链成形:三个仓把 Numba CUDA Python 转 TileLang,再把受限 PrimFunc 落到单 hart RVV C(10-31)
社区研究者 zishu0131-png 在 10-31 同日整理出三个仓:tilelang-riscv 作总入口,说明研究目标、两段编译链路与版本组合;numba2tilelang 负责把 Numba CUDA Python 静态解析后转成 TileLang 源码;tilelang2rvv 把受支持的真实 tvm.tirx.PrimFunc 转成 RVV intrinsic C,再经 Clang 编成 RV64GCV ELF,在 QEMU virt 上以单 hart 运行。作者把边界写得很清楚:这不是 TileLang 官方注册的后端,也不是通用 GPU 模型模拟器,依赖 threadIdx 的寻址、跨 hart 同步与未支持的 IR 一律拒绝,block 顺序执行,shared 与 fragment 映射为单 hart 局部数组;研究环境为 TileLang 0.1.15、Clang 23.1.3、QEMU 11.1.0,不保证其他版本兼容,且明确上游线程级产物不能自动视为该后端可接受的输入。
来源:https://github.com/zishu0131-png/tilelang-riscv
4.2 PTO-ISA/tilelangGym 从 PTO-ISA/tilelang 独立成仓,主回归门是昇腾 SimtVF 路径(10-30)
这套敏感性与回归测试集从 PTO-ISA/tilelang 的 examples/ascend/t_parallel_st_suite 拆出来,理由是让测试集能独立钉住 TileLang 与 PTOAS 的版本。当前三条主路径:昇腾 SoftLib SimtVF(target=ascend、cython,主回归门,在 remat 版本上全过)、VectorizeParallelToPTO 与 VL 的 vload/vstore,以及 benches 下用 @pto.jit 写的手工 SIMD DSL 对照实现;每个用例目录里带 config.toml 与 VERDICT.md,另有周度与 CI 回归脚本。
来源:https://github.com/PTO-ISA/tilelangGym
4.3 新闻侧:Google News 中英文查询、Hacker News、arXiv 均无窗口内实质命中(11-02)
三天的 Google News 中英文查询没有返回任何落在窗口内的 TileLang 相关报道,Hacker News 与 arXiv 同样没有新条目;候选清单里的 20 条全部来自 GitHub 仓库推送,属常规代码活动而非新闻。本期的信息面因此完全由仓库侧与社区仓构成。
五、趋势观察
这一窗口最值得记录的判断来自上游自己:宁可回退正确性改进,也要把 ThreadSync 的复杂度从关键路径上撤下来。回退说明里写清了代价——商守卫、精确参与线程界、非矩形线程集这三类改进一并撤回,只保留旧的保守路径,#3732 留待 TIRX 迁移完成后再谈优化。一个编译器项目愿意公开写「我们的证明器预算不够,先退回保守实现」,比再堆多少功能都更能说明当前阶段的真实约束:TIRX 迁移是主线第一优先级,其他优化都要给它让路。
与之呼应的是昇腾侧两份「向主线收敛」的文本:#1971 的 RFC 要把同步算法换成主线核心加独立 Z3 证明服务,先调研、后替换、中间留旧 TIR 适配层;#1975 直接把整个 Ascend 仓迁到主线多后端架构。两条路都要付出迁移成本,于是近期的可见工作会偏向地基类:#1948 把每次调用的固定宿主开销从约 450 微秒压到约 106 微秒,四笔文档事实核查,CI 清理内部缓存引用,都是这类工作。
TileOPs 的节奏则从「发版」切到「调优」。72 小时 15 笔合入里 13 笔压在同一天,几乎每笔都附 H200 实测表与「manifest 无回退」声明,并且开始用选择遗憾这类指标描述调度规则的好坏(#2808 从 1.818 到 1.017)。这意味着算子库的竞争点已经从「有没有这个算子」变成「每条 shape 上是否都挑对了程序」,而这恰好是 TileLang 这类 DSL 相对手写 kernel 的长期价值所在。
生态侧出现两条新分叉。一条是硬件方言:海光把 gfx948 当 gfx946 兼容、昇腾继续补 PTO 文档与符号形状。另一条是语言出口:RISC-V 研究链把受限 PrimFunc 降到单 hart RVV C,星标很低、作者也明确声明不是官方后端,但它把 TileLang 的 IR 当作可移植资产来用,与昇腾、海光的仓级分支适配是同一类信号。多条分叉同时存在,上游 TIRX 迁移能否给出稳定的 IR 边界,将决定这些下游投入是沉淀还是返工。
附、素材与核查说明
信源核查表
| 信源 | 核查结果 |
|---|---|
| tile-ai 组织(29 仓库) | 窗口内推送 10 个仓库 |
| 主仓 tilelang | 主分支 7 笔合入;新开 issue 4 条;文档站 3 笔 |
| TileOPs | 主分支 15 笔合入(13 笔在 11-01);新开 issue 3 条;站点 1 笔 |
| TileFoundry | 3 笔合入;站点 1 笔 |
| tilelang-ascend | 5 笔合入;新开 issue 4 条(含日测记录与 RFC) |
| tilelang-hygon | 2 笔合入;1 笔在审新增 |
| tilelang-mlir-ascend | 1 笔合入;1 笔在审新增 |
| tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack | 窗口内无推送 |
| 发布页(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / mlir-ascend) | 窗口内无新版本 |
| FlashQLA | 1 笔合入(#51) |
| TileKernels | 自 10-14 起无推送 |
| RISC-V 研究链三仓 / PTO-ISA/tilelangGym | 有推送,研究链三仓同日建链 |
| Google News / Hacker News / arXiv | 无窗口内实质命中 |
完整信源清单
- [1] tilelang ThreadSync 回退(#3733)— https://github.com/tile-ai/tilelang/pull/3733
- [2] ThreadSync 计数回归跟踪(#3732)— https://github.com/tile-ai/tilelang/issues/3732
- [3] tilelang CUDA 跨 warp 全归约(#3715)— https://github.com/tile-ai/tilelang/pull/3715
- [4] tilelang PTO 的 T.Parallel 向量化(#3512)— https://github.com/tile-ai/tilelang/pull/3512
- [5] tilelang 昇腾 DeepGEMM 写计数多缓冲(#3629)— https://github.com/tile-ai/tilelang/pull/3629
- [6] tilelang CPU 标量 GEMM 清理范围(#3709)— https://github.com/tile-ai/tilelang/pull/3709
- [7] tilelang 负索引合法化跳过不可达分支(#3712)— https://github.com/tile-ai/tilelang/pull/3712
- [8] tilelang 语言层清理 Ref 占位(#3722)— https://github.com/tile-ai/tilelang/pull/3722
- [9] tilelang TMEM 静默丢结果(#3721)— https://github.com/tile-ai/tilelang/issues/3721
- [10] tilelang ROCm 符号找不到(#3711)— https://github.com/tile-ai/tilelang/issues/3711
- [11] tilelang TF32 就近偶数舍入(#3714)— https://github.com/tile-ai/tilelang/pull/3714
- [12] tilelang 浮点转 TF32 舍入功能请求(#3713)— https://github.com/tile-ai/tilelang/issues/3713
- [13] tilelang-ascend Cython wrapper 惰性构造 Analyzer(#1948)— https://github.com/tile-ai/tilelang-ascend/pull/1948
- [14] tilelang-ascend 自动同步与 Z3 迁移 RFC(#1971)— https://github.com/tile-ai/tilelang-ascend/issues/1971
- [15] tilelang-ascend 迁移到主线多后端架构(#1975)— https://github.com/tile-ai/tilelang-ascend/pull/1975
- [16] tilelang-ascend GLM-5.3 kpool fp8_mqa_logits 示例(#1966)— https://github.com/tile-ai/tilelang-ascend/pull/1966
- [17] tilelang-ascend clamp 在 PTO 生效(#1620)— https://github.com/tile-ai/tilelang-ascend/pull/1620
- [18] tilelang-ascend bitwise_xor 文档核查(#1899)— https://github.com/tile-ai/tilelang-ascend/pull/1899
- [19] tilelang-ascend bitwise_rshift 文档核查(#1893)— https://github.com/tile-ai/tilelang-ascend/pull/1893
- [20] tilelang-ascend axpy 文档与测试(#1600)— https://github.com/tile-ai/tilelang-ascend/pull/1600
- [21] tilelang-ascend 日测 100% 通过(#1977)— https://github.com/tile-ai/tilelang-ascend/issues/1977
- [22] tilelang-ascend workflow 失败记录(#1976)— https://github.com/tile-ai/tilelang-ascend/issues/1976
- [23] tilelang-hygon MLS 边界与 WASP 同步(#31)— https://github.com/tile-ai/tilelang-hygon/pull/31
- [24] tilelang-hygon gfx948 兼容(#32)— https://github.com/tile-ai/tilelang-hygon/pull/32
- [25] tilelang-mlir-ascend CI 清理内部缓存引用(#206)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/206
- [26] tilelang-mlir-ascend Codex 工作流(#208)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
- [27] TileOPs GQA 稠密前向程序选择(#2808)— https://github.com/tile-ai/TileOPs/pull/2808
- [28] TileOPs 分页 MHA decode 扩 head dim(#2807)— https://github.com/tile-ai/TileOPs/pull/2807
- [29] TileOPs GQA 分页前向去掉旧 decode 核(#2797)— https://github.com/tile-ai/TileOPs/pull/2797
- [30] TileOPs coop2 单 warp 与 192 宽块(#2806)— https://github.com/tile-ai/TileOPs/pull/2806
- [31] TileOPs W4A8 部分一波一 CTA 一 tile(#2805)— https://github.com/tile-ai/TileOPs/pull/2805
- [32] TileOPs 编译缓存按目标键控(#2802)— https://github.com/tile-ai/TileOPs/pull/2802
- [33] TileOPs 算子层重叠检查(#2801)— https://github.com/tile-ai/TileOPs/pull/2801
- [34] TileOPs Norm 精度与溢出缺陷(#2812)— https://github.com/tile-ai/TileOPs/issues/2812
- [35] TileOPs 主分支提交记录 — https://github.com/tile-ai/TileOPs/commits/main
- [36] TileFoundry issuing mesh 与窗口切分(#282)— https://github.com/tile-ai/TileFoundry/pull/282
- [37] TileFoundry 发射决策收拢(#280)— https://github.com/tile-ai/TileFoundry/pull/280
- [38] TileFoundry 0.3 调度文章发布(#4)— https://github.com/tile-ai/TileFoundry.github.io/pull/4
- [39] TileFoundry README 通告调度文章(#281)— https://github.com/tile-ai/TileFoundry/pull/281
- [40] FlashQLA SM120 Dual-CP(#51)— https://github.com/QwenLM/FlashQLA/pull/51
- [41] 深度求索 TileKernels — https://github.com/deepseek-ai/TileKernels
- [42] RISC-V 研究链总入口 — https://github.com/zishu0131-png/tilelang-riscv
- [43] Numba 到 TileLang 转换仓 — https://github.com/zishu0131-png/numba2tilelang
- [44] TileLang PrimFunc 到 RVV C 消费器 — https://github.com/zishu0131-png/tilelang2rvv
- [45] PTO-ISA/tilelangGym — https://github.com/PTO-ISA/tilelangGym
- [46] tilelang 发布页 — https://github.com/tile-ai/tilelang/releases
- [47] TileOPs 发布页 — https://github.com/tile-ai/TileOPs/releases
- [48] TileFoundry 发布页 — https://github.com/tile-ai/TileFoundry/releases
- [49] tilelang-metax — https://github.com/tile-ai/tilelang-metax
- [50] tilelang-musa — https://github.com/tile-ai/tilelang-musa
- [51] Google News 检索 — https://news.google.com/