调研窗口:2026-10-23 07:00 ~ 2026-10-26 07:00(北京时间,72 小时,覆盖周五至周一)。

信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileFoundry / tilelang-ascend〔ascendc_pto〕/ tilelang-hygon / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 三个文档站)+ 合入与在审真值检索(is:merged 与 is:open 两组)+ 新增 issue 全量核对 + 提交级 diff 与发布页 tag 核查 + PyPI 发布状态核查 + 采用方仓核查(TileKernels / FlashQLA)+ 社区新仓检索 + Google News 中英文多组查询(走代理)+ Hacker News + arXiv + 产业媒体报道。


本期索引

  • 今日重点:72 小时合入 78 笔、在审更新 76 笔;TileOPs 首次 PyPI 发布只剩最后一项,CHANGELOG 已写定 0.1.0(10-26);主仓在审的 TIRx 替换动 1043 个文件;昇腾仓 situ_glu 合入、每日测试单日 2903 项全通
  • 一、核心项目进展
    • 1.1 主仓:partial-sync 命名屏障 id 在两类存储域间重号(#3643)(10-25)
    • 1.2 主仓:clamp 的 NaN 传播下降被回退,T.clamp 回到 min/max(#3628)(10-23)
    • 1.3 主仓:int64 原子 max/min 与 AtomicStore 的 int64、bf16、fp16 崩溃(#3307)(10-23)
    • 1.4 主仓:VectorizeLoop 不再静默忽略循环步长(#3345)(10-23)
    • 1.5 主仓:NVRTC 命中磁盘缓存后自动调优结果写不回(#3193)(10-23)
    • 1.6 在审:用独立打包的 TIRx 替换内置 TVM,改动 1043 个文件(#3666)(10-25)
    • 1.7 在审:sm_100 的 TMEM 释放缺 CTA 屏障,warp 特化前导语句漂移(#3645 / #3663 / #3664)(10-25)
    • 1.8 新增缺陷:sm_120 上归约下降失败、三张注意力示例超共享内存(#3647 / #3646)(10-25)
  • 二、多后端适配(昇腾 / 海光 / 新增 CPU 与 ROCm)
    • 2.1 昇腾主仓:DeepGEMM 示例支持 FP4×FP4 打包输入(#3589)(10-23)
    • 2.2 昇腾主仓:DMA 保留缓冲基偏移、区域界先提升、周期缓冲完整生命周期(#3627 / #3625 / #3606)(10-23)
    • 2.3 昇腾主仓:CI 镜像升到 cann9.3.0 周更版(#3630)(10-23)
    • 2.4 昇腾在审:E4M3 的 UB→L1 拷贝、片上内存溢出提前拒绝(#3639 / #3655)(10-24~25)
    • 2.5 昇腾在审:PTO 目标默认走 tvm_ffi 执行后端(#3516)(10-24)
    • 2.6 昇腾仓:situ_glu 算子合入,对齐 Kimi K3 门控激活(#1891)(10-23)
    • 2.7 昇腾仓在审:稀疏 MLA、GLM-5.3 kpool 与三张规范化示例(10-23~25)
    • 2.8 昇腾仓:每日测试 2903 项全通,次日测试作业报失败(#1941 / #1944)(10-25~26)
    • 2.9 海光:主干三笔合入,寄存器流水线重新进在审(#26 / #18 / #27)(10-23~25)
    • 2.10 新增后端在审:Intel AMX 的 BF16/INT8 GEMM、ROCm 的 FP4 与 E8M0(#3637 / #3665)(10-24~25)
  • 三、生态与采用方(TileOPs / TileFoundry / 发行版)
    • 3.1 TileOPs:首发清理八项完成七项,CHANGELOG 写定 0.1.0(#2738)(10-25)
    • 3.2 TileOPs:公共名字与路径一次定稿,打包改元数据并支持 3.13(#2741 / #2742 / #2746)(10-25)
    • 3.3 TileOPs:写算子的入口收口为 authoring,eager 调用收敛到单执行器(#2743 / #2744)(10-25)
    • 3.4 TileOPs:算子的共享语义一次写清,输入校验由清单生成(#2751 / #2758 / #2760)(10-25)
    • 3.5 TileOPs:发布文档、第三方声明与 pip 安装入口(#2753)(10-25)
    • 3.6 TileOPs 性能:GDN 预填充分区宽度按调用自身长度选,KDA 原生门控进 kernel(#2745 / #2683)(10-23~25)
    • 3.7 TileOPs 性能:合并注意力态、分页 GQA 与 indexer 尾波(#2739 / #2752 / #2735 / #2730)(10-24~25)
    • 3.8 TileOPs 性能:GEMM 尾波窄瓦片、乒乓 BMM 分级暂存(#2733 / #2734 / #2720)(10-23~25)
    • 3.9 TileOPs 性能:MoE 的 A 箱装载与原生宽度 WGMMA,逐元素倒数与行广播(#2711 / #2715 / #2710 / #2725)(10-23~24)
    • 3.10 TileOPs 性能:归约与归一化改 16 字节 evict-last 读(#2757 / #2759 / #2748 / #2737)(10-25)
    • 3.11 TileOPs 正确性与 CI:SSD 扫描尾部、GLA 段展开、H200 发布校验(#2716 / #2750 / #2749 / #2718 / #2713)(10-23~25)
    • 3.12 TileFoundry:sm_89 的 E4M3 warp MMA 与逐指令能力校验(#261 / #262 / #264)(10-23~24)
    • 3.13 TileFoundry:cp.async 目标布局寻址与 MLA 反向布局(#263 / #260 / #257)(10-23)
    • 3.14 TileFoundry:拆成 core / nvidia / apple 三个发行包,测试输入只装载一次(#268 / #269)(10-25)
    • 3.15 分发版:Sunrise 0.1.15+sunrise.1.2.0 发布 PR 合入(#9)(10-23)
  • 四、社区、教程与活动
    • 4.1 社区:tilelang-puzzles-ascend 做 PTO 对 Ascend 的量化 kernel 教程(10-23 推送)
    • 4.2 新仓:中文 TileLang 算子面试练习站,9 模块 18 题(10-23)
    • 4.3 社区技能与教程仓继续推送,文档站四笔跟随首发改版(10-23~25)
    • 4.4 讨论:TileScale 用户询问所用 TileLang 版本能否升级(#72)(10-25)
  • 五、趋势观察

今日重点

72 小时里 tile-ai 组织合入 78 笔(TileOPs 51、tilelang 10、TileFoundry 10、文档站 4、tilelang-ascend 1、tilelang-hygon 1、tilelang-sunrise 1),在审更新 76 笔,新开 issue 19 条。最重的一条在发布队列:TileOPs 的「首发清理」跟踪单(#2738)八项完成七项,只剩 TestPyPI 试跑与首个 tag,而 CHANGELOG 已在 10-25 23:17 写下首条正式版本记录「0.1.0 - 2026-10-26」,声明算子按 tileops.<family>.<Op> 导入、tileops.backend 与 tileops.authoring 分别供后端方与第三方写算子,验证配置锁定为 SM90(H200)、Python 3.12/3.13、PyTorch 2.13、TileLang 0.1.12、CUDA 13.2;但 PyPI 上该包最新仍是 2025-10-09 的 0.0.1.dev2,仓库无任何 tag 与 release,首发处在「文档就绪、制品未上传」的最后一格。主仓 10 笔合入集中在静默错:跨存储域的命名屏障重号(#3643)、T.clamp 的 NaN 传播下降被回退(#3628)、int64 原子 max/min 与 AtomicStore 崩溃(#3307)、VectorizeLoop 丢掉循环步长(#3345)。在审最值得注意的是 #3666:删掉 3rdparty/tvm、把前端与下降整体迁到独立打包的 TIRx,1043 文件 +71428/-104176。昇腾线同日合入 Kimi K3 门控激活算子 situ_glu(#1891),该仓每日测试 10-25 报 2903 项全通、10-26 凌晨又报测试作业失败(#1944)。


一、核心项目进展

1.1 主仓:partial-sync 命名屏障 id 在两类存储域间重号(#3643)

合于 10-25 14:30,2 文件 +57/-1,修 #3642。ThreadSync 对 shared 与 shared.dyn 各跑一次,每次新建一个改写器、各自从 kFirstUsedBarrier 重新编号;第二次运行跳过第一次已改写的同步,却不知道那些 id 已占用,于是把同一个 bar.sync id, count 分给互不重叠的两组线程。修法是让编号跨两次运行连续,并为跨域复用加断言。这类缺陷编译期不报错、只在大 kernel 偶发。

来源:https://github.com/tile-ai/tilelang/pull/3643

1.2 主仓:clamp 的 NaN 传播下降被回退,T.clamp 回到 min/max(#3628)

合于 10-23 16:51。回退 #3205 整条链路:删除 tl.clamp 内建、各后端下降注册、codegen 分支与设备端辅助函数(含昇腾注册),T.clamp(x, lo, hi) 恢复为 T.min(T.max(x, lo), hi),理由是额外 NaN 处理挡住了 min/max 的化简与向量化。NaN 诉求未消失,同日另开 #3605 以 T.clamp_symmetric 在 CUDA SM86+ 单列实现,分工是「T.clamp 保通用、新接口保语义」。

来源:https://github.com/tile-ai/tilelang/pull/3628

1.3 主仓:int64 原子 max/min 与 AtomicStore 的 int64、bf16、fp16 崩溃(#3307)

合于 10-23 18:44。normalize_atomic_type 自 #1716 起把 int64_t 归一成 unsigned long long 以支持 atomic add,结果让 Atomic 的 max/min 及引用形式在 int64 上直接崩,AtomicStore 同样中招并额外波及 bf16 与 fp16。修复按操作为 int64 单独选型,AtomicStore 走独立类型分派。属「为一个操作改公共归一化、连带打断其余操作」的典型回归。

来源:https://github.com/tile-ai/tilelang/pull/3307

1.4 主仓:VectorizeLoop 不再静默忽略循环步长(#3345)

合于 10-23 19:02,属 #3301 跟踪的 F319。向量化器只收到 (loop_var, extent, body),步长根本没进去:T.vectorized(0, 8, step=2)(迭代集 0/2/4/6)被下降成覆盖整个 [0, 8) 的连续访问。修复在进入向量化前拒绝非单位步长并给诊断。这一批子项的共同特征是前端接受了无法下降的写法、后端给出另一个程序的机器码。

来源:https://github.com/tile-ai/tilelang/pull/3345

1.5 主仓:NVRTC 命中磁盘缓存后自动调优结果写不回(#3193)

合于 10-23 15:04。NVRTC 从磁盘加载 kernel 时不设 adapter.libpath,而调优选中配置后要凭该字段保存缓存 kernel,于是每次命中都抛 AttributeError、结论被丢弃,用户看到的是「每次启动都在重调」。修复在缓存加载路径恢复 libpath,与全新编译路径一致,补缓存重载与负向对照用例。

来源:https://github.com/tile-ai/tilelang/pull/3193

1.6 在审:用独立打包的 TIRx 替换内置 TVM,改动 1043 个文件(#3666)

10-25 19:16 提交,1043 文件 +71428/-104176。删掉 3rdparty/tvm 及 submodule 条目(保留 CUTLASS 与 HIP 头文件),把前端、原生张量与作用域 IR、下降 pass、后端流水线、执行适配整体迁到独立构建发布的 tvm_tirx,兼容层限定在 TileLang 内部,不设全局 tvm 别名、不保留第二份 TIRx。这把前几期「tvm fork 先行、主仓跟进」推到终点:底座从 vendored 源码转成外部依赖,此后可独立发版与回滚,合入影响面覆盖整仓。

来源:https://github.com/tile-ai/tilelang/pull/3666

1.7 在审:sm_100 的 TMEM 释放缺 CTA 屏障,warp 特化前导语句漂移(#3645 / #3663 / #3664)

三笔同源,都在自动 warp 特化与 Blackwell 新存储上。#3645:sm_100 上张量内存释放没有 CTA 级屏障,生产者 warp 可在消费者 tcgen05.mma/tcgen05.ld 仍读累加器时将其释放。#3663:特化会把 T.serial 外层循环之外的前导语句搬进循环体,本应只执行一次的初始化被每轮重放,输出错。#3664:if 内包特化流水且 if 后有语句时,setmaxnreg 相关下降令 ptxas 失败。三笔合看,「哪些语句属于哪个 warp、生命周期到哪结束」仍缺完整语义模型。

来源:https://github.com/tile-ai/tilelang/pull/3645、https://github.com/tile-ai/tilelang/issues/3663、https://github.com/tile-ai/tilelang/issues/3664

1.8 新增缺陷:sm_120 上归约下降失败、三张注意力示例超共享内存(#3647 / #3646)

10-25 09:08 同日两报。前者是 gdn wy_fast_bwd_split 在 sm_120 上的 ReduceOp 下降失败——源索引依赖线程持有的归约分段,不在既有规则覆盖内;后者是 attention_sink、gqa_bwd、block_causal_attention_varlen 三张示例超 101376 字节单块共享内存上限。对应修复 #3662 已开:按设备上限选更小瓦片、更少流水级与 128 线程,大设备保持原默认。消费级 Blackwell 的容量约束正成为示例与调优表的硬边界。

来源:https://github.com/tile-ai/tilelang/issues/3647、https://github.com/tile-ai/tilelang/issues/3646


二、多后端适配(昇腾 / 海光 / 新增 CPU 与 ROCm)

2.1 昇腾主仓:DeepGEMM 示例支持 FP4×FP4 打包输入(#3589)

合于 10-23 10:26。为 dense、M-grouped、K-grouped 三套 DeepGEMM 接口打开打包 E2M1 FP4×FP4 输入,复用既有 MX GEMM kernel 与启发式,从打包字节存储反推逻辑维度并按 FP4 的 64 元素对齐;后端侧只改 MX 的 K 尾填充,批式构造器也接受 FP4。公开的 batch/einsum 接口仍保 FP8 以对齐 DeepGEMM-Ascend-Nightly,FP8×FP4 反量化留待后续。

来源:https://github.com/tile-ai/tilelang/pull/3589

2.2 昇腾主仓:DMA 保留缓冲基偏移、区域界先提升、周期缓冲完整生命周期(#3627 / #3625 / #3606)

三笔同于 10-23 合入,都是布局与生命周期修正。#3627:DMA 规划丢掉 Buffer.elem_offset,UB 别名以 elem_offset=256 构造时 T.copy(别名, out) 实际从 0 号元素读;修法是把基偏移算进地址并在跨步切片(含嵌套)保留父偏移。#3625:区域最小值与长度用 Int64Promoter 提升后再构造半开区间,避免 int32 转型遮住动态尾界、令重叠证明耗尽 Z3 预算,给原地持久化 RoPE 引入多余跨迭代依赖。#3606:子缓冲生命周期摘要须完整覆盖访问才可提升,外部预载的值跨整个父周期存活,防止嵌套与手工流水下的复用覆盖待读者。

来源:https://github.com/tile-ai/tilelang/pull/3627、https://github.com/tile-ai/tilelang/pull/3625、https://github.com/tile-ai/tilelang/pull/3606

2.3 昇腾主仓:CI 镜像升到 cann9.3.0 周更版(#3630)

合于 10-23 19:20,只改 CI:ascend 作业测试容器从 cann9.3.0-weekly.20260930 提到 cann9.3.0-weekly.20261021-ubuntu22.04-py3.12。它决定昇腾侧全部回归用例实际验证的是哪版 CANN 运行时,是判断该线结论有效期的前提。

来源:https://github.com/tile-ai/tilelang/pull/3630

2.4 昇腾在审:E4M3 的 UB→L1 拷贝、片上内存溢出提前拒绝(#3639 / #3655)

#3639 为 Ascend 950 打开 float8_e4m3fn 的 UB→L1 拷贝:ND 输入经既有 scatter/post-copy 管线重排、预打包 NZ 直接搬运,footprint 与延迟分析补 FP8 元素宽度。#3655 让 MergeUBAllocations 在下降阶段按 UB、L1、L0A、L0B、L0C 各自容量校验计划占用并报诊断,而非生成越界分配——起因是两个独立循环各自合法地选中 32KB L0A 瓦片的两个版本、合并后超限。片上内存管理正从运行时报错前移到编译期拒绝。

来源:https://github.com/tile-ai/tilelang/pull/3639、https://github.com/tile-ai/tilelang/pull/3655

2.5 昇腾在审:PTO 目标默认走 tvm_ffi 执行后端(#3516)

10-24 更新的在审改动:为 pto 支持 tvm_ffi 并设为默认,execution_backend="auto" 现解析到 tvm_ffi。共享宿主机 codegen 负责 MakePackedAPI 与 C 宿主构建,昇腾运行时模块经 aclrtLaunchKernelWithHostArgs 拉起 PTODSL 编译的 aibin,与 AscendC 路径一致。与主仓 CPU 侧 #3634(统一到 TVM-FFI、移除 CPU 专用 Cython 包装)同向:执行后端在做减法。

来源:https://github.com/tile-ai/tilelang/pull/3516

2.6 昇腾仓:situ_glu 算子合入,对齐 Kimi K3 门控激活(#1891)

合于 10-23 16:04。语义与 dtype/shape/attr 约束逐条对齐 CANN ops-nn 的 activation/situ_glu,精度基准取 CPU fp64 的 PyTorch 金标准;纯 Vector 单 kernel,宿主按 dim 把输入合轴成 (rows, cut*inner) 零拷贝视图,非末维切分不跨步。算子对应 Kimi K3 Stable LatentMoE 的前向门控激活,与上一期的 situ_glu 梯度、稀疏 MLA 一并,昇腾仓的算子清单正按具体对话模型结构逐个铺开。

来源:https://github.com/tile-ai/tilelang-ascend/pull/1891

2.7 昇腾仓在审:稀疏 MLA、GLM-5.3 kpool 与三张规范化示例(10-23~25)

在审 24 笔里,算子侧代表是 sparse_mla_fwd(#1929,DSA 稀疏 MLA 前向,varlen 打包随机行 gather + 在线 softmax + QK/PV 双 GEMM,fast 路径限定 H=128、G=1、dim/tail=(512,64) 等条件)与 GLM-5.3 kpool 压缩(#1910,池内逐维 softmax 池化、过 BF16 边界、Hadamard-128 旋转、E4M3 量化并把行与 scale 散射进分页 cache)。示例侧补 rms_norm(#1558)、NMS(#1748,成对 IoU 压 1 bit 进 uint16 位图,N=4096 时 2MB 取代 64MB 矩阵)与 GroupNorm(#1607)。

来源:https://github.com/tile-ai/tilelang-ascend/pull/1929、https://github.com/tile-ai/tilelang-ascend/pull/1910、https://github.com/tile-ai/tilelang-ascend/pull/1748

2.8 昇腾仓:每日测试 2903 项全通,次日测试作业报失败(#1941 / #1944)

机器人账号 10-25 05:15 报每日测试 2903 项全部通过、通过率 100%;10-26 04:30 又报 workflow 的 test 作业失败,10-24 05:08 另有一次失败记录。#1944 未附失败用例明细,但从时间序列看,每日全量已稳定在 2900 量级并承担回归门禁角色,一次失败即代表某条新合入或环境变更打断既有用例。

来源:https://github.com/tile-ai/tilelang-ascend/issues/1941、https://github.com/tile-ai/tilelang-ascend/issues/1944

2.9 海光:主干三笔合入,寄存器流水线重新进在审(#26 / #18 / #27)

主干在 10-23 17:57 一次性落地三笔(作者日期 10-14、10-16、10-20,属批量合入):修正 MLS 槽位与线性 DS 读取、CI 支持 DTK 26.04/26.10、改用 wave64 shuffle 与向量化 BF16 存储;合入的 #26 另修 FP4 的 MLS warp 划分句柄。在审 #18 重新引入寄存器流水线(可与共享内存流水并存,按流水级插入 commit 与 barrier),并在 warp 数超 4 时启用 warp diverge,目标把 example_gemm 打到与 ROCm 相当,10-25 06:00 有更新;#27 处理内建下降、精度与整数向量 codegen。与上一期「合并分支移除 enable_register_pipeline」对照,同一能力正以 PR 形式重做。

来源:https://github.com/tile-ai/tilelang-hygon/pull/26、https://github.com/tile-ai/tilelang-hygon/pull/18、https://github.com/tile-ai/tilelang-hygon/pull/27

2.10 新增后端在审:Intel AMX 的 BF16/INT8 GEMM、ROCm 的 FP4 与 E8M0(#3637 / #3665)

#3637 为 CPU 路径加 Intel AMX 后端:目标探测、编译选项、GEMM 分派,BF16 与 INT8 kernel 走 AVX-512/VNNI 打包与对齐复用暂存,执行侧按 1 核、40 核、120 核自适应 OpenMP。此前 TileFoundry 已在 Apple AMX 上跑通预填充示例,这一笔把 AMX 直接做进主仓 CPU 路径。#3665 修 HIP 模板两处缺口:T.cast(x, T.float4_e2m1fn) 在 HIP 上被打印成 (fp4_e2_t)x,而该类型只有显式构造函数,浮点被截成整数存储(1.0 变 0.5、2.0 变 1.0),同时补 E8M0 类型并在 gfx94x 上使能软件 FP4。

来源:https://github.com/tile-ai/tilelang/pull/3637、https://github.com/tile-ai/tilelang/pull/3665


三、生态与采用方(TileOPs / TileFoundry / 发行版)

3.1 TileOPs:首发清理八项完成七项,CHANGELOG 写定 0.1.0(#2738)

10-25 10:45 建跟踪单「Release clean-up」,八项已勾选七项:命名、打包与仓库文件、Python 3.13、公共 API 与扩展点、算子行为、输入校验、文档站;唯一未完成的是「第三方声明、PyPI 安装文档、dev 依赖组、CHANGELOG 与发布步骤,TestPyPI 试跑并打首个 tag,由维护者执行」。CHANGELOG 同期落地首条版本记录「0.1.0 - 2026-10-26」,注明首次发布、docs/stability.md 说明何为公开接口。核对 PyPI:tileops 最新仍是 2025-10-09 的 0.0.1.dev2;核对 GitHub:无 tag、无 release。另有 #2756 提前开出「Post-release clean-up」。

来源:https://github.com/tile-ai/TileOPs/issues/2738、https://github.com/tile-ai/TileOPs/blob/main/CHANGELOG.md

3.2 TileOPs:公共名字与路径一次定稿,打包改元数据并支持 3.13(#2741 / #2742 / #2746)

#2741 是 739 文件 +6452/-6740 的一次性改名:tileops、各算子族模块、tileops.backend、tileops.manifest、tileops.trace 与新增 tileops.authoring 成为公开面,其余下沉为 _ops、_kernels、_utils、_perf,删掉 tileops.ops 惰性导出层,一个算子只留一条导入路径。#2742 修打包元数据(bench 从 extra 移到 dependency group,dev 去重并补 pre-commit 与 pytest-timeout,声明 Linux/CUDA 13/Python 3,随包发 py.typed),新增 CHANGELOG 与缺陷模板。#2746 把 requires-python 定为 >=3.12 并去上界,在 Python 3.13.15 + torch 2.13.0 + PyPI 版 tilelang 0.1.12 的 H200 上跑通全部 43 项打包用例。

来源:https://github.com/tile-ai/TileOPs/pull/2741、https://github.com/tile-ai/TileOPs/pull/2742、https://github.com/tile-ai/TileOPs/pull/2746

3.3 TileOPs:写算子的入口收口为 authoring,eager 调用收敛到单执行器(#2743 / #2744)

#2743 定扩展点:仓外代码按 tileops.authoring.Op 子类化写算子,把清单条目声明为类属性 spec,类创建时由校验器检查,随之获得签名检查、形状推导、roofline 与编译边界;只有 tileops 内部类取同名条目并由后端直接服务。#2744 把每次 eager 调用重复推导的签名判定、调用规格、kernel 程序与调用记录收敛为「每个特化一份 plan」,并发下只解析一次、构建失败回滚已发布的调优配置。算子作者与调用者的接口都从隐式约定变成显式契约。

来源:https://github.com/tile-ai/TileOPs/pull/2743、https://github.com/tile-ai/TileOPs/pull/2744

3.4 TileOPs:算子的共享语义一次写清,输入校验由清单生成(#2751 / #2758 / #2760)

#2751 与 #2758 把「每个概念一种约定」落成文档并逐算子执行:所有 attention 的 LSE 统一为 float32 自然对数、稠密 [B,H,S] 与打包 [H,T] 按该布局存储(GQADenseBwdOp 此前读的是 2 为底),无可见键的行输出 0 与 -inf;out= 一律为仅关键字参数并返回该对象,反向算子与对应前向成对。#2760 把 validate_inputs 的打包偏移检查从逐算子手写改为由清单生成,顺带修掉四个线性注意力算子按参数位置取 cu_seqlens 的隐患。

来源:https://github.com/tile-ai/TileOPs/pull/2758、https://github.com/tile-ai/TileOPs/pull/2760

3.5 TileOPs:发布文档、第三方声明与 pip 安装入口(#2753)

11 文件 +103/-32:THIRD_PARTY_NOTICES 收录 Qwen 团队代码的 MIT 声明(两个 GDN 预填充 kernel 改写自该处),作为 license 文件随 wheel 与 sdist 发布;README 首屏改为 pip install tileops 并给出核对过的配置(Linux、Python 3.12/3.13、PyTorch 2.13、TileLang 0.1.12、CUDA Toolkit 13.2 与 nvcc、SM90)。提交说明要求打 tag 前一刻合入,与跟踪单最后一项互为先后。

来源:https://github.com/tile-ai/TileOPs/pull/2753

3.6 TileOPs 性能:GDN 预填充分区宽度按调用自身长度选,KDA 原生门控进 kernel(#2745 / #2683)

#2745 修 GDNFwdOp 预填充:此前分区宽度取 key head 数,q35-397b-tp8-32768 把每条 walk 切成 64 段跨 4 波,夜间基准只有 FlashInfer 的 0.988;改为按调用自身序列长度与 value head 求解分区后,bf16 从 289.3 降到 200.6 微秒(1.434 倍 FlashInfer)、fp16 从 285.0 降到 195.2 微秒(1.468 倍)。#2683 让 KDA 前向直接吃原始门控(含可选 dt_bias 与有界门控),16 文件 +3095/-458,新增 28 条工作负载并在每条 KDAFwdOp 行超过 FLA 基线,强衰减 chunk 改算成对 Gram 项以避免 NaN。

来源:https://github.com/tile-ai/TileOPs/pull/2745、https://github.com/tile-ai/TileOPs/pull/2683

3.7 TileOPs 性能:合并注意力态、分页 GQA 与 indexer 尾波(#2739 / #2752 / #2735 / #2730)

#2739 给 MergeAttentionStatesFwdOp 加无行守卫的程序与合并后一次性 LSE 存储,此前在 llama-70b-cascade 上只有 FlashInfer 原生 merge_state 的 0.985、另两行 0.988 至 0.999。#2752 让它同时支持自然对数与以 2 为底(exp2/log2),单波 16 行网格改 512 线程 CTA 处理 64 行。#2735 修 qwen35 的 RoPE 暂存溢出,head dim 256 的 FP8 缓存从单 warpgroup 64 行瓦片改双 warpgroup 128 行(此前只有最快对比实现的 0.16 至 0.57 倍),flat dim-64 行改三消费者 warpgroup。#2730 把 FP8 indexer 最后不足一波的块按 2 的幂切到最多 16 个 CTA,但实测长行因 ptxas 的 FMNMX 排序反慢 2 至 7 个百分点。

来源:https://github.com/tile-ai/TileOPs/pull/2739、https://github.com/tile-ai/TileOPs/pull/2735、https://github.com/tile-ai/TileOPs/pull/2730

3.8 TileOPs 性能:GEMM 尾波窄瓦片、乒乓 BMM 分级暂存(#2733 / #2734 / #2720)

#2733 针对持久化网格最后一波:最忙 CTA 的瓦片列数决定整体时间,ds-v3 若干行此前落后 cuBLASLt 或 quack 0.6 至 2.3 个百分点;改法是宽瓦片跑完后用更窄瓦片处理最后一行、按最忙 CTA 列数最少选布局,176 宽瓦片配 5 级环与双缓冲 epilogue。#2734 让 BmmPingPongKernel 的存储经两块缓冲交替分级(K 跨 16 段以上取 16 列、以下取 32 列),避免最后一个瓦片的 epilogue 串行等 cast 与 store。#2720 在两波以内的短 K 批量 GEMM 改用 128×128 乒乓瓦片,让换瓦与 epilogue 藏在另一组主循环下。

来源:https://github.com/tile-ai/TileOPs/pull/2733、https://github.com/tile-ai/TileOPs/pull/2734、https://github.com/tile-ai/TileOPs/pull/2720

3.9 TileOPs 性能:MoE 的 A 箱装载与原生宽度 WGMMA,逐元素倒数与行广播(#2711 / #2715 / #2710 / #2725)

#2711 修解码期 MoE:一个专家只持 1 至 7 行,但 64 行整块 TMA 装载把下一个专家的行也拉进 L2,gpt-oss-120b-decode 的 gate/up 因此 L2 受限(3.56GB DRAM 读对应 6.11GB L2 流量);改为按 128 字节 K 原子切 16 行 A 箱、只发活跃行可达的箱。#2715 让 192 与 224 宽模板瓦片各发一条 16K 的 WGMMA——此前宽度不整除 256 的 T.wgmma_gemm 被拆成 gcd(width,256) 宽多条、每条重读 A,慢 27 至 35 个百分点。#2710 把倒数改成一条近似指令加每线程一个向量(此前 1600 万行 fp16/bf16 上落后 torch.compile 1 至 1.4 个百分点)。#2725 在网格只占一波时让一线程走四行、共用一份 bias 向量。

来源:https://github.com/tile-ai/TileOPs/pull/2711、https://github.com/tile-ai/TileOPs/pull/2715、https://github.com/tile-ai/TileOPs/pull/2710

3.10 TileOPs 性能:归约与归一化改 16 字节 evict-last 读(#2757 / #2759 / #2748 / #2737)

#2759 把只读一次的 16 字节装载统一走 tileops::load16_evict_last(argreduce 共享内存段、逐线程整向量 argreduce、InstanceNorm 行归一化、8MiB 以内行折叠),超过阈值才回 evict-first。#2757 让 leading-axis 归约下行 pass 改由 1024 线程各持一个 16 字节向量,ProdFwdOp 在 (4,128,4096) fp16 的 dim 0 上从 3.33 微秒降到与 torch.compile 持平。#2748 让 LayerNorm 线程数按行数选(8192 以上且每 SM 不超过一块时 256、否则 128),同一算子服务解码与预填充各编译一次。#2737 停掉 AdaptiveMaxPool2d 的调优——按 CUDA 事件计时选出的配置轮次间可慢 1.7 倍,改跑派生默认配置并新增寄存器版全局池化 kernel。

来源:https://github.com/tile-ai/TileOPs/pull/2759、https://github.com/tile-ai/TileOPs/pull/2757、https://github.com/tile-ai/TileOPs/pull/2748

3.11 TileOPs 正确性与 CI:SSD 扫描尾部、GLA 段展开、H200 发布校验(#2716 / #2750 / #2749 / #2718 / #2713)

#2716 在 CUDA codegen 前展开 SSD chunk 扫描流水的固定尾部,修掉异步拷贝等待参数里残留算术表达式导致的整型下溢编译错(Mamba 用例此前无法编译)。#2750 显式展开打包 GLA 的两段流水,避开新版 CUDA 代码生成器拒绝符号化等待计数。#2749 在向量化因果卷积预填充装载前对齐可选 bias 存储——连续视图可能保留未对齐首址。#2718 收紧 H200 发布校验:用例缺失、报告为空、测试未成功或整行未核验都判失败,每次装轮验证用独立缓存,CUDA Graph 覆盖 23 个测试、9 个算子。#2713 让夜间基准在独占 GPU 上串行采集——此前提前预热会在另一基准占用独占设备时初始化 CUDA 上下文,致一次夜间 1061 个用例报错。

来源:https://github.com/tile-ai/TileOPs/pull/2716、https://github.com/tile-ai/TileOPs/pull/2718、https://github.com/tile-ai/TileOPs/pull/2713

3.12 TileFoundry:sm_89 的 E4M3 warp MMA 与逐指令能力校验(#261 / #262 / #264)

#261 声明 sm_89 上的 FP8 warp MMA(mma.sync 的 m16n8k32 E4M3 形式):此前 #259 建立的 sm_89 目标只能选 BF16,寄存器里的 E4M3 乘积无指令可落。#262 把能力校验从「只查 MMA」扩到每条指令:sm_89 上用 bulk copy、TMA copy 或 mbarrier 事务屏障形式此前能通过验证、到 codegen 或 ptxas 才失败且报错不提缺失硬件,改后 admit_fields 在接入阶段逐条对照目标能力。#264 用一条 16×8×128 的 E4M3 乘积把该 atom 从调度跑到 GPU、K 切片经共享内存暂存——正是这条用例暴露了随后 #263 修的 cp.async 寻址缺陷。

来源:https://github.com/tile-ai/TileFoundry/pull/261、https://github.com/tile-ai/TileFoundry/pull/262、https://github.com/tile-ai/TileFoundry/pull/264

3.13 TileFoundry:cp.async 目标布局寻址与 MLA 反向布局(#263 / #260 / #257)

#263 修 ops::copy_async 按元素下标寻址目标、源按向量遍历的不一致:线程切片含多个向量且连续段不在第一个模式上时,后续向量落到错误地址——sm_89 E4M3 GEMM 中每线程切片 (2,1,1,8)、步长 (256,-,-,1),第二个 8 字节向量写到了第 4 字节而非第 256 字节。#260 收口 MLA 反向的布局与分析缺口(保留作者仿射映射、用 CuTe 图元做切片与指令投影、补 StMatrix 与累加式 CopyReduceAsyncTensor 语义),关掉 #251。#257 修下降后 TIR 访存路径三处缺陷:ops::reduce 把线程轴误判为 cta 级,warps_per_group 取成整个网格实例数(RMSNorm 教程里 256 而非 8),跨 warp 工作区索引读越界。

来源:https://github.com/tile-ai/TileFoundry/pull/263、https://github.com/tile-ai/TileFoundry/pull/260、https://github.com/tile-ai/TileFoundry/pull/257

3.14 TileFoundry:拆成 core / nvidia / apple 三个发行包,测试输入只装载一次(#268 / #269)

#268 是 739 文件 +5489/-4518 的结构调整:核心与厂商发行版拆开、共用一套版本契约,CUDA 与 AMX 的实现、硬件文档、规格、模板、运行时头文件各自打包,发射与链接实现经 target 绑定解析,安装后的目标在核心初始化之后加载;模型资源归 target、共享因果语言模型编排归 core,评审后净减 274 行。#269 针对测试耗时(主干 314 至 509 秒、运行时头文件变动 651 秒),把同一输入经 CLI 装载器只跑一次——此前一份调度 HIR fixture 被装载 12 次、校验 9 次、每族分析一次,且测试私有装载器接受了 CLI 会拒绝的程序。

来源:https://github.com/tile-ai/TileFoundry/pull/268、https://github.com/tile-ai/TileFoundry/pull/269

3.15 分发版:Sunrise 0.1.15+sunrise.1.2.0 发布 PR 合入(#9)

10-23 10:43 合入。该 PR 把 Sunrise 分发版更新到 0.1.15+sunrise.1.2.0,纳入 TileLang 0.1.15 及配套评审过的 TileOPs、TileKernels、TVM、tvm-ffi 变更,源导入与合并适配分开提交,最终适配 TANG 后端 wheel 并修正切片 GEMM 回归(两个累加器标志都对公开模板校验,保留 stride、origin 与 offset 断言)。上一期记录过该 PR 在 S2 上装包后 import 冒烟在 tangGetDeviceCount 处失败;本窗口该仓无后续提交,未见针对该失败的修复记录。

来源:https://github.com/tile-ai/tilelang-sunrise/pull/9


四、社区、教程与活动

4.1 社区:tilelang-puzzles-ascend 做 PTO 对 Ascend 的量化 kernel 教程(10-23 推送)

10-23 推送,自述为「PTO 对 Ascend 的量化 kernel 教程与练习」,格式沿用 tile-ai/tilelang-puzzles。同组织的 tilelangGym 用于测不同基准与微基准在 PTO 下降下的性能敏感度,已于上一期收录。两仓同属 PTO-ISA 组织,与主仓在审的 PTO 向量化(#3512,把 T.Parallel 向量化成 PTO VMI 指令)形成「编译器—练习—敏感度」三件套,说明 PTO 这条昇腾指令线正在补齐开发者入口。

来源:https://github.com/PTO-ISA/tilelang-puzzles-ascend

4.2 新仓:中文 TileLang 算子面试练习站,9 模块 18 题(10-23)

10-23 建仓,自述为中文算子面试练习站:9 个模块、18 道题、21 份教学示例,含交互练习、题解与性能分析。此前社区内容以 cookbook 与技能包为主,这是首次出现「按面试/课程组织」的产物。

来源:https://github.com/zhachar14h-cell/tilelang-operator-interview

4.3 社区技能与教程仓继续推送,文档站四笔跟随首发改版(10-23~25)

sablin39/tilelang-cuda-skills(146 星)、RuneFang/tvm_tilelang_cookbook(36 星)、cc-c122/TileLang-Remote-Autotuning-Agent(10-24 推送)继续更新;kevin3113/models-tilelang(10-22 建仓、10-23 推送),其作者同期向主仓提交昇腾跨核同步的 schedmode 支持(#3614)。文档侧,TileOPs 文档站四笔跟随主仓改版:#79 删掉已退休的 GQAPrefillPagedWithKVCacheFwdOp 并在首页挂发布公告,#81 跟随 26 个算子类改名与内部包加下划线,#82 跟随公共 API 契约与 bench 依赖组变更,#83 为首次 PyPI 发布整理首页与安装页。

来源:https://github.com/tile-ai/TileOPs.github.io/pull/83、https://github.com/sablin39/tilelang-cuda-skills、https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent

4.4 讨论:TileScale 用户询问所用 TileLang 版本能否升级(#72,10-25)

10-25 05:17 有用户在 tile-ai/tilescale 提问能否升级 TileScale 依赖的 TileLang 版本。该仓最后推送停在 08-25,文档与示例仍指向较早版本。这类提问通常出现在上游接口进入快速改名期之后——本窗口主仓在审的 TIRx 替换(#3666)与 TileOPs 的 26 个算子类改名,都会放大下游跟版成本。

来源:https://github.com/tile-ai/tilescale/issues/72


五、趋势观察

发布工程成了主战场。 78 笔合入里 TileOPs 占 51 笔,其中十余笔属改名、打包、语义统一、文档与 CI 门禁——不是性能也不是新算子,而是把接口冻住。跟踪单七成、CHANGELOG 写定 0.1.0 并标注 10-26,此后 0.x 的兼容承诺开始生效;同期 TileFoundry 做成 core/nvidia/apple 三包、Sunrise 合入 0.1.15 分发版,三条线在同一周做同一件事:把「能跑」变成「能发」。

正确性重心从形状校验转向新硬件与新调度。 上周主仓合入多是形状与几何校验,本窗口换成生命周期与屏障:跨存储域屏障重号、sm_100 张量内存释放缺 CTA 屏障、warp 特化搬错前导语句、setmaxnreg 触发 ptxas 失败。共同点是编译通过而结果偶发错,且都落在自动 warp 特化与新架构存储上——把调度权交给编译器之后必然要付的账。

昇腾线从跑通走到按模型铺算子加每日回归。 situ_glu 对齐 Kimi K3 Stable LatentMoE 门控激活,在审的 sparse_mla_fwd 与 GLM-5.3 kpool 同样是特定模型结构的部件。每日全量 2903 项全通、次日又报失败,说明用例规模已足以当门禁;配合 CI 镜像追到 cann9.3.0 周更版,这条线的结论有效期以周计。

执行后端做减法,后端清单做加法。 CPU 统一 TVM-FFI、PTO 默认 tvm_ffi,加 CuTeDSL 修复与 NVRTC 编译,方向一致:减少执行路径分叉。后端清单反向扩张——Intel AMX 进 CPU 路径、ROCm 补 FP4 与 E8M0、海光重提寄存器流水线。分发形态正从「一套代码多后端分支」变成「一套 IR、多个打包」,与 #3666 把 TVM 变外部依赖同路。


附录:素材与核查说明

信源核查表

信源 核查结果
tile-ai 组织仓库 29 仓,窗口内推送 11 个;合入 78 笔,在审更新 76 笔,新增 issue 19 条
tilelang 主仓 10 笔合入(9 笔在 10-23、1 笔在 10-25);VERSION 仍 0.1.15,无新 tag;在审 50 笔更新,含 TIRx 替换(#3666)
TileOPs 51 笔合入,围绕首发清理与性能;无 tag、无 release;PyPI 最新仍为 0.0.1.dev2(2025-10-09)
TileFoundry 10 笔合入(#257 至 #269);v0.0.3 发布说明已在 10-22 那期收录,本窗口无新 tag
tilelang-ascend 1 笔合入(situ_glu,默认分支 ascendc_pto),在审 24 笔更新;每日测试 2903 项全通后次日作业失败
tilelang-hygon 主干 3 笔合入(作者日期 10-14 至 10-20,属批量合入);在审 2 笔(#18 / #27)
tilelang-sunrise 1 笔合入(0.1.15+sunrise.1.2.0 发布 PR),合入后无后续提交
tilelang-mlir-ascend 默认分支 main 窗口内 0 笔提交;推送发生在 update0.1.15 / ci-0.1.15 / probe-ci 三个分支,属版本适配与 CI 试探
tilelang-metax / musa 窗口内无推送,分别停在 10-16 / 09-30
TileRT / tilescale / DeepStack 无推送;tilescale 收到一条版本升级询问(#72)
采用方 TileKernels / FlashQLA 均停在 10-14,无推送
CANN 社区 ops-tilelang 算子仓 默认分支最后提交为 10-22,窗口内无新增
Google News(中英文多组,走代理) 精确命中 1 条(SGLang 博客),核对原文全文未提 TileLang,按无关剔除
Hacker News 命中 2 条,均为 AI 投资类讨论,与主题无关
arXiv 以 TileLang 为标题或摘要关键词检索,最新为 2026-07,窗口内 0 条
PyPI tilelang 0.1.15(09-30)、tilefoundry 0.0.3(10-21);tileops 仍为 0.0.1.dev2
社区新仓 1 个新仓(tilelang-operator-interview)、1 个首次推送仓(tilelang-puzzles-ascend)与 4 个活跃仓,均已核对建仓日期与推送时间

完整信源清单