TileLang 每日动态报告(2026-11-04)
调研窗口:2026-11-03 07:00 ~ 2026-11-04 07:00(北京时间,24 小时;上一期截点为 11-03 07:00)。
信源:GitHub tile-ai 组织逐仓核查(29 仓库,窗口内推送 9 个)+ 合入与在审真值检索(合入 8 笔、在审更新 48 笔)+ 新增 issue 全量核对(7 条)+ 主分支逐笔提交核查(tilelang 5、TileOPs 1、TileFoundry 1、VS Code 扩展 1、文档站 3)+ 发布页与 tag 核查(5 仓,窗口内无新版本)+ 采用方两仓 + Google News 中英文查询(走代理)+ Hacker News + arXiv。
本期索引
- 今日重点:ThreadSync 跨存储域复用部分屏障,把 #3643 修正确性时引入的性能回退补回(#3761);清理 3384 行遗留 TVM 解析器分支(#3736);自动调优输入供给改按当前设备(#3546);TileFoundry 同步计划正式落到 CUDA lowering(#285);24 小时合入 8 笔、在审更新 48 笔、新增 issue 7 条
- 一、核心项目进展
- 1.1 主线:ThreadSync 允许跨存储域复用部分屏障,B300 上屏障占用从 12 回到 16 块/SM(#3761)(11-04)
- 1.2 主线:移除遗留的 TVMScript AST 与解析器分支,一次删掉 3384 行(#3736)(11-03)
- 1.3 主线:自动调优的输入供给改用当前设备,昇腾与 MPS 上的调优恢复可用(#3546)(11-03)
- 1.4 在审:条件中的标量 cast 不再压缩数据向量宽度,承接昨日 #3734(#3759)(11-03)
- 1.5 在审:两处用户可见错误改为明确报错(#3765 / #3767)(11-04)
- 1.6 在审:带 padding 的共享 tile 无法整体拷贝时回退,2-CTA 的 TMA 字节按对等 leader 记账(#3738 / #3766)(11-03)
- 1.7 在审:B300 的 sm_103a 适配两条线,tcgen05 的 int8 选择与 CuTeDSL 的 TMA 架构门(#3770 / #3772)(11-03)
- 二、多后端适配
- 2.1 昇腾:3D GM 到 L1 的批量 ND2NZ 拷贝合入,一次描述符覆盖整批(#3689)(11-03)
- 2.2 昇腾在审:TileLang-Ascend 整仓迁入主线多后端架构,2887 个文件(#1975)(11-04)
- 2.3 昇腾在审:copy_pipe_to_l1 的同步标注按执行引擎拆分,修 A2/A3 上 10/10 复现的错结果(#1981)(11-03)
- 2.4 昇腾在审:首 Parallel 轴广播走标量读、Chunk O 前向迁移、MHA 示例(#1979 / #1866 / #1879)(11-03)
- 2.5 mlir-ascend:Codex 工作流与任务交接合入 harness 分支(#208)(11-03)
- 2.6 海光:Distributed 在审 PR 有更新,主分支无新提交(#19)(11-03)
- 2.7 沐曦、摩尔线程窗口内无推送(11-03)
- 三、生态与采用方
- 3.1 TileFoundry:同步计划落成 CUDA lowering 并反向校验,合入主线(#285)(11-04)
- 3.2 TileFoundry 在审:wgmma 同步计划、FSDP 与 HSDP 训练语义进 HIR 求值器(#286 / #284)(11-03)
- 3.3 TileOPs:方差合并从第一个部分量起算,常量输入的方差不再出 NaN(#2821)(11-03)
- 3.4 TileOPs 在审:GEMM 加可选 bias、FP8 分组 GEMM 支持连续布局(#2823 / #2824)(11-03)
- 3.5 TileOPs 在审:W4A16 与稠密 GEMM 的选择器改成单一代价模型重拟合(#2825 / #2826)(11-03)
- 3.6 TileOPs 在审:短 varlen 按 KV 组打包、L2 flush 改为只读(#2819 / #2822)(11-03)
- 3.7 采用方:FlashQLA 与 TileKernels 无推送,但 #3761 修的是 TileKernels 的 fused 内核(11-04)
- 四、社区、教程与活动
- 4.1 单日 7 条新 issue:一条生产级 CUDA graph 问题,四条来自同一外部报告人(11-03)
- 4.2 昇腾 950 的 vcvt 饱和语义与 mlir-ascend 日测(#3762 / #1980)(11-03)
- 4.3 文档站三次更新,VS Code 扩展改名 vsix(11-03)
- 4.4 新闻侧:三通道无窗口内命中,外媒的昇腾训练报道口径不一(11-03)
- 五、趋势观察
- 附、素材与核查说明
今日重点
这一天的主线是一笔「把性能要回来」的修复。上一窗口合入的 #3643 解决了一个真实的正确性问题:分别调用的 ThreadSync("shared") 与 ThreadSync("shared.dyn") 可能把同一个命名屏障 ID 分配给互不相交的线程组。修法是让后一次调用保留前一次已经用掉的 ID,但这条保留规则同时挡住了另一件事——同一组线程在两个作用域之间本来可以复用同一个 ID。代价直接落在 TileKernels 的 mhc_pre_big_fuse 内核上:这个内核每块 96 个线程,线程 32 到 95 跑两段式 cp.async 流水线,异步等待之后的同步在第一遍拿到 ID 3,同一批 64 个线程的缓冲依赖同步在第二遍拿到 ID 4,多出来的那个 ID 让编译产物声明的屏障资源从 4 变成 5;在 B300 上跑隐藏维 1280 时,屏障占用上限因此从 16 块/SM 降到 12 块/SM,寄存器、共享内存与同步指令数都没有变化。#3761 的处理办法是让 ThreadSync 接受一个有序的作用域列表,CUDA 流水线改成 ThreadSync(["shared", "shared.dyn"]),同步仍按原来的作用域顺序规划与插入,但部分屏障的 ID 在两个作用域都处理完之后才统一分配;跨作用域共用一个 ID 之前,用一个全新的分析器证明参与谓词在整个候选线程范围内成立,分支与 Bind 上下文由 ConstrVisitor 提供,线程相关的谓词按证明义务处理而不是默认成立,已证密的密集组直接用包围盒体积、跳过部分线程数的枚举,未证明的组则把作用域留在键里并让参与者计数一起入键,显式编号的屏障(含更早一次 ThreadSync 分配的 ID)继续保留。改动 401 行、涉及 4 个文件。
作者在描述里做了一个很干净的归因实验:单独改编译产物里声明的屏障资源数、不动任何可执行指令,就能复现同样的变慢。这类实验值得记一笔——它把「性能回退」从观感变成了可指认的单一变量,也让修复的验收标准变得明确:不是笼统地跑分变快,而是资源声明回到 4、占用回到 16 块/SM。
第二条是供给侧的减法与口径统一。#3736 把 tilelang/language/ast/、tilelang/language/parser/ 与 tilelang/language/tir/entry.py 一并删掉,这些是早期本地 TVMScript 解析器分叉留下的模块,约 3380 行 Python,仓内已无调用方:语言层现在通过 tilelang.language.tir.common 使用上游的 tvm.tirx.script.parser 接口,AST 与语句覆写注册在 tilelang.language.overrides,T.prim_func 与 T.macro 由 tilelang.language.eager 导出。实际改动是新增 1 行、删除 3384 行、10 个文件,同时移除了 docs/conf.py 里针对 *language/ast* 的忽略规则;公开的语言层导出由方言测试与前端口测试兜住。#3546 修的是自动调优的设备口径:set_autotune_inputs 安装的供给闭包从一个调用点从不传入的参数上解析基准设备,实际传过去的是内核参数列表,于是 CUDA 或 CPU 的兜底分支在没有 CUDA 的机器上把每个捕获张量搬到 CPU,每一轮调优都因启动器设备不匹配失败;现在闭包改用 get_current_device(),与默认张量供给同一策略,CUDA 行为不变、昇腾与 MPS 改为在当前设备上供给、其余后端仍退到 CPU。验收在双卡 Ascend950 上:三笔昇腾自动调优回归修前全败、修后全过,主仓 testing/python/autotune/ 12 项通过、28 项跳过,另有 3 项因该机没有 CUDA 而失败,作者明确留给 CUDA CI 验证。
第三条在昇腾侧:GM 到 L1 的拷贝从「一个描述符一个矩阵」扩到整批。#3689 允许把 [B, M, N] 区域写成一个 T.copy,lowering 生成批量 ND2NZ 描述符,保留源侧的批与行 stride,目的侧的批间距按物理布局推导,转置批量走 DN2NZ;超出 GM 边界的源批直接跳过并保持目的存储不动,M/N 的尾部补零只作用于合法批。既有二维与单矩阵行为不变,实现只针对 AscendC 后端、且要求 L1 目的端是规范矩阵布局,PTO 不在本次范围内。验证在 Ascend950DT 与 CANN 9.2 上:50 个 AscendC 示例测试与 1733 个昇腾单测通过,其中 28 个设备用例覆盖批与 M/N 尾块、整批越界、非连续源 stride、目的存储更大与矩形转置;CUDA 与 PTO 用例被排除,CI 镜像用的 CANN 9.3 未在本地测到,这一点作者写在了描述里。改动 579 行、10 个文件。
第四条是 TileFoundry 的一条暗线落地。上一期报告里它还挂在「在审」,这一窗口合入:#285 让 sync 分析规划出的提交、等待与栅栏真正由 CUDA lowering 发射,并逐条反向校验实际发射与计划是否一致,不一致抛 LoweringError。语料上的结果是默认流程 14 个程序加上同步、拒绝 31 个,这 14 个的 TIR 与主线差别只有新增语句,另有 30 个与主线完全相同;单张 sm_89 上跑了 7 个能构建且带同步的程序,6 个与参考完全一致、synccheck 与 memcheck 全干净,而同样这 7 个按主线的方式 lower 有 4 个结果错误。作者在动机里给的直接症状是 sm_89 的 E4M3 GEMM 最大偏 659。这一笔 788 行新增、214 行删除、25 个文件,把「分析能证明」推进到「发射必须与证明一致」。
一、核心项目进展
1.1 主线:ThreadSync 允许跨存储域复用部分屏障,B300 上屏障占用从 12 回到 16 块/SM(#3761)
合于 11-04 00:50,新增 401 行、删除 72 行,4 个文件。背景是上一个窗口的 #3643:ThreadSync("shared") 与 ThreadSync("shared.dyn") 分开调用时,可能给互不相交的线程组分配同一个命名屏障 ID,这是一个真实的正确性问题,修法是保留更早调用已经用掉的 ID。副作用是同一组线程跨作用域也不能复用,回到 TileKernels 的 mhc_pre_big_fuse 内核上就是每块 96 线程的一次多余分配:线程 32 到 95 的两段式 cp.async 流水线在第一遍拿到 ID 3、第二遍拿到 ID 4,编译产物声明的屏障资源从 4 变 5,B300 隐藏维 1280 的屏障占用上限从 16 块/SM 掉到 12 块/SM,而寄存器、共享内存与同步指令数完全没变。本笔让 ThreadSync 接受有序的作用域列表,CUDA 流水线改成 ThreadSync(["shared", "shared.dyn"]):同步仍按原顺序规划插入,部分的屏障 ID 在两个作用域都处理完后统一分配;跨作用域共用一个 ID 之前用新分析器证明参与谓词在候选线程范围内恒成立,分支与 Bind 上下文由 ConstrVisitor 提供,线程相关谓词作为证明义务而不是假设,已证密的密集组按包围盒体积计算、跳过部分线程数枚举,未证明的组把作用域留在键里并让参与者计数入键,显式编号的屏障(含更早调用分配的 ID)继续保留,不相交线程组仍然隔离——#3643 的修复没有被回退。字符串形式的旧调用保持兼容,CuTeDSL 也走同一条 CUDA pass 流水线。
来源:https://github.com/tile-ai/tilelang/pull/3761
1.2 主线:移除遗留的 TVMScript AST 与解析器分支,一次删掉 3384 行(#3736)
合于 11-03 13:05,新增 1 行、删除 3384 行,10 个文件。被删的是 tilelang/language/ast/、tilelang/language/parser/ 与 tilelang/language/tir/entry.py——早期本地 TVMScript 解析器分叉的残留,仓内没有任何调用方:语言层通过 tilelang.language.tir.common 使用上游的 tvm.tirx.script.parser 接口,TileLang 自己的 AST 与语句覆写注册在 tilelang.language.overrides,T.prim_func 与 T.macro 由 tilelang.language.eager 导出,tilelang/language/tir/ 下只有 common、exports、ir、op 四个模块被消费。保留这个孤立的解析器分叉的风险写得很具体:一旦有人直接 import tilelang.language.parser.operation,它会在 tirx.PrimExpr 与 tirx.IterVar 上注册互相冲突的算子。本笔同时删掉 tilelang/language/tir/__init__.py 里对 prim_func 的无用再导出,以及 docs/conf.py 中针对 *language/ast* 的忽略规则(*parser* 保留,因为 tilelang/language/overrides/parser.py 还在用)。验收口径是公开导出不变,由 test_tilelang_language_dialect.py 与语言前端口测试覆盖。
来源:https://github.com/tile-ai/tilelang/pull/3736
1.3 主线:自动调优的输入供给改用当前设备,昇腾与 MPS 上的调优恢复可用(#3546)
合于 11-03 13:19,新增 153 行、删除 12 行,3 个文件,修 #3545。set_autotune_inputs 安装的冻结供给闭包从「调用点从不传入」的某个参数上解析基准设备——按供给程序的契约,实际传进来的是内核参数列表——于是在没有 CUDA 的机器上,CUDA 或 CPU 的兜底分支把每个捕获张量搬到 CPU,每一轮调优都因启动器设备不匹配而失败。现在闭包改用 get_current_device(),与默认张量供给同一策略:CUDA 上行为不变,多卡基准 worker 仍按卡拷贝;昇腾与 MPS 上捕获张量在各自当前设备供给,调优因此可用;其余后端仍退到 CPU。新增的昇腾用例放在 testing/ascend/autotune/,覆盖装饰器路径、AutoTuner 的基准供给设备、以及跨设备捕获的输入三种情形,双卡 Ascend950 上修前全败、修后三项全过;主仓 testing/python/autotune/ 12 项通过、28 项跳过,另有 3 项(eager 模式、标量输入、带输入)因该机无 CUDA 失败,作者注明留待 CI 的 CUDA runner 验证。
来源:https://github.com/tile-ai/tilelang/pull/3546
1.4 在审:条件中的标量 cast 不再压缩数据向量宽度,承接昨日 #3734(#3759)
11-03 14:09 更新,新增 148 行、删除 29 行,2 个文件。这是对昨日合入的 #3734 的收尾:#3734 在 DecoupleTypeCast 之后正确收紧了数据 cast 的宽度限制,但也让「标量条件里的 cast」开始限制数据向量宽度——在 b5cb02633 上,一个受守卫的 BF16 拷贝可能因为条件里有一个 int64 cast,从 16 个 lane 被压到 4 个,哪怕该条件在 16 个 lane 上完全一致。修法是把规划器已有的条件一致性分析复用到 cast 宽度豁免上,把 if 语句、tirx.if_then_else 与标量谓词的 Select 条件从它们的子句里分出来单独访问,条件里的 cast 不再施加数据类型宽度限制,子句里的 cast 仍然施加,并在带逐 lane 向量谓词的目标上保留条件 cast 的约束;#3579 的内存守卫行为与 #3734 修好的守卫 FP32 转 BF16 限制都保持不变。回归覆盖 8 lane 与 16 lane 的一致守卫、三种条件形式、子句内的 FP32 cast、向量谓词与舍入位分组,九个 int64 守卫用例修前失败,最终合法化与 cast 舍入两组套件共 74 项通过。
来源:https://github.com/tile-ai/tilelang/pull/3759
1.5 在审:两处用户可见错误改为明确报错(#3765 / #3767)
两笔都在 11-04 01:10 前后更新,同一个报告人提交。#3765 对应 #3653:由单个线程索引边界守卫的部分屏障,当该边界是 warp 大小的整数倍时,整 warp 的线程对条件判断一致;但通用算术证明漏掉了外层分支使用按位 & 时的 thread_binding < 448,ThreadSync 直接中止。本笔在线程维是整 warp 数时显式接受这种情形(同时覆盖边界等于或超出该维的情况),其余条件仍走既有证明路径。 #3767 对应 #3005:fragment 跨线程分布,同一 T.Parallel 体内用两个结构不同的索引读写同一 fragment(例如 buf[i] 与 buf[(i + 1) % N] 这种跨 lane 模板)无法在不做跨线程交换的前提下 lower,此前 ParallelOpNode::RecordBufferAccess 撞上 StructuralEqual 的 ICHECK,暴露成编译器内部断言;现在改成 ValueError,把两个索引表达式都写进报错并说明 fragment 不能这样访问。两笔都属于「把内部断言换成用户能读懂的诊断」,对使用者体感的影响比改动规模大。
来源:https://github.com/tile-ai/tilelang/pull/3765、https://github.com/tile-ai/tilelang/pull/3767
1.6 在审:带 padding 的共享 tile 无法整体拷贝时回退,2-CTA 的 TMA 字节按对等 leader 记账(#3738 / #3766)
两笔都在 11-03 21:31 到 21:47 之间更新。#3738 修 #3704:从比 padding 后共享 tile 更小的全局区域做 T.copy(例如 K=100 拷进 BK=112)会被 ClassifyWarpSpecializedCopy 与 SelectCopyInstForLowering 判成 TMA 整块拷贝,然后在 lowering 硬失败,报 s_extent == g_extent (112 vs. 100)。原因是选择逻辑只检查了最后一维对齐、dtype 配对与全局 stride,没有检查 padding 后的范围能否构成一次整块拷贝;本笔补上这一判据,不满足时回退到普通拷贝路径。#3766 修的是 cluster 场景下 TMA 字节计数:cta_group::2 的 TMA load 会清掉屏障地址的对等位,一对 CTA 的字节都落在偶数 CTA 的屏障上,而带 cluster 屏障的 tma_copy 原先只对 rank 0 期待 cluster_size × tile_bytes——在 (2,1,1) 与 (1,2,1) 上碰巧对,在 (4,1,1) 与 (2,2,1) 上 rank 0 只收到 2 份却期待 4 份,另有字节落到 rank 2 的屏障上。两笔都不是新功能,属于「边界条件一破就报内部错误」的那类缺陷。
来源:https://github.com/tile-ai/tilelang/pull/3738、https://github.com/tile-ai/tilelang/pull/3766
1.7 在审:B300 的 sm_103a 适配两条线,tcgen05 的 int8 选择与 CuTeDSL 的 TMA 架构门(#3770 / #3772)
B300 报出的目标架构是 sm_103a,两条在审 PR 分别处理它带来的两个卡点。#3770 对应 #3769:tcgen05.mma.kind::i8 只存在于架构专属的 sm_100a 与 sm_110a,既不在 sm_103a(B300)上,也不属于家族可移植的 sm_10xf,而 AllowTcgen5Mma 对 TargetIsSm100(arch 100 到 110)一律接受 int8/uint8 操作数,内核顺利通过 lowering 后在 ptxas 报 Feature '.kind::i8' not supported on .target 'sm_103a';本笔让 TileLang 在这些目标上不再选择 kind::i8。#3772 处理 CuTeDSL 的 TMA 架构门:B300 的 sm_103a 在进入 DSL 之前就被拒,报错列出 sm_90/90a/100a/110/120/120a/121/121a 而不含 103a;在主线 fa69f7fa 加 nvidia-cutlass-dsl 4.7.0 与 CUDA 13.3 的组合下,B300 上有六个 CuTeDSL 测试失败,包括 tma_load 的多个 JIT 用例。同批还有 #3768 修 SM100 flash attention 反向示例:T.tcgen05_gemm 现在把 mbar 定为必填的仅关键字参数且不再隐式等待,而示例仍按同步 GEMM 的写法调用,两个脚本在编译前就报 missing 1 required keyword-only argument;作者同时指出 examples/flash_attention_sm100/ 目录下没有 test_*.py,这是问题一直没被发现的原因。
来源:https://github.com/tile-ai/tilelang/pull/3770、https://github.com/tile-ai/tilelang/pull/3772、https://github.com/tile-ai/tilelang/pull/3768
二、多后端适配(昇腾 / 海光 / 摩尔线程 / 沐曦)
2.1 昇腾:3D GM 到 L1 的批量 ND2NZ 拷贝合入,一次描述符覆盖整批(#3689)
合于 11-03 17:07,新增 579 行、删除 39 行,10 个文件。此前昇腾的 GM 到 L1 拷贝一个 DMA 描述符只处理一个矩阵,拷一批要写 for b in T.serial(B) 逐批调用;本笔允许把 [B, M, N] 区域写成一次 T.copy,lowering 生成批量 ND2NZ 描述符,保留源的批与行 stride,目的侧的批间距按物理布局推导,转置批量走 DN2NZ,并复用既有的矩阵轴边界处理:超出 GM 边界的源批被跳过、目的存储保持不动,M/N 的尾部补零只作用于合法批。既有二维与单矩阵行为不变;实现只针对 AscendC 后端,且要求 L1 目的端是规范矩阵布局,PTO 支持不在本次改动内。验证环境是 Ascend950DT 加 CANN 9.2:50 个 AscendC 示例测试与 1733 个昇腾单测通过,其中 28 个 BF16/FP32 设备用例覆盖批与 M/N 尾块、整批越界、非连续源 stride、目的存储大于源、矩形转置拷贝。作者注明 CUDA 与 PTO 用例被排除,CI 镜像使用的 CANN 9.3 未在本地验证。这条与上一期合入的昇腾 950 FlashMLA 内核是同一方向的配套:注意力与 GEMM 的批量搬运都需要一次描述符覆盖整批,逐批调用在生产形状上会摊出可观的发射开销。
来源:https://github.com/tile-ai/tilelang-ascend/pull/3689
2.2 昇腾在审:TileLang-Ascend 整仓迁入主线多后端架构,2887 个文件(#1975)
11-04 00:34 更新,改动规模达到新增 607317 行、删除 97595 行、2887 个文件,是本期体量最大的一笔,也是整个组织在审 PR 里最值得跟踪的一条。目标是按提案 #1911 把 TileLang-Ascend 搬到主线的多后端架构上:原有昇腾组件移入主线对应目录的 ascend/a2a3 下(tilelang/ 对 tilelang/ascend/a2a3,src/ 对 src/ascend/a2a3),原 Ascend C 与 PTO A2A3 后端在主线独立注册,与主线的 Ascend、PTO A5 正交,pass 与 codegen 互不干扰、只复用少量基础设施;两边共用 target="ascend" 与 "pto",但 A2A3 后端附带 arch="dav-2201" 标识(由 NPU 名称自动识别)按此分流;A2A3 前端原语改用 import tilelang.ascend.a2a3.language as T,与 A5 的 import tilelang.ascend.language as T 并存,算子写法相对原仓不变;pass 与 codegen 直接用从 TileLang-Ascend 迁来的版本以保持原编译流程,执行后端复用主线 cython 并补 workspace 索引等昇腾独有适配,examples 与 testing 分别落到 examples/ascend_a2a3 与 testing/ascend_a2a3,CI 暂用原仓的 ci_cd.yml 并确认适配新路径后仍能跑通。作者在描述里把待办写得很实在:三个目录下文件与原文件的差异尚待收敛、静态检查与格式化 CI 待恢复、文档与 Agent 相关待迁移,并注明「涉及文件较多,Agent 批量操作难免有疏漏」,请求协助检视。这种把迁移中的不确定直接摊在 PR 描述里的做法,比等一个「已完成」的结论更利于评审。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1975
2.3 昇腾在审:copy_pipe_to_l1 的同步标注按执行引擎拆分,修 A2/A3 上 10/10 复现的错结果(#1981)
11-03 20:36 更新,对应 #1780。根因写得很清楚:A2/A3 的 PTO 跨核 V 到 C 通信经 GM 中转,三段搬运分属三个引擎——AIV 的 MTE3 做 UB 到 GM 的发送、AIC 的 MTE2 做 GM 到 L1 的接收(TPOP 内部 TLOAD)、AIC 的 MTE1 做 L1 到 L0A;而 copy_pipe_to_l1 在 operation_config.h 里被标成 PIPE_MTE3,自动同步 pass 因此在它和 copy_l1_to_l0a 之间生成 set_flag(PIPE_MTE3, PIPE_MTE1, …) 加对应的 wait_flag——等的却是 AIC 本核的 MTE3,而 AIC 的 MTE3 队列为空、事件立刻置位,MTE1 读 L1 于是可能在 MTE2 写完之前放行,L0A 拿到旧数据、MMA 结果错误。错误标注随 #1238 引入,同一张表里 copy_pipe_to_ub 与 copy_l0c_to_pipe 都按执行侧引擎标注,只有这一条沿用了发送端视角。复现证据两条:issue 作者把生成源码里的 PIPE_MTE3 改成 PIPE_MTE2 后 10/10 全部恢复精确,本地在最新主干上 10/10 次运行错误、单次最多 16384/16384 个元素不匹配。修法按 issue 的建议拆名字而不是直接替换:copy_pipe_to_l1 标注改为 PIPE_MTE2(A2/A3 的 GM 到 L1 加载),新增 copy_pipe_to_l1_bind 标注 PIPE_MTE1(A5 直连路径上等待 L1 就绪并绑定地址),沿用 copy_pipe_to_ub 与 copy_pipe_to_ub_V 的既有先例,避免修好 A2/A3 却让 A5 出错。同源问题也覆盖了 WAR 方向:MTE1 仍在读时,MTE2 的 TLOAD 可能覆写同一 L1 缓冲。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1981
2.4 昇腾在审:首 Parallel 轴广播走标量读、Chunk O 前向迁移、MHA 示例(#1979 / #1866 / #1879)
三笔都在 11-03 下午集中更新,属于昇腾算子里程碑式的一批。#1979 对应 #1885:失败族可以归成一个模式——由第一根 T.Parallel 轴(或外层 serial 变量)索引的一维 UB 操作数出现在二元表达式左侧,issue 的收缩模式 dk_i = sum_j v[j] * A[j, i] 正是这一形态,21 点二分矩阵里 5 个 FAIL 全部属于这一族,与归约方向、广播写法无关;根因是 AscendLowerParallelToVector 把二维并行嵌套降成行循环时,入口 HandleSimpleCase 只对右侧操作数做了标量访问检测,左侧走向量路径后撞上 UB 对齐陷阱;本笔让首轴的广播操作数也走标量读。#1866 迁移 Chunk O 前向(基于 chunk 的门控线性注意力 O 前向,GDN)到昇腾 NPU,用开发者模式加四项 pass 配置(自动 C/V 合并、自动 C/V 同步、自动同步、内存规划),单个 @tilelang.jit 内核加固定 20 核的 grid-stride 循环,做到片上融合、零 GM workspace、零宿主 ATen 算子。#1879 新增 MHA 前向示例:BSHD 布局、Expert 模式的三分派 MMA 流水,按 wrapper、共享层与三形态模块拆成 5 个 Python 文件,wrapper 不含任何 T.* 代码、只做元数据校验与 JIT 分发;is_causal 采用右下角对齐掩码,全掩码行经 sumexp + 1e-30 保护输出 0;cann-bench 20 个用例精度全过,对比 torch-npu 基线平均加速比 0.97 倍、综合得分 73.13(本地 910B2 实测口径)。0.97 倍意味着还没有净收益,但作为示例它把「昇腾上怎么写一个完整注意力算子」的骨架给出来了。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1979、https://github.com/tile-ai/tilelang-ascend/pull/1866、https://github.com/tile-ai/tilelang-ascend/pull/1879
2.5 mlir-ascend:Codex 工作流与任务交接合入 harness 分支(#208)
合于 11-03 17:50,新增 2715 行、删除 39 行,66 个文件,合入目标分支是 harness 而不是 main(主线最新提交仍停在 10-31 的 ci: remove internal cache-service references),这一点在引用时不要写成「合入主线」。内容是给算子 harness 加第一套 Codex 工作流:项目初始化时用迭代提问、把已有算子仓接入、以及从持久化上下文里恢复开发、调优或分析;三套技能共享 .workspace/ 下可编辑的 Markdown 记录,任务交接区分「当前候选」与「已验证的最佳候选」,保留失败与回退过的实验,让子代理拿到聚焦上下文而主代理维护共享状态;本地安装器负责注册技能并把项目条目追加进 AGENTS.md(保留既有内容、忽略 .workspace/、把 Git 项目解析到 worktree 根)。README、依赖、复用说明与路线图描述的都是 TileLang 先行、昇腾为范围的设定;工作流参考了 GSD 并保留来源修订与 MIT 声明,NPU 执行适配器与 npm 分发仍在待办里。这一笔与主仓、昇腾仓里的 Agent 相关目录是同一件事的三个面:把「智能体写算子」的流程沉淀成仓库内的可复用资产。
来源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
2.6 海光:Distributed 在审 PR 有更新,主分支无新提交(#19)
tilelang-hygon 在窗口内有推送记录(11-03 11:06 北京时间),但主分支没有新提交,最新一笔仍是 11-02 的 #32(把 gfx948 按 gfx946 兼容处理,先让既有指令可测)。推送来自在审的 #19(标题只有 Distributed,描述为空),11-03 11:06 更新,从时间与标题看是在推进分布式相关的适配,但 PR 描述没有内容,本报告不对实现细节作推断。该仓近期的常规节奏是每两三天一笔 HCU(海光计算单元)侧修正:10-30 的 #31 改 MLS 边界与持久 WASP 同步、#30 补普通拷贝的 LDS swizzle、10-27 的 #29 做 scale 缓冲的 bank 划分。
来源:https://github.com/tile-ai/tilelang-hygon/pull/19
2.7 沐曦、摩尔线程窗口内无推送(11-03)
tilelang-musa(摩尔线程)最后推送停在 09-30 09:50,tilelang-metax(沐曦)停在 10-26 07:08,两仓本窗口都没有提交,也没有在审 PR 更新。本期的后端动态因此集中在昇腾与海光两侧:合入 8 笔里有 1 笔属于昇腾(#3689)、1 笔属于 mlir-ascend(#208,合入 harness 分支),在审的 48 笔更新中昇腾仓占 15 笔、主仓占 23 笔。
来源:https://github.com/tile-ai/tilelang-musa、https://github.com/tile-ai/tilelang-metax
三、生态与采用方(TileOPs / TileFoundry / TileKernels / FlashQLA)
3.1 TileFoundry:同步计划落成 CUDA lowering 并反向校验,合入主线(#285)
合于 11-04 00:19,新增 788 行、删除 214 行,25 个文件。sync 分析规划出共享内存访问所需的提交、等待与栅栏,但 finalize 此前一条都不发射,经共享内存暂存的调度程序因此存在竞争,卡上表现为 sm_89 的 E4M3 GEMM 相对 torch 最大偏 659。本笔把分析步骤挂到它们所前导的 HIR 节点上,以 SyncMetadata 按标签索引:节点自身语句的标签为空,InsertSlice 加入的拷贝用 base 与 update,循环携带拷贝用 ("entry", 槽位) 与 ("backedge", 槽位);核心 lowering 在 lowering 某节点的语句之前调用新钩子 before(owner, ctx, tag),由 CudaLowering 在那里发射该节点的步骤;schedule finalize 负责跑 sync 分析。契约写进了四份规范文档,其中分析规范 §1.2.5 描述 SyncMetadata、标签、规则以及分析仍会拒绝的情形。风险面作者也标了:语料里所有在主线能 finalize 的程序仍然能 finalize,只多出语句,例外是 hir/staged_dynamic_inner——它的 lane 在一个「行程数等于 lane」的循环里读写同一块共享 tile,没有栅栏能把它们汇合;另外触碰共享内存的保留调用按「未知字节的一次访问」处理,两个这样的调用在没有线程网格的区域里需要 TIR 无法放置的栅栏,分析会拒绝它们,即便字节互不相交;CopyAsyncTensor、CopyReduceAsyncTensor 与 wgmma 还没有发射器。
来源:https://github.com/tile-ai/TileFoundry/pull/285
3.2 TileFoundry 在审:wgmma 同步计划、FSDP 与 HSDP 训练语义进 HIR 求值器(#286 / #284)
#286 在 11-04 02:28 更新,新增 1907 行、删除 310 行,52 个文件,是 #285 的直接续笔:WGMMA 在发射之后仍可能继续读共享内存与寄存器,CUDA 同步计划需要把后续访问排在这些作用之后;本笔为每种操作数效应声明完成语义,把 WGMMA 的栅栏、提交与等待挂到需要它们的 HIR 节点上,lowering 在指令的发射点消费这些步骤,若后续作用域放不下集合等待就在发射作用域里完成该操作;同时修正寄存器 Reshard 别名,让元素在持有副本的每个线程上留在同一个寄存器里,计划用这一存储身份跟踪寄存器依赖。作者明确写了这条风险:矩阵指令发射器仍不可用,本笔不建立 WGMMA 的执行正确性或性能,TMA 完成语义仍单独处理。#284 在 11-03 14:34 更新,处理训练语义:此前的求值器把 reshard 当全局张量,漏掉了不同参与者的贡献;本笔求值参与者本地值、GPU 坐标、调用与循环、gather/scatter 以及按轴的 Partial(sum) 解析,并让 Reshard 的别名证明带所有权语义,扩展 #260 与 #170 的结构化生命周期;配套加了一个受 modded-nanogpt 启发的两 block GQA 研究,含显式前向反向、动量 SGD 与独立的 PyTorch 参考,覆盖 FSDP 与 2×2、4×4 两种 HSDP 驻留策略,报告集合通信组、载荷与存储。验收数字是重构后核心套件 775 项、后端套件 778 项(两卡 sm90 上 13 项跳过),十组 HIR 配置在 CPU 上通过,两种策略在四卡 sm90 上与真实 2×2 PyTorch HSDP 跑两个更新步、本地梯度、参数与动量全部对齐。
来源:https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/pull/284
3.3 TileOPs:方差合并从第一个部分量起算,常量输入的方差不再出 NaN(#2821)
合于 11-03 18:04,新增 43 行、删除 4 行,3 个文件。缺陷是方差合并从一个空的累加器起步:把一个有限的大均值合进空累加器时,平方差在乘以零权重之前就已经溢出,常量输入的方差因此变成 NaN;受影响的是在保留轴前后各有一段归约时对 var、std 与 var_mean 的调用。修法是让合并从第一个真实的部分量初始化,后续部分量沿用既有的 Chan 递推;行归约、分发、输出布局与公开 API 都不变。验收:三笔新增回归在原始内核上失败、修复后通过,候选过 155 项回归测试与 24 项额外数值检查,tests/ops/test_reduce.py 从 155 增到 158;三项自动调优测试被排除,原因是先前的基线尝试超过了 60 秒测试上限。这类「空累加器 + 大偏移」的数值坑在归约族里并不罕见,值得记的是它只在保留轴前后同时归约的形状上才暴露,普通整轴归约测不到。
来源:https://github.com/tile-ai/TileOPs/pull/2821
3.4 TileOPs 在审:GEMM 加可选 bias、FP8 分组 GEMM 支持连续布局(#2823 / #2824)
#2823 给 GEMMFwdOp.forward 加了可选 bias[N]:bias 加进 float32 的累加结果、输出只做一次舍入,与 torch.nn.functional.linear 的口径一致;每种稠密结构都在自己的尾声里加 bias——warp 专用(TMA 与行存储)、coop2(含 176/192 宽与窄 tile)、coop2_swap、ping-pong、coop2s、simple、swap_ab、GEMV 各带与 cp.async 内核,split-K 路径在归约里只加一次。#2824 让 MoEGroupedGEMMFP8FwdOp 除掩码分块外也接受连续专家行(紧凑打包或对齐打包,带段末偏移 physical_psum 或逐行专家 id),语义对齐 deep_gemm.m_grouped_fp8_gemm_nt_contiguous;SM90 内核按段末遍历每个专家的运行段,紧凑的逐行 id 在生产者与消费者拆分之前用二分转成段末,对齐的逐行 id 按行 tile 逐块走、从首行读该 tile 的专家,紧凑专家最后一块不足 block_m 行时逐行存储,绝不写到下一个专家的行上,SM89 拒绝连续布局;掩码调用仍构建与之前相同的程序,35 个掩码构建的 TIR 与主线字节一致。基准侧新增三行 Qwen3-30B-A3B 的 EP8 形状,用的是从真实路由测出的组大小(该模型 128 专家、top-8,取 48 层里偏斜居中的第 44 层,选 EP8 中分到行数最多的那个 rank 的 16 个专家),prefill 行路由 4096 个 token(gate/up 与 down),decode 行取 128 条序列的末 token。
来源:https://github.com/tile-ai/TileOPs/pull/2823、https://github.com/tile-ai/TileOPs/pull/2824
3.5 TileOPs 在审:W4A16 与稠密 GEMM 的选择器改成单一代价模型重拟合(#2825 / #2826)
两笔同属 #2133 的子任务 2,分别在 11-03 20:54 与 20:54 更新,一个管 W4A16、一个管稠密,且后者只改文档。#2825 把 _select_config 改成在 _legal_configs 上取 _config_cost 的 argmin,删掉 tile 优先、最宽 K 的并列窗口与 joint_switch_margin 三条例外规则;_config_cost 新增 stage_fill 项——更深的环能藏住每个 K tile 的更多延迟,但填充本身更慢,没有这一项时在 K=4096 上模型会把 256 宽四阶段排到 512 宽两阶段之前,而后者实测最多快 11%;mma_big、iter_ws 与 launch_ms 被删掉,因为数据不足以钉住它们。常数在基准的 SM 时钟下按忙时重拟合,拟合形状取清单之外的形状、清单本身留作留出检验,_Calibration 的文档串写明扫描、拟合与检查三步,从中性常数重跑能复现出厂值。内核侧还有一处实打实的省内存:split-K 的 CTA 只暂存自己那一片的 scale 与 zero 组,而不是整个 K,切片对齐时用向量化加载;_smem_bytes 与 _legal_configs 按每个 split 计共享内存,而主线上每个 split 的 CTA 都暂存了整个 K 的组。另一个细节是 _ConfigSpace.smem_bytes 放在比可选上限低 1 KB 的位置——正好压在上限时,17 个被扫描的配置启动失败,因为分配器会把动态共享内存向上取整、而屏障是静态的。#2826 则把稠密选择器的常数如何重拟合写进文档串(扫描、拟合、检查三步,含软退火与留出检验),不改代码。
来源:https://github.com/tile-ai/TileOPs/pull/2825、https://github.com/tile-ai/TileOPs/pull/2826
3.6 TileOPs 在审:短 varlen 按 KV 组打包、L2 flush 改为只读(#2819 / #2822)
#2819 在 11-03 23:37 更新,是 #2496 的一部分,给 SM90 上的 GQAVarlenFwdOp 加了一个打包的 warp 专用程序:当一个请求用它的 KV 头分组按位填满一个 tile 更省 tile 时,键与值只按 KV 头读一次;打包内核的拒绝逻辑用一个代价模型衡量调用的批、打包后的 query 总数与组大小,只有当形状保证的节省超过实测的 9% 单头 tile 开销时才接手;宽窄两类 item 在同一次发射里交替领取,不做宿主回读、不额外加同步;当设备上数出的未拆分 item 少于 SM 数时,单 tile 的打包请求把键扫描拆成多段、经 fp32 工作区合并。窗口侧还有两笔实测收益:窗口行省 1% 到 9%,交错 RoPE 行省 3% 到 6%。同时修了主线上三个问题:带 sink 且 score scale 大于约 2.4 时无可见键的行返回 NaN(单头 warp 专用与 FP8 warp 专用两个程序)、后备的 query 旋转在请求首个键之前的行上先读 RoPE 表、以及被拒的负控制用例让调用方帧留在引用环里导致一个测试的 CUDA graph 在另一个测试的捕获期间被销毁。测试节点从 173 增到 197。#2822 在 11-03 19:01 更新,把 bench_kernel 的 L2 flush 从「对 int8 缓冲做 zero_()」改成「读一个私有 int32 缓冲,大小是报告 L2 的两倍」,缓冲由固定种子的 CPU 生成器一次抽出;CUDA 上用 Triton 内核做这次读(每元素一次加载、每 16 KB 一块一次 atomic_max),不用 torch 自己的归约是因为把它当 flush 会拖慢下一次 Marlin w4a16 调用 1.3 到 1.5 微秒。收益在 H200、1500 MHz、CUPTI 内核时间(3 次运行各 30 次调用的中位数)上:2048×11008 f32 的 Sum 从 27.81 降到 21.38 微秒,1024×11008 f32 的 Ge(torch-compile)从 31.31 降到 25.67,1024×4096 f32 从 6.30 降到 5.47;旧 flush 会留下约 120 MB 脏行,其回写落进下一次计时调用,新 flush 自身只写 71 到 445 个 L2 扇区(2 到 14 KB)。顺带修了一个主线 bug:L2 大小未知时日志写 256 MB 却分配 512 MB。
来源:https://github.com/tile-ai/TileOPs/pull/2819、https://github.com/tile-ai/TileOPs/pull/2822
3.7 采用方:FlashQLA 与 TileKernels 无推送,但 #3761 修的是 TileKernels 的 fused 内核(11-04)
QwenLM/FlashQLA 最后推送停在 10-30 03:02,deepseek-ai/TileKernels 停在 10-14 05:49,两仓在本窗口都没有提交,采用方侧没有新代码。不过本期的主线头条与采用方有直接关系:被 #3761 修掉的那笔性能回退,触发点正是 TileKernels 的 mhc_pre_big_fuse 内核——同一个内核在 B300、隐藏维 1280 上丢掉了 25% 的屏障占用预算。这说明采用方的内核已经成为编译器侧的回归探针:TileKernels 不推送代码的这几周,它的内核仍通过主线的改动在被动地「被测试」。对跟踪生态的人来说,这是一个值得留意的反向信号——采用方静默期里,编译器的改动仍然可能在它们的形状上出问题。
来源:https://github.com/QwenLM/FlashQLA、https://github.com/deepseek-ai/TileKernels、https://github.com/tile-ai/tilelang/pull/3761
四、社区、教程与活动
4.1 单日 7 条新 issue:一条生产级 CUDA graph 问题,四条来自同一外部报告人(11-03)
窗口内主仓新增 7 条 issue(编号 #3762 到 #3775,另有一条昇腾仓的自动日测报告)。技术含量最高的是 #3775,报告人是带生产负载的团队:环境为从 pip 安装的 tilelang 0.1.12、Python 3.12.3、两张 H200 加 NVLink、驱动 580.178.04、CUDA 13 容器、torch 2.14.1+cu130,服务侧是单进程 TP2(每个设备一个工作线程、各自一个 CUDA 上下文)的 sglang,模型为 GLM 系 MoE 加 DSA 注意力,路径上的 TileLang 内核是 mhc_pre_big_fuse_with_norm、mhc_post 与 DSA prefill 内核。把执行后端强制成 TILELANG_EXECUTION_BACKEND=nvrtc(默认 auto 会解析到 tvm_ffi)后,tilelang 的 JIT 内核在 CUDA graph 捕获期间运行会让捕获损坏,而且损坏出现在 tilelang 发射之后的下游无关内核上,三次字节一致的配置复现出三个不同的首个失败点:一次是 cudaError700(无效的 TMA 描述符句柄)在 TMA 描述符初始化处抛出、之后在无关发射点回响,一次是 cublasGemmEx 报 CUBLAS_STATUS_EXECUTION_FAILED(纯 BF16 进出),一次是宿主框架 JIT 出来的逐元素内核里的非法内存访问;三次都发生在 decode 的 CUDA graph 捕获期间(批大小桶 8),都在服务启动后约 100 秒,且 /proc/environ 与干净对照的差别只有那一个环境变量。这属于最难定位的一类:现象随机、失败点漂移、诱因只有一个环境变量,报告人把三次运行的对照条件写全了。
另四条来自同一位外部报告人,环境统一为主分支或 0.1.15 pip 包、WSL 下 Linux、Python 3.13.5、RTX 4060 Laptop(sm_89)、驱动 566.26、CUDA 12.6 运行时,且都附了可复现脚本与完整输出:#3774 是 16×16 的共享内存置换只做了行位移、丢了列偏移(256 个元素错 223 个,且与「丢掉列位移」的结果完全一致),而同样写法的 32×32 是正确的;#3773 是均匀 while 循环里做三次共享内存拷贝加 GEMM 会得到不稳定结果(串行循环对照 100 次 0 错,while 版本 100 次错 21 次、出现 10 种不同输出);#3763 是嵌套整数向下取整 (x // -2) // 4 在 int16、64 元素、32 线程下报 Let cannot bind the same var to two different values。这三条加上 #3775,构成了「消费级卡 + WSL + pip 安装」这一组合的集中报障面。剩下的两条:#3769 是 B300 上 tcgen05 的 int8 GEMM 被 ptxas 拒绝(已有在审的 #3770 修),#3762 是昇腾 950 上 T.simd.vcvt 在 sat=False 时仍然饱和,详见下条。
来源:https://github.com/tile-ai/tilelang/issues/3775、https://github.com/tile-ai/tilelang/issues/3774、https://github.com/tile-ai/tilelang/issues/3773
4.2 昇腾 950 的 vcvt 饱和语义与 mlir-ascend 日测(#3762 / #1980)
#3762 由昇腾侧使用者提交,环境为 Ascend950PR(NpuArch 3510 / dav-3510)、CANN 9.2.0-beta.2、驱动 25.7.rc1、torch_npu 2.12.0、从源码装的 TileLang(10-21 的 e84ce3de)。问题是 T.simd.vcvt(src, "uint8", sat=False) 对 int32 向量静默忽略 sat=False,永远饱和:负值变 0、超过 255 的变 255,既不报错也不告警,而调用方以为自己在用回绕语义。文档串对这种情形写得很明确——sat=False 时越界值回绕——期望值是 -1 变 255、-128 变 128、300 变 44,实测是 -1 变 0、-128 变 0、300 变 255。影响落在 int8 量化尾声那类 f32 转 int8 的路径上:一旦出现负的中间值,语义就从回绕变成截断到零。#1980 是昇腾仓的自动日测,11-03 09:35 那轮 3030 项全部通过、通过率 100%,工作流运行编号 41553942797。这两个数字放在一起看有意思:昇腾仓的日测是满绿的,而日测覆盖不到的语义差异(比如 sat 参数的实际行为)仍然要靠人工报告。
来源:https://github.com/tile-ai/tilelang/issues/3762、https://github.com/tile-ai/tilelang-ascend/issues/1980
4.3 文档站三次更新,VS Code 扩展改名 vsix(11-03)
文档站 tilelang.github.io 在窗口内有三次 Update docs 提交(11-03 13:11、13:28、11-04 00:59 北京时间),属于随主线合入自动构建的常规更新,内容对应本窗口合入的几笔——#line 分级映射、ThreadSync 作用域列表与昇腾批量拷贝的文档相应位置。TileOPs.github.io 本窗口没有推送,最新仍是 11-01 的「跟随 TileOPs 主线:重叠与对齐保证、构建说明」。VS Code 扩展仓在 11-04 02:19 有一笔「更新文档并改名 vsix」,接在 11-02 的「把 CUDA 调试对齐到 TileLang 行映射」之后:上一期主线合入的逐行锚定改变了生成代码的行号来源,扩展侧的调试配置需要跟着改,这一笔把产物名与文档一并调整。
来源:https://github.com/tile-ai/tilelang.github.io、https://github.com/tile-ai/tilelang-vscode-edition
4.4 新闻侧:三通道无窗口内命中,外媒的昇腾训练报道口径不一(11-03)
Google News 的中英文组合查询、Hacker News 与 arXiv 三个通道在本窗口都没有 TileLang 相关的实质命中:候选清单 15 条里 12 条来自 GitHub 仓库发现,2 条是与主题无关的 Show HN(侦探游戏与硬件设计提示词工具),1 条是 Wccftech 关于 DeepSeek 与昇腾的报道。arXiv 上 all:tilelang 的最新命中仍是 2026-07-24 的 TileSight(以 tile 为中心的 GPU 性能分析模型),与 TileLang 的联合查询在窗口内无新论文。唯一值得记的行业侧消息有口径冲突:Wccftech 称 DeepSeek 将在昇腾 950DT 构成的 Atlas 950 SuperPoD 上从零训练 V5,而同期 The Information 的报道是「最早 2026 年底开始在 950DT 上做预训练」,Bloomberg 与金融时报的早期报道则称 950DT 的 16 万片订单主要用于推理、此前在昇腾上训练受挫后回退到 Nvidia。三种说法对「是否以及何时用于训练」并不一致,且均未获华为或 DeepSeek 官方确认。把它放在这里只作背景:不论训练是否落地,昇腾侧的编译器与算子投入(本期 #3689、#1975、#1981)都是在为这条路径铺地基。
来源:https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods、https://arxiv.org/
五、趋势观察
第一,性能回退开始被当成一等公民,而且修法在方法论上比结论更值得记。#3643 为修正确性引入的 ID 保留规则,本身是必要且正确的;#3761 没有把它回退,而是让保留与复用在不同层级上并存——跨作用域共用一个 ID 之前先证明参与谓词恒成立,证明不了就退回带作用域的计数。配合作者那个「只改声明的屏障资源数」的隔离实验,这条修复链把「正确性」与「资源声明」两件事拆开了:后者是可以单独复现、单独验收的变量。这种拆分能力,恰恰是编译器项目从「能跑」走向「可维护」的标志。
第二,昇腾的工程化在两条腿上同时推进。一条是架构:#1975 用 2887 个文件、60 万行新增把整个 TileLang-Ascend 搬进主线多后端框架,靠 arch="dav-2201" 这样的标识做后端分流,目标是让 A2A3 与 A5 在同一个仓里正交共存;另一条是流水线:mlir-ascend 的 #208 把 Codex 的初始化、接入、恢复三套工作流和任务交接写进 .workspace/,主仓与昇腾仓也各自有 Agent 目录。把「迁移仓」与「用智能体迁移仓」放在同一天看,能看出昇腾侧的瓶颈已经从「缺算子」转向「缺一致的工程骨架与可复用的迁移流程」。
第三,算子库的自证文化在加深。TileOPs 本期在审的几笔,几乎都把「证据」写在 PR 里:真实路由测出的 Qwen3-30B-A3B EP8 组大小(取 48 层里偏斜居中的第 44 层)、35 个掩码构建的 TIR 与主线字节一致、L2 flush 改动附 46 行实现的对照表并说明旧 flush 留下 120 MB 脏行、W4A16 的代价模型明确写出「哪些项被删掉是因为数据钉不住」。TileFoundry 走的是另一条路——把同步做成可证明的分析,再让 lowering 反向校验,分析拒绝的情形直接列出来(TMA 与 wgmma 的完成语义尚未声明、跨线程边界不等的循环)。两条路合起来说明同一个取向:宁可承认「没证明」,也不把「跑通了」当成「对了」。
第四,风险面在往安装形态上集中。一天 7 条新 issue 里,四条来自同一位使用 pip 包、跑 WSL、用消费级 sm_89 卡的外部报告人,全部附了可复现脚本;#3775 则是生产环境里「pip 包 + 多 CUDA 上下文 + nvrtc 后端 + CUDA graph 捕获」的组合。这些不是编译器的深层缺陷,却决定了新用户第一次尝试时的体验。对照昇腾仓日测 3030 项全绿而 sat=False 的语义偏差仍靠人工报告,可以看出测试覆盖的形状与真实使用形状之间还有距离。
最后一点观察:本期新闻侧依旧空白,15 条候选里 12 条是仓库发现。TileLang 的对外声量仍然绑定在「DeepSeek 与昇腾」这类事件上,而日常进展只发生在提交里。对跟踪这个项目的人来说,这意味着日报的主体只能是 PR 与 issue,而不是报道——这已经是连续多期的稳定形态。
附、素材与核查说明
信源核查表
| 信源 | 核查结果 |
|---|---|
| tile-ai 组织(29 仓库) | 窗口内推送 9 个仓库 |
| 主仓 tilelang | 主分支 5 笔合入;新开 issue 6 条;在审更新 23 笔;文档站 3 笔自动构建 |
| TileOPs | 主分支 1 笔合入;在审更新 7 笔 |
| TileFoundry | 主分支 1 笔合入;在审更新 2 笔 |
| tilelang-ascend | 默认分支 ascendc_pto 窗口内无新提交;在审更新 15 笔;日测 3030 项全过 |
| tilelang-mlir-ascend | 1 笔合入 harness 分支(非主线) |
| tilelang-hygon | 主分支无新提交;在审 #19 有更新 |
| tilelang-vscode-edition | 1 笔合入(文档与 vsix 改名) |
| tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / tilelang-puzzles | 窗口内无推送 |
| 发布页(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) | 窗口内无新版本,tilelang 最新仍为 09-30 的 v0.1.15 |
| FlashQLA / TileKernels | 窗口内无推送 |
| Google News(中英文六组查询)/ Hacker News / arXiv | 无窗口内实质命中 |
完整信源清单
- [1] tilelang ThreadSync 跨存储域复用部分屏障(#3761)— https://github.com/tile-ai/tilelang/pull/3761
- [2] tilelang 移除遗留 TVMScript AST 与解析器(#3736)— https://github.com/tile-ai/tilelang/pull/3736
- [3] tilelang 自动调优输入供给设备(#3546)— https://github.com/tile-ai/tilelang/pull/3546
- [4] tilelang 条件 cast 保留向量宽度(#3759)— https://github.com/tile-ai/tilelang/pull/3759
- [5] tilelang warp 对齐线程边界的部分屏障(#3765)— https://github.com/tile-ai/tilelang/pull/3765
- [6] tilelang 多索引 fragment 访问报错(#3767)— https://github.com/tile-ai/tilelang/pull/3767
- [7] tilelang padded 共享范围无法整块拷贝时回退(#3738)— https://github.com/tile-ai/tilelang/pull/3738
- [8] tilelang 2-CTA TMA 字节按对等 leader 记账(#3766)— https://github.com/tile-ai/tilelang/pull/3766
- [9] tilelang tcgen05 kind::i8 不用于无此特性的目标(#3770)— https://github.com/tile-ai/tilelang/pull/3770
- [10] tilelang CuTeDSL 接受 sm_103a 与 sm_107a(#3772)— https://github.com/tile-ai/tilelang/pull/3772
- [11] tilelang SM100 flash attention 反向示例(#3768)— https://github.com/tile-ai/tilelang/pull/3768
- [12] tilelang CuTeDSL 示例缺口补齐(#3685)— https://github.com/tile-ai/tilelang/pull/3685
- [13] tilelang nvrtc 后端损坏 CUDA graph 捕获(#3775)— https://github.com/tile-ai/tilelang/issues/3775
- [14] tilelang 共享内存置换丢列偏移(#3774)— https://github.com/tile-ai/tilelang/issues/3774
- [15] tilelang while 循环内共享内存 GEMM 不稳定(#3773)— https://github.com/tile-ai/tilelang/issues/3773
- [16] tilelang 嵌套 int16 向下取整 Let 冲突(#3763)— https://github.com/tile-ai/tilelang/issues/3763
- [17] tilelang B300 上 tcgen05 int8 被 ptxas 拒绝(#3769)— https://github.com/tile-ai/tilelang/issues/3769
- [18] tilelang 昇腾 950 上 vcvt 忽略 sat=False(#3762)— https://github.com/tile-ai/tilelang/issues/3762
- [19] tilelang-ascend 3D GM 到 L1 批量 ND2NZ 拷贝(#3689)— https://github.com/tile-ai/tilelang-ascend/pull/3689
- [20] tilelang-ascend 整仓迁入主线多后端架构(#1975)— https://github.com/tile-ai/tilelang-ascend/pull/1975
- [21] tilelang-ascend copy_pipe_to_l1 同步标注拆分(#1981)— https://github.com/tile-ai/tilelang-ascend/pull/1981
- [22] tilelang-ascend 首 Parallel 轴广播走标量读(#1979)— https://github.com/tile-ai/tilelang-ascend/pull/1979
- [23] tilelang-ascend Chunk O 前向迁移(#1866)— https://github.com/tile-ai/tilelang-ascend/pull/1866
- [24] tilelang-ascend MHA 前向示例(#1879)— https://github.com/tile-ai/tilelang-ascend/pull/1879
- [25] tilelang-ascend 每日测试 3030 项全过(#1980)— https://github.com/tile-ai/tilelang-ascend/issues/1980
- [26] tilelang-mlir-ascend Codex 工作流与任务交接(#208)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
- [27] tilelang-hygon 分布式在审 PR(#19)— https://github.com/tile-ai/tilelang-hygon/pull/19
- [28] TileFoundry CUDA lowering 发射同步计划(#285)— https://github.com/tile-ai/TileFoundry/pull/285
- [29] TileFoundry wgmma 同步计划(#286)— https://github.com/tile-ai/TileFoundry/pull/286
- [30] TileFoundry FSDP 与 HSDP 训练语义(#284)— https://github.com/tile-ai/TileFoundry/pull/284
- [31] TileOPs 方差合并从首个部分量起算(#2821)— https://github.com/tile-ai/TileOPs/pull/2821
- [32] TileOPs GEMM 可选 bias(#2823)— https://github.com/tile-ai/TileOPs/pull/2823
- [33] TileOPs FP8 分组 GEMM 连续布局(#2824)— https://github.com/tile-ai/TileOPs/pull/2824
- [34] TileOPs W4A16 选择器代价模型(#2825)— https://github.com/tile-ai/TileOPs/pull/2825
- [35] TileOPs 稠密 GEMM 选择器重拟合文档(#2826)— https://github.com/tile-ai/TileOPs/pull/2826
- [36] TileOPs 短 varlen 按 KV 组打包(#2819)— https://github.com/tile-ai/TileOPs/pull/2819
- [37] TileOPs L2 flush 改为只读(#2822)— https://github.com/tile-ai/TileOPs/pull/2822
- [38] Wccftech 关于 DeepSeek 与昇腾 950DT 的报道 — https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods