TileLang 每日动态报告(2026-11-06)
调研窗口:2026-11-05 07:00 ~ 2026-11-06 07:00(北京时间,24 小时;上一期截点为 11-05 07:00)。
信源:GitHub tile-ai 组织逐仓核查(29 仓库,窗口内推送 10 个)+ 合入与在审真值检索(窗口内合入 18 笔、在审更新 57 笔)+ 新增 issue 5 条 + 主分支逐笔提交核查(tilelang 5、TileOPs 9、TileFoundry 1、tilelang-ascend 2、tilelang-musa 1)+ 非默认分支核查(tilelang-hygon 的 feat/hcu-scaled-cast 分支 3 笔)+ 发布页核查(tilelang / TileOPs / TileFoundry 三仓,窗口内无新版本)+ 采用方两仓 + 文档站与夜间仓核查 + Google News 中英文四组查询(走代理)+ Hacker News + arXiv。RISC-V 四源专项巡检本窗口无变化。
本期索引
- 今日重点:TileOPs 一天九笔合入,把「广播就地读、单波次启动、瓦片选择」三件事同时推进;主仓五笔合入之外,一条从 0.1.12 起就存在的静默错误 CUDA 缺陷被报出;昇腾修掉一个段错误,同时被指出 threads=2 下 sin/cos 有 97.1% 元素算错
- 一、核心项目进展
- 1.1 主线:TF32 转换支持 round-to-nearest-even,SM90 与 SM100 分开诊断(#3714)(11-05)
- 1.2 主线:host-mapped 缓冲可显式注解,CPU 张量不再被设备检查拒绝(#3802)(11-05)
- 1.3 主线:标量条件里的 cast 不再限制数据向量宽度(#3759)(11-05)
- 1.4 主线:Z3 的 rlimit 可按 kernel 配置,默认仍为十万(#3752,依赖 tile-ai/tvm #91)(11-05)
- 1.5 主线:嵌套 T.macro 复用转换模板,前端重复编译开销收敛(#3760)(11-05)
- 1.6 新 issue:条件标量化后缺 __syncthreads,0.1.12 至 0.1.15 静默生成错误 CUDA(#3808)(11-05)
- 1.7 新 issue:无 warp specialization 时 T.Pipelined 在 tcgen05 等待前重载 stage(#3800)(11-05)
- 1.8 在审:迁到原生 Apache TIRx 的完整迁移草案(#3803);共享转置改走 ldmatrix/stmatrix(#3809)(11-05)
- 二、多后端适配
- 2.1 昇腾:bitwise_rshift 的 64 位 dtype 改为显式报错,不再段错误(#1954)(11-05)
- 2.2 昇腾:A2/A3 性能技能对齐 Ascend950 工作流(#1932)(11-05)
- 2.3 昇腾:threads=2 下 tmp workspace 被无条件减半,sin/cos 在间接寻址 tile 上 97.1% 元素算错(#1991 / #1992)(11-05)
- 2.4 海光:分支三笔 LDS 相位修,并首次开出两个 PR(#33 / #34)(11-05)
- 2.5 摩尔线程:README 文档更新,结束自 09-30 以来的静默(11-05)
- 2.6 沐曦与 mlir-ascend:沐曦连续第三个窗口无推送;mlir-ascend 三项 TLIR 算子在审(11-05)
- 三、生态与采用方
- 3.1 TileOPs:九笔合入,广播操作数就地读、GEMM 单波次填充、reduction 布局缓存(11-05)
- 3.2 TileOPs:roofline 计价规则重写,三条不可能的带宽行回到合理区间(#2830)(11-05)
- 3.3 TileOPs:两笔合入相隔 2.5 分钟导致 main 的 GPU Smoke 失败,当日补修(#2837)(11-05)
- 3.4 TileOPs 在审:三阶段瓦片因子、TN 单元素收缩、FP8 分组 GEMM 的 TMA 缩放加载(#2836 / #2838 / #2839)(11-05)
- 3.5 TileFoundry:WGMMA 同步计划合入(#286),部分归约的完成语义在审(#288)(11-05)
- 3.6 TileSight:H20 的 WGMMA 能力声明修正,为仓库首个 PR(#1)(11-05)
- 3.7 采用方与第三方:TileKernels、FlashQLA 无推送;两个新的第三方仓出现(11-05)
- 四、社区、教程与活动
- 4.1 文档站三次更新;夜间仓无主分支提交;昇腾每日测试 3030 项全过(11-05)
- 4.2 新闻侧:四组查询、Hacker News 与 arXiv 均无窗口内命中(11-05)
- 五、趋势观察
- 附、素材与核查说明
今日重点
这一天最密集的产出来自算子库而不是编译器。TileOPs 在 21:06 到 22:46 的一个多小时里连续合入六笔,加上 14:01 与 21:31 的两笔,一天九笔,全部是 perf、fix、ci 标签。三件事同时推进:一是广播操作数就地读,lerp、where、clamp、masked_fill 五类算子不再把 stride 0 的广播张量先展开成输出形状,实测广播行耗时降到主线的 0.13 到 0.44 倍(#2832);二是单波次启动,一元逐元素在「体内最多两条便宜指令、读写总量不超过 L2 的四分之三」时缩成一个 256 线程的波次,并且顺手修正了 ReLU 对 NaN 返回 0 的老问题(#2829);三是瓦片选择器,给小 M 的 NT GEMM 造了一个宽度能填满一波的原生宽度瓦片,decode-down 从落后变成领先最快对比项 3%(#2834)。三件事的共同点不是「更快」,而是每笔都带着可复现的 A/B 表、都标出自己仍然落后的那一行。
第二条是可信度问题,而且横跨两个仓。主仓新开 issue #3808 报告:从 0.1.12 到 0.1.15,条件相关的标量化之后,共享内存 epilogue 少了一道 __syncthreads(),编译器不报错、生成的 CUDA 是错的;报告者给了完整复现脚本与四个版本号,属于「静默错代码」而非性能退化。同一窗口里 #3800 报的是另一种静默错:不开 warp specialization 时,T.Pipelined 会在 T.tcgen05_gemm 的 barrier 等待之前重载下一 stage,于是 MMA 读的 operand 已经被覆盖。这两条与主仓当天合入的 #3759 正好构成一组对照——#3759 修的是「标量条件里的 int64 cast 把数据向量宽度从 16 lane 压到 4 lane」,同一个标量化/向量化交界处,一边在修宽度,一边在漏同步。
第三条在昇腾侧。tilelang-ascend 合入 #1954,把 T.tile.bitwise_rshift 在 int64/uint64 上的编译期段错误改成显式 ValueError,只允许 int16/uint16/int32/uint32;这不是新增 64 位右移支持,而是把「编译器崩掉」变成「编译器说不」。同一仓当天新开两条同题 issue(#1991、#1992),报告在 910_9362 / CANN 9.0.0、threads=2(C:V=1:2)下,AscendVidReduction 把 kernel 内所有 UB 分配无条件减半,连注入的 tmp workspace 也减半,结果是 sin/cos 作用在间接寻址的 skip-set buffer 上时 3978/4096(97.1%)个元素算错,sigmoid/round 则违反 AscendC 的 calCount <= tmp 契约、越界写落在布局后方的 scratch 上。与之并置的是同日 #1988 的每日测试报告:3030 项全过、通过率 100%。这两组数字放在一起说明一件事——现有测试覆盖不到「threads=2 加间接寻址」这种组合,通过率不能当作正确性的替代指标。
第四条是生态面的动静。海光仓从「功能分支上攒了四笔、未开 PR」变成正式进入评审:feat/hcu-scaled-cast 分支当天又落三笔 LDS 相位修(默认 swizzle 适配 64 bank、MN 切片消费端恢复 swizzle 规划、K 分区的布局感知读),并开出 #33(scaled cast)与 #34(reduction 的 lane swap 与约束传播)两个 PR。摩尔线程结束了自 09-30 以来的完全静默,主分支落了一笔 README 的安装与算子开发说明更新。加上 TileSight 仓库出现首个 PR(修正 H20 的 WGMMA 能力声明)与主仓 #3803 的 TIRx 迁移草案,可以看到组织层面的三条线在同时展开:国产后端从自研分支进入公共评审、性能建模从论文走向仓库、编译器底座准备换掉 vendored TVM。
一、核心项目进展
1.1 主线:TF32 转换支持 round-to-nearest-even,SM90 与 SM100 分开诊断(#3714)
合于 12:09。此前 T.cast 到 tfloat32 只有默认的 TF32 就近舍入路径,想要 ties-to-even 得自己导入 CUDA helper。这一笔给 round="rn" 加了原生 lowering:非饱和形式发 cvt.rn.tf32.f32,饱和形式发 cvt.rn.satfinite.tf32.f32,并且绕开默认的 TF32 舍入构造函数、直接保留转换后的位模式。标量与向量两条路径都覆盖,TF32 的 lane 访问走打包载体。架构诊断按形式分开:非饱和形式要求 SM90 以上,饱和形式要求 SM100 以上,不满足就在编译期报错;没有显式指定 rounding 的 cast 行为不变。报告者是关掉 #3713 的需求方,属于「用户不必再自带 helper」这类接口收敛。
来源:https://github.com/tile-ai/tilelang/pull/3714
1.2 主线:host-mapped 缓冲可显式注解,CPU 张量不再被设备检查拒绝(#3802)
当天创建、当天合入(12:08 提交、16:02 合入)。CUDA kernel 可以直接访问映射进 GPU 地址空间的 host 内存,但 TileLang 原有的设备检查假定张量参数都住在目标设备上,于是这类 CPU 张量在启动前就被拒。新增的 T.annotate_host_mapped_buffers(*buffers) 把被注解参数的索引写进 PrimFunc 的 tl.host_mapped_params 属性,支持多缓冲、合并重复声明,并且能穿过 IR 序列化与 kernel 缓存恢复。Cython 与 TVM-FFI 两个 launcher 对被注解参数接受 CPU 张量,其余 dtype、形状、秩、步长检查照旧;设备侧参数的检查仍对齐当前执行设备,host 张量则不提供 CUDA 设备号。运行时新增一个原生 helper,用 cuMemHostGetFlags 校验注册状态、用 cuMemHostGetDevicePointer 解析设备可见地址,普通可分页内存仍会被拒。
来源:https://github.com/tile-ai/tilelang/pull/3802
1.3 主线:标量条件里的 cast 不再限制数据向量宽度(#3759)
合于 19:27,标签里带 ascend。起因是 #3734 在 DecoupleTypeCast 之后正确地收紧了数据 cast 的宽度限制,但把条件里的 cast 也算进去了:一段 T.assume(n % 16 == 0) 保护的 BF16 拷贝,因为条件里有 T.Cast("int64", ...),在 b5cb02633 上从 16 lane 被压到 4 lane,即使条件在 16 个 lane 上一致。修法是把 planner 已有的条件一致性分析复用起来,把 cast 宽度豁免从缓冲索引扩展到条件:if 语句、tirx.if_then_else 与标量谓词的 Select 的条件单独访问,分支仍然受限;在 lane 级向量谓词的目标上仍保留条件 cast 的约束。stochastic rounding 的 rbits 分组在跳过宽度约束时也被保留。这一笔与窗口内的两个新 issue 落在同一片区域,说明标量化与向量化的交界是当前缺陷的高发带。
来源:https://github.com/tile-ai/tilelang/pull/3759
1.4 主线:Z3 的 rlimit 可按 kernel 配置,默认仍为十万(#3752,依赖 tile-ai/tvm #91)
合于 23:05,前置是 fork 里的 tvm #91(23:03 合入)。改动把 Z3 求解器的 rlimit 从写死的 100,000 变成可按 kernel 定制:不配置时保持旧默认值。作者的表述很直接——「这间接修掉了很多与 Z3 证明器相关的问题」。这类改动的意义在于把「证明超预算」从不可调的环境行为变成可调的编译参数,与昨天昇腾侧收紧 AutoSchedule 周期下界(把 Z3 求解时间降 10.04%)是同一方向的两种做法:一个改搜索起点,一个改搜索预算。
来源:https://github.com/tile-ai/tilelang/pull/3752、https://github.com/tile-ai/tvm/pull/91
1.5 主线:嵌套 T.macro 复用转换模板,前端重复编译开销收敛(#3760)
合于 23:38。问题在外层函数每次运行时,嵌套的 T.macro 定义都会重新做源码解析、AST 重写与 Python 编译,即使 kernel 本身已经命中缓存。改法是引入一个有界的进程内缓存,只存与上下文无关的转换结果(未绑定的代码与重写后的源码),每个宏实例再绑定当前的 globals 与闭包值;依赖上下文的转换走原路径,被包装过的函数绕开缓存。宏调用点与 kernel 缓存都不变。回归覆盖新闭包、globals、捕获的缓冲、注解、上下文相关回退与闭包生命周期;CPU 侧前端检查通过,GPU 执行本地未验证——作者把这条边界写清楚了。
来源:https://github.com/tile-ai/tilelang/pull/3760
1.6 新 issue:条件标量化后缺 __syncthreads,0.1.12 至 0.1.15 静默生成错误 CUDA(#3808)
21:14 开出,是本窗口最重的缺陷报告。复现是一个共享内存 epilogue 的 kernel:0.1.12、0.1.13、0.1.14、0.1.15 四个版本都会在条件相关的标量化之后丢掉一道 __syncthreads(),编译器不报任何错,生成的 CUDA 直接算错。报告者给了完整可运行的 Python 复现脚本与环境(Python 3.14.9、CUDA 2.15.0+cu134)。这类问题的破坏力在于它跨了四个发布版本且无任何告警——凡是依赖共享内存在阶段之间做交接的 kernel 都可能踩到,而用户侧看到的只是「结果不对」。
来源:https://github.com/tile-ai/tilelang/issues/3808
1.7 新 issue:无 warp specialization 时 T.Pipelined 在 tcgen05 等待前重载 stage(#3800)
08:44 开出。触发条件是关掉 warp specialization(tl.disable_warp_specialized=True)后跑 T.Pipelined 加 T.tcgen05_gemm 的循环:stage 在 mbarrier_wait_parity 之前就被下一轮拷贝覆盖,MMA 读到的是被改写的 operand。报告者把 lowering 出来的 kernel 源码贴了出来,指认 As 的写入与 tcgen05mma_ss 的读之间缺少顺序约束。与 #3808 同属「生成的代码看起来没问题、实际不成立」这一类,区别在于这条发生在 Blackwell 的 tensor memory 路径上。
来源:https://github.com/tile-ai/tilelang/issues/3800
1.8 在审:迁到原生 Apache TIRx 的完整迁移草案(#3803);共享转置改走 ldmatrix/stmatrix(#3809)
#3803 是本期最有分量的在审项:把 vendored 的 TVM 编译器换成独立的 Apache TIRx,前端、IR 变换、代码生成、JIT 适配与打包全部迁到原生对象模型,并 pin 了 apache/tvm-tirx 的 7e49a148 与 apache/tvm-ffi 的 3e906bc6 两个上游提交,依赖树保持干净上游、不注入补丁。作者把它标成 draft,本身就是一次「底座换代」的公开准备。
#3809(01:27 提交)处理另一件具体的事:T.transpose 在两个共享缓冲之间一直走通用 SIMT 循环,8 位与 16 位瓦片用不上共享内存带宽,examples/deepseek_v32 里的 FP8 稀疏 MLA 只能手写 16x16 寄存器暂存来绕开;同时 T.Pipelined 给源区加了一维 stage 轴后,秩不匹配会让 MakeIndices 直接失败。改法是在支持 stmatrix 的目标(SM90 及以上、排除 CuTeDSL)先试 warp-matrix lowering:16 位瓦片每个 8x8 块用一次 ldmatrix 加一次 stmatrix.trans,8 位瓦片按 16x16 字节块用两次 ldmatrix.trans 加 __byte_perm 选择拼出四行,其余区域回退到原 SIMT 路径。
来源:https://github.com/tile-ai/tilelang/pull/3803、https://github.com/tile-ai/tilelang/pull/3809
二、多后端适配
2.1 昇腾:bitwise_rshift 的 64 位 dtype 改为显式报错,不再段错误(#1954)
合于 16:36(ascendc_pto 分支)。A2/A3 上通过公开 API T.tile.bitwise_rshift 传 int64/uint64 操作数时,编译器会在 OptimizeForTarget 阶段直接段错误。修法是在构造 intrinsic 之前检查 dst 与 src0 的 dtype,只允许 int16/uint16/int32/uint32,其余抛带实际 dtype 与支持列表的 ValueError;位移量 scalarValue 不限表达式 dtype;Buffer 与 BufferRegion 统一走 _underlying_buffer()。新增直接 API 诊断测试与子进程崩溃回归测试。作者明确写了边界:本次不新增 64 位右移支持,也不改直接构造底层 intrinsic 的路径。关联 #1720。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1954
2.2 昇腾:A2/A3 性能技能对齐 Ascend950 工作流(#1932)
15:36 合入,纯技能与流程重构:把 A2/A3 的性能调优技能与 A5(Ascend950)工作流结构对齐并统一,新增芯片类型识别与芯片信息检索两项技能。单看没有功能价值,但它意味着性能调优的「操作手册」开始按芯片代际分叉——同一套 skill 要能识别自己跑在哪一代上,否则调优建议会跨代误用。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1932
2.3 昇腾:threads=2 下 tmp workspace 被无条件减半,sin/cos 在间接寻址 tile 上 97.1% 元素算错(#1991 / #1992)
17:03 与 17:25 连开两条同题 issue,标题与正文一致(内容上可视作同一条被提交了两次)。报告环境是 910_9362 / CANN 9.0.0、threads=2(C:V=1:2)。根因在 AscendVidReduction:它把 kernel 内所有 UB 分配无条件减半,包括 InjectTmpBuffer 注入的 tmpspace,而 skip-set buffer(GM 索引含 UB load 的间接寻址,如 T.copy(KV[b_i, idx_ub[bi_i], g_i, :D], kv_ub))的形状保持整块、本来不参与 vid 减半。后果分两档:sin/cos 直接算错(upstream HEAD 上 3978/4096 个元素错,97.1%);sigmoid/round 违反 AscendC 的 calCount <= tmp 契约,数值碰巧正确但越界写落在布局后方的 scratch 上,布局一变就是静默数据污染。报告者说明这是 #1897 review 时发现的,而 #1897 的修法只覆盖了 count 侧语义(HasModifiedBuffer 跳过 workspace 槽位),workspace 侧仍未处理。
来源:https://github.com/tile-ai/tilelang-ascend/issues/1991、https://github.com/tile-ai/tilelang-ascend/issues/1992
2.4 海光:分支三笔 LDS 相位修,并首次开出两个 PR(#33 / #34)
主分支仍停在 11-02,当天三笔都落在 feat/hcu-scaled-cast 分支,作者是海光侧(提交域 hygon.cn):
- 10:06
fix(hcu): adapt default LDS swizzle to 64-bank phases - 14:04
fix(hcu): restore LDS swizzle planning for MN slice consumers - 19:49
fix(hcu): support K-partitioned layout-aware LDS reads
更重要的是评审状态的变化:09:08 开出 PR #33(scaled cast,即分支主体),19:52 开出 PR #34(reduction 的 XOR shuffle 用 DPP 与 lane swap 优化,外加「布局推断先传播已知约束再分配自由候选」的修法)。上一期报告写的是「四笔都在功能分支上、尚未开 PR」,本窗口这两笔 PR 把海光从自研分支推进到公共评审,是国产后端里第一个走完这条路径的。
来源:https://github.com/tile-ai/tilelang-hygon/pull/33、https://github.com/tile-ai/tilelang-hygon/pull/34、https://github.com/tile-ai/tilelang-hygon/commit/07451160e5aa758419201adb598725cbef70e3a4
2.5 摩尔线程:README 文档更新,结束自 09-30 以来的静默(11-05)
19:46 主分支落了一笔 docs: update README installation and operator development references,是本仓自 09-30 以来第一次提交。同期 6 个遗留 PR(#1 至 #7 中的依赖升级、TileOPs 子模块、MUSA 集成的证据门槛文档)时间戳被一并刷新,但内容仍是 04 月到 08 月的老提案,本窗口没有代码变更。判断上宜保守:文档提交说明仓库有人在维护,不足以证明后端开发恢复;真正的信号要看后续是否出现算子或 lowering 提交。
来源:https://github.com/tile-ai/tilelang-musa/commit/4cac45b859108a86ac9a0b1cbfdeb33da7711697
2.6 沐曦与 mlir-ascend:沐曦连续第三个窗口无推送;mlir-ascend 三项 TLIR 算子在审(11-05)
tile-ai/tilelang-metax 最后推送仍是 10-26,连续第三个窗口静默。tile-ai/tilelang-mlir-ascend 主分支无新提交,但三项基于 TLIR 方言的算子 PR 在推进:#203 给 TLIR 加 tl.sub/mul/div/max/min 并补 TL 到 HIVM 的 lowering,01:42 更新;#211 加 tl.shl、#212 加 tl.shr,两笔在 19:44 与 20:58 先后提交。这条线在做的事是把 TileLang 的语义一层层落进 MLIR 方言,节奏慢但方向单一。
来源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-mlir-ascend/pull/203、https://github.com/tile-ai/tilelang-mlir-ascend/pull/211、https://github.com/tile-ai/tilelang-mlir-ascend/pull/212
三、生态与采用方
3.1 TileOPs:九笔合入,广播操作数就地读、GEMM 单波次填充、reduction 布局缓存(11-05)
九笔集中在 14:01 与 21:06 至 22:46 两个时段,逐笔看主干:
- #2832 广播操作数就地读:
LerpTensorFwdOp、WhereFwdOp、ClampTensorFwdOp、MaskedFillScalarFwdOp、MaskedFillTensorFwdOp五类算子不再先把 stride 0 的广播张量展开成输出形状,而是在一个 kernel 里按 stride 0 直接读。一个共享 builder 服务全部五种,操作数形状合并成一个循环嵌套,16 字节以内的向量留在一行内层,每个算子只写自己的逐元素体。广播行耗时是主线的 0.13 到 0.44 倍,与 torch-compile 打平或更好;其余逐元素行与主线逐字节一致。 - #2829 一元逐元素单波次:体内最多两条各自降为一条指令的操作、读写总量不超过 L2 的四分之三、结果与操作数同宽时,缩成 256 线程的一波。顺带修了 ReLU 对 NaN 返回 0 的问题(改为返回 NaN,16 位走 torch 位模式、float32 用规范 NaN)。H200 上 327 项 GPU 回归全过,13 项新用例过了 Compute Sanitizer memcheck。
- #2834 小 M NT GEMM 的原生宽度瓦片:
ds-v3-decode-down(128x7168x2048,NT)在主线上只用 132 个 SM 里的 112 个,112 宽的瓦片在 0.1.12 上被拆成七个 n16 WGMMA。改成每 16 个 K 发一次瓦片自身宽度的 WGMMA 之后,decode-down 领先最快对比项 3%,gate-up-b32 比主线快 1.4%(仍落后 cuBLAS 0.5 到 1%)。新增的fill瓦片只对 96 < M <= 128 且 N 能整除 8 的形状启用,结构因子 0.58 单独拟合。 - #2835 reduction 布局缓存:归约 kernel 每次调用都重建轴置换与输出形状,而
ReduceCall是固定的。改为每个 kernel 实例只算一次、与既有张量适配器共享纯形状 helper,不引入全局缓存。六组 FP32/BF16VarMeanFwdOp配置的完整热调用耗时下降 4.33% 到 6.17%,生成的 CUDA 与主机源码完全一致——省的是调用准备成本,不是 GPU 时间。 - #2833 CI 修:GitHub 废弃 runner 2.336.0 之后
gpu-smoke全部排队,entrypoint.sh又带--disableupdate无法自升级,RUNNER_VERSION从 2.336.0 提到 2.338.2。
来源:https://github.com/tile-ai/TileOPs/pull/2832、https://github.com/tile-ai/TileOPs/pull/2829、https://github.com/tile-ai/TileOPs/pull/2834、https://github.com/tile-ai/TileOPs/pull/2835、https://github.com/tile-ai/TileOPs/pull/2833
3.2 TileOPs:roofline 计价规则重写,三条不可能的带宽行回到合理区间(#2830)
夜间任务 41873301494 抓出三行物理上不可能的 roofline:GQAVarlenFwdOp 的两个 gpt-oss-serve-window 配置读到 7.48 与 7.50 TB/s,GroupedGEMMFwdOp 的 qwen3-30b-ep8-empty 读到 5.30 TB/s。这一笔把计价规则按「这次调用实际读了多少」重写:GQA 的稠密、varlen、分页前向只对「有 query 能看到的 key 行」计 K/V、对「能看到 key 的 query 行」计 q、对 key 实际到达的截断位置计 RoPE 表、对看到 key 的请求计 FP8 缩放;GroupedGEMM 通过新的 grouped_gemm_roofline 按 batch_sizes 给某行的组计价;NSA 压缩、NSA 稀疏、DSA decode、MHA 分页与 FP8 lightning indexer 都只在真正读的位置计费,DSA 还认 is_causal=False。修正后三行变成 0.62、0.62、2.65 TB/s。附加收益是修掉了空 K/V 存储与空分页池的除零。测试节点加 16。
来源:https://github.com/tile-ai/TileOPs/pull/2830
3.3 TileOPs:两笔合入相隔 2.5 分钟导致 main 的 GPU Smoke 失败,当日补修(#2837)
22:46 合入的这笔本身就是一条纪律记录。#2824 在 21:06 合入时把 grouped GEMM 的 roofline 行改成必须声明 sizes 并从它推导 G;#2830 在 21:08 合入,写的还是只有 G 没有 sizes。两笔相隔约 2.5 分钟,各自 CI 都是绿的,合并后 main 的 GPU Smoke 立刻报 RowError: row misses ['sizes']。修法是把 _grouped 的 G: 2 换成 sizes: [1, 0],断言的字节数(24、32)不变。这类事故的成因不是谁写错了,而是并行合并缺少 post-merge 门禁:单笔 CI 通过无法覆盖「两笔语义相交」的情况,只能靠合入后立刻跑全量。
来源:https://github.com/tile-ai/TileOPs/pull/2837、https://github.com/tile-ai/TileOPs/pull/2824、https://github.com/tile-ai/TileOPs/pull/2830
3.4 TileOPs 在审:三阶段瓦片因子、TN 单元素收缩、FP8 分组 GEMM 的 TMA 缩放加载(#2836 / #2838 / #2839)
#2836 给 coop2_swap 的三阶段瓦片单独拟合因子(1.36),因为按 FLOP 算它并不比四阶段快,但按字节与 K 步数被高估了约 10%,结果把 160 宽的瓦片挤掉、夜间成绩从 1.005 到 1.011 倍落到 0.985 到 0.995 倍 QuACK。因子只在那 42 个 off-manifest 形状上拟合(平台区 1.344 到 1.392),44 形状扫描上选择器的 regret 从 0.89% 降到 0.06% 几何均值、最大从 5.04% 降到 0.99%。#2838 修的是分派器把 K 当成连续维、于是否决 TN 布局 K=1 的合法调用(A=[1,64]、B=[1,64] 求 AᵀB 本来可以用 cp.async 内核算),改为回退检查两侧操作数的物理连续维是否各有至少 4 字节行宽,新增 24 个测试节点。#2839(00:56 提交)让 FP8 分组 GEMM 用 TMA 加载 A 的缩放(每盒四个 K 步、stage_scales 2、在三行上标定),并给 K <= 1536 的非融合紧分组 BF16/FP16 GEMM 换 64x256 的 ping-pong 瓦片;FP8 qwen3-30b-ep8-skew-up 从 0.973 升到 1.015 倍 DeepGEMM,BF16 skew-down 从 0.962 升到 1.099 倍 QuACK。#2840(06:51 提交)补 SM89 的 prefill sinks,并把 sink 的有无纳入编译实例身份。
来源:https://github.com/tile-ai/TileOPs/pull/2836、https://github.com/tile-ai/TileOPs/pull/2838、https://github.com/tile-ai/TileOPs/pull/2839、https://github.com/tile-ai/TileOPs/pull/2840
3.5 TileFoundry:WGMMA 同步计划合入(#286),部分归约的完成语义在审(#288)
13:40 合入的 #286 是上一期在审项,处理 WGMMA 的异步性:指令发出后仍会继续读共享内存与寄存器,后续访问必须排在它的效果之后。做法是给每个操作数效果声明完成条件,把 fence、commit、wait 挂到需要它们的 HIR 节点上,lowering 在指令发射点消费这些步骤;若后续作用域放不下集体等待,就在发射作用域里做完。同时修正寄存器 Reshard 的别名,让每个元素在持有副本的每个线程上留在同一个寄存器,并用这个存储身份追踪寄存器依赖;SM90 上以 sm_90a 单独编译发出 WGMMA 组步骤的 CUDA 单元。作者保留了两条风险声明:矩阵指令发射器仍未接上,所以这笔不证明执行正确性、也不给性能;TMA 的完成语义另行处理。20:12 另开 #288,处理「已完成的部分归约」在分析中的计价。
来源:https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/pull/288
3.6 TileSight:H20 的 WGMMA 能力声明修正,为仓库首个 PR(#1)
14:37 提交、17:26 更新,是 tile-ai/TileSight(24 星)的第一个 PR。该仓库是 07-24 arXiv 论文《TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters》的开源实现:从瓦片形状、数据移动、软件流水与硬件档案出发,在 CPU 上估计算子代价、缓存流量与执行时间,CUDA/HIP 探针单独测目标 GPU。这笔改动很小但方向明确:H20 是 Hopper/GH100,带 requires_arch=("wgmma",) 的计划原先因为架构档案没有声明 WGMMA 支持而抛 UnsupportedError,现在把 support_wgmma=True、并显式保持 support_utcmma=False,于是这类计划被接受、而 tcgen05 与 Tensor Memory 需求仍被拒。作者同时指出容量要按 SKU 分开处理(NVIDIA 同时列 96 GB 与 141 GB 的 H20),本次只动能力元数据,容量、带宽、频率、吞吐与利用率数值全部不变。验证是 CPU 上 847 项全过,明确声明「不是 GPU 性能验证」。
来源:https://github.com/tile-ai/TileSight/pull/1、https://arxiv.org/abs/2607.22432
3.7 采用方与第三方:TileKernels、FlashQLA 无推送;两个新的第三方仓出现(11-05)
deepseek-ai/TileKernels 仍停在 10-14,QwenLM/FlashQLA 停在 10-30,本窗口都没有提交,FlashQLA 最新发布仍是 09-30 的 v0.1.3。第三方侧出现两个新的小仓:Tianchen-ai/tilelang-weft-bench(「TileLang-Weft full CPU operator benchmarks using the shared Weft Bench engine」,0 星,建立 CPU 侧算子基准的统一引擎)、flytiger-eco/tilelang-wrapper-for-sail(0 星,暂无描述)。两者都很早期,先记入观察名单,不宜当技术进展写。
来源:https://github.com/deepseek-ai/TileKernels、https://github.com/QwenLM/FlashQLA、https://github.com/Tianchen-ai/tilelang-weft-bench、https://github.com/flytiger-eco/tilelang-wrapper-for-sail
四、社区、教程与活动
4.1 文档站三次更新;夜间仓无主分支提交;昇腾每日测试 3030 项全过(11-05)
tile-ai/tilelang.github.io 在 16:07、19:33、23:43 各推一次 Update docs,对应主仓当天的接口变化(host-mapped 注解、TF32 rounding、Z3 配置)。TileOPs-nightly 只有分支推送、主分支无新提交;TileOPs.github.io 在 09:02 有一次推送。昇腾仓的每日测试 issue #1988 记录 11-05 08:16 一轮 3030 项全部通过、通过率 100%,与同日新开的 #1991/#1992 形成对照(见 2.3)。
来源:https://github.com/tile-ai/tilelang.github.io、https://github.com/tile-ai/tilelang-ascend/issues/1988
4.2 新闻侧:四组查询、Hacker News 与 arXiv 均无窗口内命中(11-05)
Google News 四组查询(英文 "TileLang" OR "tile-lang"、TileLang GPU kernel,中文 TileLang 算子、TileLang 编译器)里,最新一条命中是 376 小时前的 SemiAnalysis《The Huawei-DeepSeek Symbiosis》,其余集中在 09-30 那一波「DeepSeek 开源昇腾版 TileLang、DeepGEMM、DeepEP」的报道(约 880 小时前),全部早于本窗口。Hacker News 以 tilelang 检索 24 小时内 0 条,抓到的三条 Show HN 是终端数据查看器、工作流引擎与 macOS 文件管理器,与主题无关。arXiv 的 all:tilelang 最新命中是 11-03 的《Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving》(窗口外,属引用 TileLang 的工作),再往前是 07-24 的 TileSight 本体。报告主体因此仍全部来自代码侧。
来源:https://news.google.com/、https://news.ycombinator.com/、https://arxiv.org/
五、趋势观察
第一,算子库的竞争点已经从「算子数量」彻底转到「选择器与准备成本」。今天 TileOPs 九笔里,真正改 GPU kernel 的只有 #2832(广播就地读)与 #2834(原生宽度瓦片),#2835 省的是调用准备(4.33% 到 6.17%,生成源码不变),#2836 改的是一个拟合因子,regret 从 0.89% 降到 0.06%。这说明性能差距的剩余空间主要不在「怎么算」,而在「选哪个瓦片、少做哪些准备工作」。这条路走到头会得到一个用标定契约描述的代价模型,而不是一堆手调参数——今天 #2836 的「单独拟合、只动一个常数」正是这个形态的雏形。
第二,编译器侧当前的主要缺陷类型是「静默生成错代码」,而不是编译失败或性能不足。#3808 跨 0.1.12 到 0.1.15 四个版本少一道 __syncthreads(),#3800 在 Blackwell 路径上让 MMA 读到被覆盖的 operand,#1991/#1992 在昇腾 threads=2 下让 sin/cos 算错 97.1% 的元素——三条都不报错、都能通过常规测试、都要靠人写出针对性的复现才能抓到。这类缺陷对用户的伤害最大,因为它把「编译器对不对」这个前提变成了需要验证的假设。与之对应的工程动作应当是扩大形状组合的差分测试(对照组用 reference 实现逐元素比),而不是继续加单元测试数量。
第三,昇腾侧的节奏从「架构改动后清回归」转向「清正确性边界」。#3512 那类结构性改动带来的四类回归在上一个窗口刚清完,本窗口的两笔合入(段错误改报错、性能技能分代)和两条新 issue(tmp workspace 减半、calCount <= tmp 违约)都属于边界收敛。值得注意的是 #1991 报告者是 #1897 的 review 方,而 #1897 的修法只覆盖了 count 侧——同一处布局逻辑的另一个侧面在两周后才被指出,说明昇腾后端的布局不变量还没有形成可检查的书面契约。
第四,国产后端的「公共化」在提速。海光从自研分支走到两个正式 PR,摩尔线程恢复文档维护,mlir-ascend 用三个小 PR 把算术算子一个个搬进 TLIR 方言。三条线的共同点是都不再等「一次性合入一个大后端」,而是把可评审的最小单元推到公共仓库。这对上游是好事:评审者能早就布局与 lowering 的设计提意见,而不是在三万行一次性合入时才看。
最后一条风险面:主仓 #3803 的 TIRx 迁移如果推进,3rdparty/tvm 的 vendored 结构会被替换成两个 pin 的上游提交,届时所有依赖 TVM 内部 API 的后端(昇腾、海光、mlir-ascend 的胶水层)都要跟着改。草案里明确写了「不注入依赖补丁」,这既是承诺也是约束——今天 mlir-ascend 把 BishengIR 适配器搬进 contrib、昇腾把性能技能按芯片代际分叉,都可以看作是为这类底座切换预留的可替换面。
附、素材与核查说明
信源核查表
| 信源 | 核查结果 |
|---|---|
| tile-ai 组织(29 仓库) | 窗口内推送 10 个仓库 |
| 主仓 tilelang | 主分支 5 笔合入;新开 issue 2 条(#3808、#3800);在审更新以标量化、CuTeDSL、Profiler、Ascend 四簇为主 |
| TileOPs | 主分支 9 笔合入;在审 4 笔新增或更新;无新 issue |
| TileFoundry | 主分支 1 笔合入(#286);在审 1 笔(#288) |
| TileSight | 首个 PR(#1),无主分支提交 |
| tilelang-ascend | 默认分支 ascendc_pto 2 笔合入(#1932、#1954);新开 issue 3 条(#1991、#1992 同题,#1988 为每日测试报告) |
| tilelang-mlir-ascend | 主分支无新提交;#203、#211、#212 三笔 TLIR 算子 PR 推进 |
| tilelang-hygon | 主分支无新提交(停在 11-02);feat/hcu-scaled-cast 分支 3 笔;新开 PR #33、#34 |
| tilelang-musa | 主分支 1 笔文档提交(自 09-30 以来首次);6 个遗留 PR 时间戳刷新,内容未变 |
| tilelang-metax | 窗口内无推送(连续第三个窗口) |
| TileOPs-nightly / TileOPs.github.io / tilelang.github.io | 夜间仓主分支无提交;TileOPs 文档站 1 次推送;tilelang 文档站 3 次 Update docs |
| 发布页(tilelang / TileOPs / TileFoundry) | 窗口内无新版本;tilelang 最新仍为 09-30 的 v0.1.15,TileOPs 为 10-27 的 v0.1.0,TileFoundry 为 10-25 的 v0.0.4rc1 |
| TileKernels / FlashQLA | 窗口内无推送 |
| TileLang×RISC-V 四源巡检 | 无变化(四个上游的引用与 PR 状态与上期快照一致) |
| Google News(中英文四组)/ Hacker News / arXiv | 无窗口内实质命中,最新命中为 376 小时前 |
完整信源清单
- [1] tilelang TF32 round-to-nearest-even 转换(#3714)— https://github.com/tile-ai/tilelang/pull/3714
- [2] tilelang host-mapped 缓冲注解(#3802)— https://github.com/tile-ai/tilelang/pull/3802
- [3] tilelang 标量条件 cast 保留向量宽度(#3759)— https://github.com/tile-ai/tilelang/pull/3759
- [4] tilelang Z3 rlimit 可配置(#3752)— https://github.com/tile-ai/tilelang/pull/3752
- [5] tile-ai/tvm Z3 资源预算配置(#91)— https://github.com/tile-ai/tvm/pull/91
- [6] tilelang 宏转换模板复用(#3760)— https://github.com/tile-ai/tilelang/pull/3760
- [7] tilelang 条件标量化后缺同步屏障(#3808)— https://github.com/tile-ai/tilelang/issues/3808
- [8] tilelang 无 warp specialization 时 stage 提前重载(#3800)— https://github.com/tile-ai/tilelang/issues/3800
- [9] tilelang 迁到原生 Apache TIRx(#3803)— https://github.com/tile-ai/tilelang/pull/3803
- [10] tilelang 共享转置走 ldmatrix/stmatrix(#3809)— https://github.com/tile-ai/tilelang/pull/3809
- [11] tilelang 多缓存目录(#3807)— https://github.com/tile-ai/tilelang/pull/3807
- [12] tilelang-ascend bitwise_rshift dtype 校验(#1954)— https://github.com/tile-ai/tilelang-ascend/pull/1954
- [13] tilelang-ascend A2/A3 性能技能对齐(#1932)— https://github.com/tile-ai/tilelang-ascend/pull/1932
- [14] tilelang-ascend tmp workspace 被减半(#1991)— https://github.com/tile-ai/tilelang-ascend/issues/1991
- [15] tilelang-ascend tmp workspace 被减半(#1992)— https://github.com/tile-ai/tilelang-ascend/issues/1992
- [16] tilelang-ascend 每日测试 3030 项全过(#1988)— https://github.com/tile-ai/tilelang-ascend/issues/1988
- [17] tilelang-hygon scaled cast PR(#33)— https://github.com/tile-ai/tilelang-hygon/pull/33
- [18] tilelang-hygon reduction lane swap PR(#34)— https://github.com/tile-ai/tilelang-hygon/pull/34
- [19] tilelang-hygon LDS 相位提交 — https://github.com/tile-ai/tilelang-hygon/commit/07451160e5aa758419201adb598725cbef70e3a4
- [20] tilelang-musa README 文档提交 — https://github.com/tile-ai/tilelang-musa/commit/4cac45b859108a86ac9a0b1cbfdeb33da7711697
- [21] tilelang-mlir-ascend TLIR 算术算子(#203)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/203
- [22] tilelang-mlir-ascend tl.shl(#211)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/211
- [23] tilelang-mlir-ascend tl.shr(#212)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/212
- [24] TileOPs 一元逐元素单波次(#2829)— https://github.com/tile-ai/TileOPs/pull/2829
- [25] TileOPs 广播操作数就地读(#2832)— https://github.com/tile-ai/TileOPs/pull/2832
- [26] TileOPs 小 M NT GEMM 原生宽度瓦片(#2834)— https://github.com/tile-ai/TileOPs/pull/2834
- [27] TileOPs reduction 布局缓存(#2835)— https://github.com/tile-ai/TileOPs/pull/2835
- [28] TileOPs CI runner 升级(#2833)— https://github.com/tile-ai/TileOPs/pull/2833
- [29] TileOPs roofline 计价重写(#2830)— https://github.com/tile-ai/TileOPs/pull/2830
- [30] TileOPs grouped GEMM 行声明 sizes 修复(#2837)— https://github.com/tile-ai/TileOPs/pull/2837
- [31] TileOPs FP8 分组 GEMM 连续布局(#2824)— https://github.com/tile-ai/TileOPs/pull/2824
- [32] TileOPs 三阶段瓦片因子(#2836)— https://github.com/tile-ai/TileOPs/pull/2836
- [33] TileOPs TN 单元素收缩(#2838)— https://github.com/tile-ai/TileOPs/pull/2838
- [34] TileOPs FP8 分组 GEMM 的 TMA 缩放加载(#2839)— https://github.com/tile-ai/TileOPs/pull/2839
- [35] TileOPs SM89 prefill sinks(#2840)— https://github.com/tile-ai/TileOPs/pull/2840
- [36] TileFoundry WGMMA 同步计划(#286)— https://github.com/tile-ai/TileFoundry/pull/286
- [37] TileFoundry 部分归约的完成语义(#288)— https://github.com/tile-ai/TileFoundry/pull/288
- [38] TileSight H20 的 WGMMA 能力声明(#1)— https://github.com/tile-ai/TileSight/pull/1
- [39] TileSight 论文(arXiv 2607.22432)— https://arxiv.org/abs/2607.22432
- [40] 第三方仓 tilelang-weft-bench — https://github.com/Tianchen-ai/tilelang-weft-bench
- [41] 第三方仓 tilelang-wrapper-for-sail — https://github.com/flytiger-eco/tilelang-wrapper-for-sail
- [42] TileKernels(采用方,10-14 后无推送)— https://github.com/deepseek-ai/TileKernels
- [43] FlashQLA(采用方,10-30 后无推送)— https://github.com/QwenLM/FlashQLA
- [44] tilelang 发布页(最新 v0.1.15,09-30)— https://github.com/tile-ai/tilelang/releases
- [45] TileOPs 发布页(最新 v0.1.0,10-27)— https://github.com/tile-ai/TileOPs/releases
核查说明
本期所有合入、在审、issue 与提交数量均来自 GitHub 检索与原子源的实测输出:合入与在审走 search API(合入 18 笔、在审更新 57 笔、新开 issue 5 条),提交标题走各分支的 commits atom(12 个分支、21 笔落在窗口内),发布版本走 releases atom。组织级匿名 API 在采集后段触及小时限额,主分支之外的组织清单以网页与 atom 源交叉核对,未使用估算值。#1991 与 #1992 两条 issue 标题与正文一致,按「同题两条」记录,未合并为一条。海光三笔与摩尔线程一笔属未合入提交,按分支级动态表述,不代表主线已合入。TileSight 与 TileFoundry 的 WGMMA 相关改动均由作者自述限定为能力元数据或同步计划,未声称执行正确性或性能,报告沿用其限定。Google News、Hacker News 与 arXiv 三条新闻通道在本窗口均无命中。