调研窗口：2026-11-04 07:00 ~ 2026-11-05 07:00（北京时间，24 小时；上一期截点为 11-04 07:00）。

信源：GitHub tile-ai 组织逐仓核查（29 仓库，窗口内推送 8 个）+ 合入与在审真值检索（合入 12 笔、在审更新 42 笔）+ 新增 issue 4 条 + 主分支逐笔提交核查（tilelang 3、TileOPs 7、TileFoundry 1、tilelang-ascend 1、mlir-ascend 2）+ 非默认分支核查（tilelang-hygon 的 feat/hcu-scaled-cast 分支 4 笔）+ 发布页与 tag 核查（5 仓，窗口内无新版本）+ 采用方两仓 + Google News 中英文四组查询（走代理）+ Hacker News + arXiv。

本期索引

  • 今日重点：一笔 10-14 合入的昇腾 PTO 框架改动（#3512）进入集中清账期，一天冒出五个互不依赖的补丁；TileOPs 单日合入 7 笔性能改动，创节后最密的一天；主仓三笔合入全部与「度量口径」有关，L2 刷写从写零改成只读
  • 一、核心项目进展
    • 1.1 主线：CuTeDSL 的 TMA 架构门接受 sm_103a 与 sm_107a，B300 与 gr100 上六个用例恢复（#3772）（11-04）
    • 1.2 主线：L2 刷写由写零改为只读归约，do_bench 去掉 fast_flush 参数（#3534）（11-04）
    • 1.3 主线：AscendC CI 的自动调优供给用例显式同步两台设备，异步拷贝错误不再伪装成 arange 失败（#3784）（11-04）
    • 1.4 在审：ThreadSync 的 warp 一致性证明补上由 warp 索引拼出的组合谓词（#3787）（11-04）
    • 1.5 在审：T.copy 目的区域越界改为编译期报错，全复制 fragment 的包含证明补线程归属校验（#3786 / #3788）（11-04）
    • 1.6 在审：SM120 的 FP8 GEMM 换用 block-scaled MMA，RTX 5090 上 8192 立方从 463 到 727 TFLOPS（#3790）（11-04）
    • 1.7 在审：Profiler 与自动调优的口径校验成组提交，六笔同日出（#3797 / #3798 / #3799 / #3716 / #3717 / #3718）（11-04）
  • 二、多后端适配
    • 2.1 昇腾 PTO：一笔架构改动引出四类回归，五个互不依赖的补丁同日提交（#3792 / #3793 / #3794 / #3795 / #3796）（11-04）
    • 2.2 昇腾：Cube 与 Vector 的核心比可配置，不再固定 1:2（#3783）（11-04）
    • 2.3 昇腾：AutoSchedule 的周期下界收紧，Z3 求解时间降 10.04%（#3776）（11-04）
    • 2.4 昇腾：AI Core 异常转储兼容 CANN 9.2.0，版本门槛从 9.3 下放（#1956）（11-04）
    • 2.5 海光：feat/hcu-scaled-cast 分支四笔，E8M0 缩放转换与 64 bank 的 LDS 读规划（11-04）
    • 2.6 mlir-ascend：BishengIR 工具链适配器迁入 contrib（#209）（11-04）
    • 2.7 沐曦与摩尔线程窗口内无推送（11-04）
  • 三、生态与采用方
    • 3.1 TileOPs：单日 7 笔合入，GQA varlen 打包、ArgMax/ArgMin 每线程八向量、默认缓存策略（11-04）
    • 3.2 TileOPs 在审：一元逐元素单波次启动、FP8 分组 GEMM 连续布局、GEMM 可选 bias（#2829 / #2824 / #2823）（11-04）
    • 3.3 TileFoundry：FSDP 与 HSDP 的训练语义进 HIR 求值器（#284）（11-04）
    • 3.4 TileFoundry 在审：wgmma 同步计划（#286），另开一条命名统一的 issue（#287）（11-04）
    • 3.5 采用方与第三方：TileKernels 与 FlashQLA 无推送，三个第三方仓在 24 小时内活跃（11-04）
  • 四、社区、教程与活动
    • 4.1 新增 issue 4 条：warp 尺寸 64 的证明、PTO 有符号原子、CuTeDSL 的 NVFP4 量化器（11-04）
    • 4.2 新闻侧：四组查询、HN 与 arXiv 均无窗口内命中（11-04）
  • 五、趋势观察
  • 附、素材与核查说明

今日重点

这一天的主线不是「又合了什么」，而是「昨天的地基开始被验收」。10-14 合入的 #3512 把 T.Parallel 向量化到 PTO 的 VMI 流水线上，是昇腾后端一次结构性改动：整数类型从无符号语义的 i* 改成显式带符号的 si*，SIMD 的 lowering 改由新的 PTO 校验、合法化、向量化三段流程接管，显式 SIMD 并行循环的判定也重写了一遍。三周之后，这套改动在真实算子的形状上被逐个踩出四类回归：整数原子操作因为 pto.si32 不被接受而编译失败（#3789 报告、#3792 修复），T.clear 在 T.SimdVF 里走到只支持浮点的广播路径（#3795），标量 cast 的边界事实在 VMI 内存合法化里被过早绑定（#3793），显式 SIMD 的并行循环被当成标量 lane 访问而拒绝（#3796），另有一笔把整数原子适配到无符号载体（#3794）。五个补丁全部基于主线、彼此不依赖，作者在描述里反复写同一句「This PR is based directly on main and does not depend on the other PTO regression fixes」——这是有意的：让五笔各自可独立合入，避免像过去那样攒成一个巨型回滚包。验证也不再是「跑通了」，而是带具体用例数的硬件实测：47、18、16、37 项，四个独立算子（SwiGLU、Top-k、Cache、Engram 反向）作为代表用例。

第二条是度量口径的自我修正，而且两个仓在同一天做了同一件事。主仓 #3534 把 do_bench 的 L2 刷写从「往缓冲里写零」改成「读一遍缓冲」：写零会在 L2 里留下整块脏行，下一轮计时期间这些行要写回，等于把刷缓存的代价摊进了被测内核。#3534 用 cache.view(-1, 4096).sum(dim=1) 分块读，并顺手删掉了 fast_flush 这个公开参数——那个参数本来只是选择用 int32 还是 int8 建缓冲。昇腾侧的 msprof 后端改得更彻底：原来用一个 SimtVF 循环把零写进 UB 再 WTS_FV 存储，现在改成按 NORMAL_FV 从全局读入 UB，类型也固定为 int32。TileOPs 的 #2822 是同一思路的孪生改动，还额外记录刷写方法、只在方法一致时才与历史读数比较。三笔改动放在一起看，是一个信号：性能数字开始要求「可跨版本比较」，而不只是「这一次比上一次快」。

第三条是硬件前沿的补票。B300 上报的目标架构是 sm_103a，而 CuTeDSL 的 TMA 辅助函数用一张白名单把架构挡在门外，于是 B300 上六个 CuTeDSL 用例在还没进 DSL 之前就抛 ValueError；#3772 把 sm_103a 与 gr100（cc 10.7）上报的 sm_107a 一起加进白名单，作者在 B300 节点上复现了 8 通过 1 跳过、在 gr100 上 13 通过 1 跳过。同一窗口里 #3790 走另一条路：SM120 上 FP8 加 fp32 累加的 GEMM 一直用 SM89 的 mma.sync 原子，在 RTX 5090 上只有 cuBLASLt 的 68%，换成带 ue8m0 单位缩放的 block-scaled MMA 后，8192 立方从 463 升到 727 TFLOPS，且输出与旧原子逐位一致。两笔的性质不同——一笔是「让新卡能跑」，一笔是「让新卡跑满」——但都指向同一件事：架构支持的边界正在从 Hopper 往 Blackwell 与国产加速器两头同时扩。

第四条留给 TileOPs 的提交节奏。单日 7 笔合入里，5 笔是性能、1 笔 CI、1 笔文档：GQA varlen 在 SM90 上按 KV 组打包请求，把短请求的 key/value 读取次数降到每个 KV 头一次；ArgMax/ArgMin 在能一波驻留的前提下让每线程抱八个 16 字节向量，(2048, 4096) 的 bf16 ArgMin 从 8.704 微秒降到 7.393 微秒，快过 FlagGems 的 8.543；一元逐元素与归一化类内核把「只读一次」的操作数改回默认缓存策略，只在真正重复读的地方保留 evict-last/evict-first。这些改动的共同点不是「更快」，而是「知道为什么快」——每笔都带对照表、都标出自己仍然落后于对比项的那一行（Softmax 的 (32768, 4096) 就从 121.792 退到 131.648 微秒）。

一、核心项目进展

1.1 主线：CuTeDSL 的 TMA 架构门接受 sm_103a 与 sm_107a，B300 与 gr100 上六个用例恢复（#3772）

合于 11-04 12:46（北京时间，下同），是本期最早的一笔合入。B300 报告的目标架构是 sm_103a，而 CuTeDSL 的 TMA 辅助函数维护一张架构白名单，遇到不在表里的值直接在进 DSL 之前抛错，错误信息把允许的取值列成 sm_90、sm_90a、sm_100a、sm_110、sm_120、sm_120a、sm_121、sm_121a。在 fa69f7fa 的主线上、配 nvidia-cutlass-dsl 4.7.0 与 CUDA 13.3，B300 上有六个用例因此失败：test_tilelang_jit_cutedsl.py 里 tma_load 路径三个（GEMM JIT、多流内核、动态形状各一），以及另一个 tcgen05 的 CuTeDSL 测试文件里 tma_store 路径的三个 tmem_copy_roundtrip 变体。DSL 本身是接受 sm_103a 的——它的 tcgen05 操作接受整个 sm_100f 家族——所以问题纯粹出在这道门禁。加上这一项之后，两份文件在 B300（cc 10.3、CUDA 13.3、cutlass-dsl 4.8.1）上是 8 通过 1 跳过。同一笔还加了 sm_107a，因为 gr100 的 bring-up 板（cc 10.7）上报这个值、会撞同一道门；两项都加上后，同样的文件在 gr100 上是 13 通过 1 跳过，其中含 #3685 引入的 gather/scatter 文件。

值得记的是作者交代的验证边界：手上没有 sm_103 硬件，改动是靠另一位报告者（0z5a，在 #3685 里首次报告并 bisect）提供的 B300 机器验证的，CI 不覆盖 sm_103a。这类「代码合入、验证借机器」的情形在架构追赶期会越来越多，读者需要知道哪些能力是有实测背书的、哪些只是加入了白名单。

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

1.2 主线：L2 刷写由写零改为只读归约，do_bench 去掉 fast_flush 参数（#3534）

合于 11-04 11:27，改动 7 个文件、新增 36 行、删除 49 行，其中三处是真正的实现，四处是跟着改的测试。原来是 cache.zero_()：用一个 cache_size MB 的缓冲写零来挤掉 L2，问题在于写零之后缓冲里全是脏行，下一轮被测内核执行时这些行要写回内存，写回的带宽被算进了被测内核的时间里——缓冲区越大，这个偏差越明显。现在改成 cache.view(-1, 4096).sum(dim=1, dtype=cache.dtype)：按 4096 列分块读，把行读进 L2 并产生归约输出，脏行的问题不存在；分块是为了避免整块归约带来的全局同步；保留 dtype 是为了不触发一次全缓冲的 int64 转换。这一改动同时落在 bench.py 的 estimate 循环、torch_bench.py 的 event、cupti、cudagraph 三条路径，以及 msprof.py 的昇腾专用刷写内核上——后者原来用 SimtVF 把零写进 UB 再以 WTS_FV 存回全局，属于典型的写式刷写，现在改成以 NORMAL_FV 从全局读入 UB，元素类型也从 int8/int32 二选一固定成 int32。

公开接口的变化需要留意：do_bench 的 fast_flush 参数被删除。它原本只是在 int32 与 int8 两种缓冲元素类型之间切换（int32 建缓冲更快），现在缓冲固定为 int32，参数没有意义。对直接调用 do_bench(..., fast_flush=False) 的脚本，这是破坏性改动；仓库内的测试已经同步修改，对外部使用者则要靠 release note 提醒。

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

1.3 主线：AscendC CI 的自动调优供给用例显式同步两台设备，异步拷贝错误不再伪装成 arange 失败（#3784）

合于 11-04 17:01。现象是 AscendC 的 CI 作业里，自动调优的 set-inputs 用例中 [other_npu] 这个参数化项在下一次调用 torch.arange() 时报 srclen=4, dstlen=4——一次 4 字节的 memcpy 尺寸不符，而这正好是跨设备拷贝探针的形状。真正的原因是跨设备拷贝失败被异步上报，错误浮出的位置比发生的位置晚一步，于是报错点看起来在一个完全无关的 arange 上。修法很朴素：在探针前后各加一次设备同步（当前设备与 npu:other），让异步错误落回 try ... except RuntimeError 的作用域内。验证是单机双卡跑该用例文件，四项全过，且当跨设备拷贝真的异步失败时，[other_npu] 这一参数化项被跳过而不是让整个用例失败。

这笔本身很小，但它接的是上一期 #3546 的尾巴——那笔把自动调优的输入供给改成按当前设备解析，验收时提到昇腾双卡上的回归，这里补的是同一场景在 CI 上的稳定性。同一功能连着两个窗口各修一次，说明昇腾侧的自动调优仍是「能跑但边界不稳」的状态。

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

1.4 在审：ThreadSync 的 warp 一致性证明补上由 warp 索引拼出的组合谓词（#3787）

这笔接手的是上一期 #3761、#3765 留下的同一片区域：ThreadSync 要证明「屏障参与线程在整个 warp 内一致」，证明不了就拒绝 lowering。触发形状是 192 个线程下一个组合谓词 tx < 160 and tx // 32 % 2 == 0——warp 0、2、4 整组参与，warp 1、3、5 整组不参与，没有任何 warp 被劈开，本该是可证明的，但主线直接抛内部错误「参与线程不是 warp 一致」。作者做了一张比对的表，把缺陷范围压得很准：单项谓词 tx // 32 % 2 == 0、tx % 64 < 32、tx < 160 在主线上都能 lowering，只有「边界项与 warp 索引取模项用 and 拼起来」这一种形态会挂；而 tx < 168（真的劈开了第 5 个 warp）在两边都正确拒绝。根因在证明器的编码：它把线程号重命名成两个假想线程，让算术求解器证明 linear1 / warp_size == linear2 / warp_size 能推出两侧谓词相等，而 FloorDiv 与 FloorMod 的编码在「上界 + 取模」这种组合下建立不起两个除法的关系。修法是把 threadIdx.x 改写成 warp_index * warp_size + lane，两个假想线程共享 warp 索引、各自一个自由 lane，正好就是「同一个 warp」这个关系；于是只通过 warp 索引读线程号的谓词在两边塌成同一个表达式，一致性可证，而真正依赖 lane 的谓词仍然不等，#3786 里列出的 tx < 168 依旧被拒——证明没有放松，只是补上了漏掉的一类。

同一作者同日还提交了 #3786：T.copy 的目的区域超出缓冲时改为编译期报错，而不是生成越界访问。两笔合起来是同一取向——把「编译器默默生成错代码」变成「编译器说不」。

来源：https://github.com/tile-ai/tilelang/pull/3787、https://github.com/tile-ai/tilelang/pull/3786

1.5 在审：全复制 fragment 的包含证明补线程归属校验（#3788）

ProveFragmentContains 在判断「大 fragment 是否覆盖小 fragment」时有一条快速路径：只要大 fragment 是完全复制的，直接返回真。这条路径在检查线程归属之前就返回了，于是出现这样的误判：小 fragment 拥有线程 0 到 127，大 fragment 完全复制但只拥有线程 0 到 95，96 到 127 根本没有被覆盖，函数却说包含成立；开启 check_forward_index 时，两个完全复制但物理槽位不同的 fragment 也会被接受。后果是布局推断把不兼容的映射当成兼容，属于「错误被静默放过」的类型，报告者是在查 #2399 时顺手发现的。修法是先看物理索引，再把小 fragment 的线程映射换算到大 fragment 的坐标系里（要考虑有界的线程范围偏移），用算术分析器证明落在复制范围与线程范围内，证不出来就返回假。验证在 A100 上对上游主线新构建，通过一个测试期编译的桥接直接调用真实的 C++ 实现，覆盖线程数 128 对 96 的拒绝、96 对 128 与 96 对 96 的接受、物理槽位、非连续归属、有界范围、符号索引与部分复制。

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

1.6 在审：SM120 的 FP8 GEMM 换用 block-scaled MMA，RTX 5090 上 8192 立方从 463 到 727 TFLOPS（#3790）

这笔是本期性能侧最有分量的一条。问题出在指令选择：SM120 上 FP8 加 fp32 累加的 T.gemm 在 RTX 5090、8192 立方的形状上只有约 465 TFLOPS，而 cuBLASLt 是 681；调 tile 与流水级数都没用，因为瓶颈是指令本身——TileLang 发射的是 SM89 时代的 mma.sync ... f32.e4m3.e4m3.f32，这种原子在这张卡上只有半速。作者做了一组只测寄存器的探针：e4m3 到 f32 走 mma.sync 是 512 TFLOPS，而走 kind::mxf8f6f4.block_scale（单位缩放）是 1021，走 f16 累加的 mma.sync 是 1022——差距全在指令上。既然 ue8m0 缩放取 0x7f（即 2 的 0 次方）时 block-scaled 与普通 MMA 结果相同，就可以在 TL_SM120_BLOCK_SCALED_MMA_ENABLED 打开且目标为 sm_120a/f、sm_121a/f 时，让四个 fp32 累加的 FP8 派发路径改走 block-scaled 原子，其余目标、f16 累加与显式非 a 后缀的 sm_120 保持原样。逐位一致性在指令级上验过：e4m3 与 e5m2 的四种组合，随机寄存器内容含 NaN 与 Inf 编码。

数字：e4m3 GEMM、bf16 输出、128×128×128 tile、3 级流水、256 线程的条件下，4096 立方从 408 升到 673（cuBLASLt 530），8192 立方从 463 升到 727（681），4096×14336×4096 从 449 升到 696（498），256×8192×8192 从 348 升到 606（457）。唯一变慢的是 M=16 配 128 行 tile 的形状，31.5 微秒退到 33.5 微秒，作者说明该形状受带宽限制、换 16 或 64 行 tile 会更快。新增用例直接检查 PTX：sm_120a 必须拿到 block-scaled 原子、sm_89 必须拿到普通原子。

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

1.7 在审：Profiler 与自动调优的口径校验成组提交，六笔同日出（#3797 / #3798 / #3799 / #3716 / #3717 / #3718）

六笔都来自同一位作者（Afloat16），主题是把「验证器与被测对象的口径」对齐。assert_allclose 现在会拒绝内核在该给张量的位置返回 None——原来的实现只要有一侧是 None 就跳过比较，于是一个返回空的内核也能「通过」（#3797）；assert_consistent 原来直接对返回值做 zip，标量张量会抛 TypeError，缺一个或多一个输出会被静默忽略，形状不同则可能靠截行或广播蒙过去，现在先归一化成输出列表再核对数量与形状，空张量也纳入（#3798）；CUPTI 基准测试里的 suppress_stdout_stderr 原来对 Python 流调用 fileno()，笔记本与 StringIO 没有文件描述符就崩，现在分别重定向原生描述符 1/2 与 Python 流，并用 ExitStack 保证异常路径也能还原（#3799）。另外三笔在自动调优侧：目标字典作为缓存键时用字符串表示，插入顺序不同的等价字典会产生不同键，改成按排序后的 JSON 序列化（#3716）；Tensor Core 检测优先用显式算力、再用显式或当前目标架构，最后才查本地设备，否则无 GPU 的交叉编译会误报不支持（#3718）；Profiler 的动态形状替换补上 m + 1 这类复合维度，并按缺失变量给出 ValueError 而不是 TypeError（#3717）。

这一组的技术含量不在改动本身，而在它们暴露的形状：验证器长期是「通过即正确」的薄弱环节，而自动调优的缓存键、目标解析这类基础设施一旦有偏差，会以「结果不可复现」的形式出现，很难归因。

来源：https://github.com/tile-ai/tilelang/pull/3797、https://github.com/tile-ai/tilelang/pull/3798、https://github.com/tile-ai/tilelang/pull/3799

二、多后端适配

2.1 昇腾 PTO：一笔架构改动引出四类回归，五个互不依赖的补丁同日提交（#3792 / #3793 / #3794 / #3795 / #3796）

先说清背景。10-14 合入的 #3512（[Ascend][PTO] Vectorize T.Parallel to PTO VMI）把并行循环的向量化从旧的 SIMD 通道迁到 PTO 的 VMI 流水线：整数类型从无符号的 i* 改成显式带符号的 si*，Lowering 走 Verify、Legalize、Vectorize 三段，显式 SIMD 循环的判定也重新实现。这五个补丁各自对应它在真实算子上踩出的一处：

  • #3792 整数原子编译失败：codegen_pto 为整数发射 pto.si32 / pto.si64，而 PTO 的原子接口只接受无符号的 i32 / i64，于是 T.atomic_add、T.atomic_max、T.atomic_min 在整数缓冲上直接编译失败（对应 issue #3789）。修法是在 tl.atomic_* 外面套一层兼容包装：指针元素类型与操作数重写为同宽无符号，整数 min/max 传正确的符号性，返回值再转回原类型。用例从 15 项扩到 18 项，覆盖有符号与无符号 32/64 位的加减取极值、负操作数、前值返回、超过 32 位的计数器。
  • #3794 有符号载体适配：与 #3792 同源但走另一条路——不改 tl.atomic_* 的语义，而是在生成的 PTODSL 调用边界上重写指针视图与操作数，保留地址空间，浮点原子保持原路径。代表用例是 SwiGLU 的 clamp 计数（bf16、hidden 2048、clamp 0.5、psum 布局），18 项原子用例全过。
  • #3795 SIMD 填充丢失：T.SimdVF 里的 int32 T.clear 走到了只支持浮点的广播路径，原因是新 PTO 流水线不再运行旧的 SIMD pass，而那个 pass 同时负责填充与并行循环。修法是抽出一个只做填充的入口，复用原有的 LowerFill，放在缓冲版本归一化之后、PTO 并行校验之前，仍发射原来的 vdup/vsts 序列，并保留连续区域、对齐、整向量以及多缓冲版本切片的检查。16 项用例通过，代表用例是 Top-k（512 token、1 专家、top-1）。
  • #3793 标量 cast 的边界事实被过早绑定：u = uint32(i - 1)、i 在 [0, 63] 这种与 VMI 地址无关的标量，会在内存合法化里抛错；急于 Bind 还会留下一个被低估的区间，等 IntSet 的 cast 分析再抛。修法是先检查整数 cast 的可表示性再安装事实，给证明引擎留安全的标量定义、给 IntSet 保守的数值区间，避免直接求 cast 端点。47 项用例通过，代表用例是 Cache（block_size=64、n=65536）。
  • #3796 显式 SIMD 并行循环被当成标量 lane：Engram 反向的初始化用了一个 Parallel 嵌套，其中显式 SIMD 存储按整向量块推进，新的 lane 布局推断把它当标量访问并拒绝（M=2, N=32, L=128）。修法是让显式 SIMD 的分类在布局推断与并行单元发现之间共享，这类嵌套保留 #3512 之前的通用布局路径，自动逐元素单元仍走新的 PTO lane 路径。37 项用例通过，代表用例是 Engram 反向（512 token、hc 4、hidden 2048）。作者同时写明一条限制：FP32/128 lane 的显式加自动混合用例只有生成源码级覆盖，当前 PTOAS/Bisheng 设备编译器在该组合上超过 6 KiB 的 VF 栈预算，没有设备运行结果。

把这五笔并排看，能读出昇腾侧的工程节奏：#3512 那样的结构性改动先落地、再按真实算子逐个暴露边界，每个补丁都带硬件实测与用例计数，并且刻意保持彼此独立。「先合框架、后清回归」这套打法在上一轮架构迁移（#1975）里已经用过一次，现在成了常规流程——代价是主干在这类窗口里处于「已知有几处不工作」的状态，外部使用者踩到的概率不低。

来源：https://github.com/tile-ai/tilelang/pull/3792、https://github.com/tile-ai/tilelang/pull/3793、https://github.com/tile-ai/tilelang/pull/3794、https://github.com/tile-ai/tilelang/pull/3795、https://github.com/tile-ai/tilelang/pull/3796

2.2 昇腾：Cube 与 Vector 的核心比可配置，不再固定 1:2（#3783）

昇腾的 AI Core 分 Cube 与 Vector 两类，AscendC 支持 mix(1,2)、mix(1,1)、mix(1,0)、mix(0,1) 等分组，而 TileLang 此前一律按 1:2 启动，等于把一个硬件自由度的上限锁死，也让未来核心比不同的芯片难以适配。这笔改动把核心比变成一个显式或按设备推导的参数：T.MixedKernel(..., sids=1) 或 sids=2 按请求的分组大小走并做能力校验；省略或传 None 时用当前设备的 vector_core_num / cube_core_num，取不到设备信息就退回 2；T.Kernel 保持原接口与核心类型判定，识别为混合内核时用准备好的向量核数，单核路径不变，显式 T.Vector(1/2) 的分组语义保留。选定的组数会一路带到 sid 域、任务分配、同步与代码生成，避免重复查询设备与重复归一化；冲突的声明与超过设备能力的请求直接拒绝。

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

2.3 昇腾：AutoSchedule 的周期下界收紧，Z3 求解时间降 10.04%（#3776）

AutoSchedule 用二分搜索找流水线的启动间隔（II），原来的下界取所有任务 II 的最大值，偏松，二分要多走几轮。这笔在资源模型下补了三条必要条件：单任务界（保留 P >= max_i II_i）、资源团界（对两两资源冲突的团按残差排序求和，得到 P >= sum_i II_i）、自依赖界（S_u - S_u >= wait - d*P 推出 P >= ceil(wait/d)，版本相关的距离按流水级数截断）。用三者最大值初始化二分，其余逻辑不动。测试环境是昇腾 950PR、CANN 9.3.0 周更版，关缓存、一暖三测交替：examples/ascend 全部编译通过且调度结果不变，Z3 求解阶段的累计耗时下降 10.04%。这类「不改结果、只改搜索起点」的改动风险最低、收益却是全局的——编译期的求解时间是每个用户都要付的。

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

2.4 昇腾：AI Core 异常转储兼容 CANN 9.2.0，版本门槛从 9.3 下放（#1956）

这笔合在 tilelang-ascend 仓的 ascendc_pto 分支（11-04 10:18），是本期昇腾仓唯一的合入。AI Core 异常转储功能原先声明依赖 CANN 9.3.0 以上，作者在 9.2.0 上做了可行性验证并修了三处：acldumpTensorInfo 的必填字段补齐（addrType=ACL_DUMP_ADDR_RAW、placement、shapeNum、originShapeNum），因为 9.2.0 会强制校验 addrType，而原实现靠 calloc 置零得到的 ACL_DUMP_ADDR_PTR 被拒绝；跳过空地址与零大小张量，9.2.0 不接受 tensorAddr 为 NULL；注册接口从已废弃的 aclrtSetExceptionInfoCallback 换成 aclrtExceptionInfoCallbackRegister。工具侧 parse_exception_dump() 兼容 9.2.0 输出的 .npy 与旧版的 .bin 两种格式，中英文教程的版本要求一并改为 CANN 9.2.0 以上。对还在 9.2 上的用户，这是从「功能不可用」到「可用」的变化。

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

2.5 海光：feat/hcu-scaled-cast 分支四笔，E8M0 缩放转换与 64 bank 的 LDS 读规划（11-04）

海光仓的主分支停在 11-02，本窗口的动静全在功能分支上，四笔提交集中在八小时内（北京时间 16:26 到 20:16），作者来自海光（huangteng@hygon.cn 与 zhouweiwang@hygon.cn）：

  • feat(hcu): add scaled_cast and unify GEMM LDS read planning：新增 E8M0 缩放的 FP8/BF8/FP4 转换，标量与打包两种 lowering，新增 src/tl_templates/hcu/scaled_cast.h（223 行）与 lower_scaled_cast.cc；同时把 GEMM 消费端元数据在矩阵读选择与 LDS 策略之间共享，支持单级 LDS 切片，主变换文件改动 650 行。
  • fix(hcu): align GEMM LDS strategies with 64-bank read phases：按 bank-word 几何推导 plain-copy 的 swizzle，为矩阵与两种原生 pack 尺寸共用的窄 b16 tile 提供显式逆变换；混合或受限相位的异步消费端优先用兼容的公共存储；选择原生 wrap 时把完整的 b128 读相位计入。12 个文件、新增 1536 行、删除 112 行，新增 5 个测试文件。
  • fix(hcu): prefer mixed LDS wraps proven at most two-way：给混合 LDS wrap 加「至多两路」的证明，只在证明成立时优先使用。
  • fix(hcu): use physical read width for mixed wrap checks：混合 wrap 检查改用物理读宽度，去掉原来按 kPack 推断 8 字节每 lane 的写法。

主题是共享内存的 bank 冲突与相位对齐——这是海光 DCU 上 GEMM 性能的关键路径，也说明该后端已经从「能编译」进入「按微架构调优」的阶段。四笔都在功能分支上，尚未开 PR，时机上可能与 TileLang 主线做一次同步后再提交。

来源：https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-scaled-cast

2.6 mlir-ascend：BishengIR 工具链适配器迁入 contrib（#209）

合入主分支（11-04 16:35），把 BishengIR 工具链的适配代码从原有位置搬到 contrib 目录。这是纯结构整理：把「随上游工具链演进而需要频繁改动的胶水层」与核心 lowering 分开，后续升级 BishengIR 时改动面收敛在 contrib 内。同一分支上还有一笔 refactor(contrib): move BishengIR toolchain adapters into contrib 的主体提交，合并提交是第二笔。这类整理单看没有功能价值，但它是把「适配层」当成独立生命周期管理的开始——昇腾工具链版本迭代快，胶水代码的可替换性比优雅更重要。

来源：https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

2.7 沐曦与摩尔线程窗口内无推送（11-04）

tile-ai/tilelang-metax 最后推送停在 10-26 15:08，tile-ai/tilelang-musa 停在 09-30 17:50，两仓在本窗口都没有提交，也没有在审 PR 更新。国产四家后端里，昇腾与海光是本窗口的两条活跃线，沐曦与摩尔线程处于静默期。上一期报告里沐曦同样无推送，已连续两个窗口；摩尔线程则是自 09-30 以来没有动静，值得在后续窗口留意是否转入内部开发。

来源：https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa

三、生态与采用方

3.1 TileOPs：单日 7 笔合入，GQA varlen 打包、ArgMax/ArgMin 每线程八向量、默认缓存策略（11-04）

7 笔全部在窗口内合入，节奏上是本季度最密的一天，逐笔看：

#2819 给 SM90 的 GQAVarlenFwdOp 加了打包的 warp 特化程序：当一个请求的 KV 头组能在一个 tile 内按位置铺满时，让请求按 KV 组填充 tile，于是每个 KV 头只读一次 key/value；是否走打包由代价模型判断，只有「形态保证能省下的量」超过实测的 9% 单头开销才用；宽窄请求在同一次启动里交替领取，不做主机回读、不加同步；当设备上统计出的不可拆项少于 SM 数量时，单 tile 的打包请求把 key 扫描拆成多段、用 fp32 工作区合并。另修了三个主线缺陷：sink 加分数缩放超过约 2.4 时无可见 key 的行返回 NaN、回退路径在请求首个 key 之前读到 RoPE 表、以及一个负向对照把调用方栈帧留在引用环里导致跨测试的 CUDA graph 被提前销毁。测试节点 173 增到 197。

#2828 给按行分块的 ArgMax/ArgMin 在「一波驻留」的前提下让每线程抱八个 16 字节向量，并把 64 线程块加入调优候选；驻留判断由 SM 的线程数、常驻 CTA 上限与寄存器共同决定。(2048, 4096) 的 bf16 ArgMin 从 8.704 微秒降到 7.393（对比项 FlagGems 8.543，快 1.155 倍），背靠背场景 7.808 降到 6.000；f16 ArgMax 从 6.816 降到 6.400（对比项 9.344，快 1.46 倍）。12 行基准里 4 行变快、0 行变慢。

#2820 把「每个元素只读一次」的操作数改回默认缓存策略：逐元素、rope、norm、softmax、归约类内核不再一律用 evict-last/evict-first，只在真正重复读的地方保留（二元内核的广播行、多批多头共享的 rope cos/sin 行、两遍 RMSNorm 的两次读、融合加法的前两遍）。同时给 RMSNorm 加了配对核与分阶段核两种实现、把 warp 行 Softmax 的 CTA 缩到刚好填满 SM 的 warp 数、2D neox RoPE 改用 128 线程，并把每架构的常驻 CTA 上限收进 RESIDENT_CTAS。代价是 Softmax 的 (32768, 4096) 从 121.792 退到 131.648 微秒（对比项 quack 126.080），作者在表里如实标出。

#2822 是 L2 刷写的只读改法，与主仓 #3534 同源，168 项测试通过，并额外记录刷写方法、只在方法一致时与历史读数比较。#2825 用单一重拟合的代价模型给 W4A16 的合法配置排序，split-K 的 CTA 只暂存自己那份量化元数据，共享内存记账把启动余量算进去。#2826 是配套文档，写清稠密与 tiny-M 的标定契约（测量方式、拟合约束、留出验证），没有运行时改动。#2827 把 off-CUDA 套件挪到夜间作业：原来 preflight 里这两步在 ubuntu-latest 上要 13 到 18 分钟、比其他检查晚十分钟出结果，而其中 tests/ops 的冒烟跑用的是算子参考实现，检查的是设备无关性与派发、不是内核结果；挪走后 preflight 只保留收集检查、丢弃测试报告与 cuda_only 必失败检查（约 3 分钟），1856 加 2789 个测试节点转到夜间，发布门禁仍要求夜间成功。

来源：https://github.com/tile-ai/TileOPs/pull/2819、https://github.com/tile-ai/TileOPs/pull/2828、https://github.com/tile-ai/TileOPs/pull/2820

3.2 TileOPs 在审：一元逐元素单波次启动、FP8 分组 GEMM 连续布局、GEMM 可选 bias（#2829 / #2824 / #2823）

#2829 让「体内只有一两条单指令操作」的一元逐元素启动跑成单波次：读写字节数不超过 L2 的四分之三、结果与操作数同宽、体内最多两个各降为一条指令的操作，满足三条才缩成 256 线程块的一波；体内条件由 TIR 变量上的节点计数判断，带 cast、除法、数学调用、守卫或不能在裸变量上求值的体一律保留默认启动。目前符合的是 relu、abs、neg、bitwise_not 与 bool 上的 logical_not。顺带修了一个真问题：ReLU 对 NaN 输入返回 0，现在返回 NaN（16 位用 torch 的位模式、float32 用规范 NaN）。基准里三行 ReLU 与一行 LogicalNotFwdOp 变快，其余配置不变；作者也标出 GPU 6 上带刷写的 decode 行只是打平。

#2824 给 SM90 的 FP8 分组 GEMM 加上连续布局（紧凑与对齐两种打包、段尾或逐行两种专家元数据），SM89 仍只支持掩码布局，并补齐按专家大小偏斜的实测负载与 DeepGEMM 连续布局的对比。#2823 给 GEMMFwdOp 加可选 bias[N]，覆盖稠密、GEMV 与 split-K 三条路径，附 cuBLASLt BIAS 对比。两笔都对应 #2130 里列的缺口，说明 TileOPs 正在按「与 torch 生态逐项对齐」的清单推进，而不是随机加算子。

来源：https://github.com/tile-ai/TileOPs/pull/2829、https://github.com/tile-ai/TileOPs/pull/2824、https://github.com/tile-ai/TileOPs/pull/2823

3.3 TileFoundry：FSDP 与 HSDP 的训练语义进 HIR 求值器（#284）

合于 11-04 16:36。此前的求值器把 reshard 当成全局张量处理，丢掉了不同参与者的贡献差异，于是训练语义（尤其是分片与复制混合的 HSDP）在 HIR 层无法正确求值。这笔实现 #255 的训练子集：按参与者求局部值、GPU 坐标、调用与循环、gather/scatter，以及沿轴的 Partial(sum) 归约，都走共享的投影规则；Reshard 的别名证明改成对归属敏感，接上 #170 的结构化生命周期，元组实参从张量叶子校验。配套加了一个受 modded-nanogpt 启发的两块 GQA 研究：显式前向反向、动量 SGD、独立的 PyTorch 参考，覆盖 FSDP、2×2 与 4×4 的 HSDP、两种驻留策略，报告集合通信组、载荷与存储占用，并嵌入可复用的 SVG，README 里的每张表都能从代码推出来。

验证：重构后的核心套件 775 项、后端套件 778 项（13 跳过，两张 sm90 卡），10 组 HIR 配置在 CPU 上通过；两种驻留策略在四张 sm90 上与真实的 2×2 PyTorch HSDP 对齐（跑两次更新，比较局部梯度、参数与动量）。

来源：https://github.com/tile-ai/TileFoundry/pull/284

3.4 TileFoundry 在审：wgmma 同步计划（#286），另开一条命名统一的 issue（#287）

#286 处理的是 WGMMA 的异步性：指令发出之后仍可能继续读共享内存与寄存器，所以后续访问必须排在它的效果之后。做法是给每个操作数效果声明完成条件，把 WGMMA 的 fence、commit、wait 挂到需要它们的 HIR 节点上，lowering 在指令的发射点消费这些步骤；若后续作用域放不下一个集合等待，就在发射作用域里把操作做完。同时修了寄存器 Reshard 的别名，让每个元素在持有副本的每个线程上都留在同一个寄存器，计划用这个存储身份追踪依赖。作者明确写了一条风险：矩阵指令的发射器还没接上，这笔不证明 WGMMA 的执行正确性，也不给性能，TMA 的完成语义另行处理。

#287 是一条文档向 issue，指出代码库里「一个概念五个名字」（mesh 的一个成员被叫作 placement 133 处、program 110 处、instance 92 处、participant 44 处，另有 wave_units 一类带 unit 的写法），而 unit 这一个词又同时指空类型、硬件计算单元与 mesh 成员；Python 与 CUDA 运行时对 local_layout 的签名还不一致。这类 issue 通常不会立刻改代码，但它是编译器项目从「能写」到「能被人长期读」的必经一步。

来源：https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/issues/287

3.5 采用方与第三方：TileKernels 与 FlashQLA 无推送，三个第三方仓在 24 小时内活跃（11-04）

两个采用方仓在本窗口都没有提交：deepseek-ai/TileKernels 停在 10-14 05:49，QwenLM/FlashQLA 停在 10-30 03:02（最新发布仍是 09-30 的 v0.1.3）。需要说明的是，采用方不推送并不等于不被影响——上一期 #3761 修的正是 TileKernels 的 mhc_pre_big_fuse 内核在 B300 上的屏障占用回退，说明编译器的改动会以「被动回归」的方式作用到它们的形状上。

第三方侧有三条由仓库发现捞到的活跃项，都是 24 小时内新建或推送：sablin39/tilelang-cuda-skills（149 星，给编写 TileLang 与用 CUDA 工具链调试提供技能集）、xwhzz/tilelang-hexagon（2 星，把 TileLang 往高通 Hexagon DSP 上带）、Index4NekoMeteor/Tilelang-AttnRes（1 星，用 TileLang 实现 Kimi K3 的注意力残差结构）。三条都很小，但方向有意思：一条做工具链侧的开发体验，两条做新目标架构与前沿模型结构的移植，属于社区自发扩展 TileLang 边界的那类工作。

来源：https://github.com/deepseek-ai/TileKernels、https://github.com/QwenLM/FlashQLA、https://github.com/sablin39/tilelang-cuda-skills

四、社区、教程与活动

4.1 新增 issue 4 条：warp 尺寸 64 的证明、PTO 有符号原子、CuTeDSL 的 NVFP4 量化器（11-04）

主仓三条。#3791 是「[BUG][ice-on-valid-code]」：ThreadSync 只在 warp 尺寸 32 下能证明 warp 一致谓词，同一份 IR 在 thread_warp_size 为 64 时被拒——这条与本期在审的 #3787 咬得很紧，说明证明器对 warp 尺寸的假设是硬编码的，而 AMD 与部分国内加速器是 64 宽 warp。#3789 报 PTO 后端为整数 T.atomic_add 发射带符号的 pto.si32/si64 导致校验失败，当天就有 #3792 与 #3794 两个方向的修法。#3781 报 CuTeDSL 的 NVFP4 block-scaled 量化器无法编译：打包 FP4 的重解释与缺失的 ceil/floor。加上 TileFoundry 的 #287（命名统一），本窗口共 4 条新 issue，全部是「内部实现细节泄漏到用户可见面」的类型，没有一条是性能抱怨。

来源：https://github.com/tile-ai/tilelang/issues/3791、https://github.com/tile-ai/tilelang/issues/3789、https://github.com/tile-ai/tilelang/issues/3781

4.2 新闻侧：四组查询、HN 与 arXiv 均无窗口内命中（11-04）

Google News 的四组查询（英文 TileLang GPU kernel、"TileLang" OR "tile-lang"，中文 TileLang 算子、TileLang 编译器）里，最新一条命中停在 352 小时前（10-21 SemiAnalysis 的华为与 DeepSeek 报道），其余集中在 09-30 那一波「DeepSeek 开源昇腾版 TileLang、DeepGEMM、DeepEP」的报道（856 小时前），全部早于本窗口；Hacker News 以 tilelang 检索最近 24 小时为 0 条；arXiv 的 all:tilelang 最新命中仍是 07-24 的 TileSight，窗口内没有新论文。GitHub 仓库发现捞到的 9 条候选里，7 条是仓库自身，1 条是 Swarm World 这类与主题无关的 Show HN，1 条是无关的 2048 类项目。可以确认本窗口新闻侧为零，报告主体只能来自提交与 PR——这已经是连续多期的稳定形态。

来源：https://news.google.com/、https://news.ycombinator.com/、https://arxiv.org/

五、趋势观察

第一，昇腾 PTO 的清账模式已经定型，它的收益与代价都值得记录。#3512 在 10-14 合入，三周内不断有真实算子把它踩出边界，今天的五笔补齐只是最新一批。这种「先落架构、后按算子清回归」的节奏在 #1975 整仓迁移时也出现过，好处是架构迭代快、单笔改动可控，坏处是主干在相当长的窗口里处于「已知有几处不工作」的状态。五个补丁刻意保持互不依赖，说明维护者有意识地在降低回滚粒度——这比攒成一个大回滚包要成熟，但也意味着需要更强的发布前回归纪律，否则「独立可合」会变成「独立可漏」。

第二，性能度量开始被当作产品特性来对待。主仓 #3534 与 TileOPs #2822 在同一天把 L2 刷写从写零改成只读，理由都是「脏行写回污染下一轮计时」；TileOPs 更进一步，记录刷写方法、只在方法一致时比较历史读数。这不是小修小补：一个基准数字如果在不同版本间不可比，那么所有「快了多少」的结论都要重算。把度量方法本身纳入版本管理，是性能工程走向可信的前提。

第三，算子库的竞争点从「算子数量」转向「选择器质量」。TileOPs 今天合入的 W4A16 单一重拟合代价模型与稠密 GEMM 的标定契约文档，加上在审的 FP8 分组 GEMM 连续布局与 GEMM bias，构成一条清晰的线：先把每个形状的合法配置枚举出来，再用一个统一的代价模型排序，标定过程写成可复现的契约。这条线做对了，算子库的性能就不再依赖个别内核的手工调优；做不对，就会退化成「某些形状特别快、其他形状靠运气」。

第四，架构支持的边界在两头扩。一头是新卡：B300 的 sm_103a、gr100 的 sm_107a 进白名单，SM120 的 FP8 换成 block-scaled MMA；另一头是国产加速器：海光在 64 bank 的 LDS 相位上做微架构级调优，昇腾在核心比与异常转储这类工程细节上补齐。这两种扩法的工作量结构完全不同——新卡主要是「让路径能走通」，国产卡更多是「让同一条路径跑出接近硬件的效率」，后者需要的实测密度高得多，也正是海光那四笔提交里测试文件占比那么大的原因。

最后一点风险面：一个窗口内 42 笔在审 PR 更新、其中 12 笔带昇腾/PTO 标签、7 笔带 CUDA/SM/CuTeDSL 标签、4 笔集中在 Profiler，这个结构说明仓库的变更压力正从「加功能」转向「修边界」。对使用者来说，好消息是缺陷在被逐个指认；坏消息是这些缺陷大多落在验证器、缓存键、证明器这类「平时看不见」的地方，等它们显形时往往已经影响到结果的可信度。

附、素材与核查说明

信源核查表

信源 核查结果
tile-ai 组织（29 仓库） 窗口内推送 8 个仓库
主仓 tilelang 主分支 3 笔合入；新开 issue 3 条；在审更新 38 笔（昇腾/PTO 12、CUDA 与 CuTeDSL 7、Profiler 4）
TileOPs 主分支 7 笔合入；在审更新 3 笔；无新 issue
TileFoundry 主分支 1 笔合入；在审更新 1 笔；新开 issue 1 条
tilelang-ascend 默认分支 ascendc_pto 1 笔合入；在审清单因匿名接口限额未取到，改以主分支与提交原子源核查
tilelang-mlir-ascend 主分支 2 笔（含 #209 合入）
tilelang-hygon 主分支无新提交（停在 11-02）；feat/hcu-scaled-cast 分支 4 笔
tilelang-metax / tilelang-musa 窗口内无推送
TileOPs-nightly 仅分支推送，主分支无新提交
发布页（tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon） 窗口内无新版本，tilelang 最新仍为 09-30 的 v0.1.15
FlashQLA / TileKernels 窗口内无推送
Google News（中英文四组查询）/ Hacker News / arXiv 无窗口内实质命中，最新命中为 352 小时前

完整信源清单

核查说明

本期所有合入、在审、issue 与提交数量均来自 GitHub 检索与原子源的实测输出，未做估算；昇腾仓与海光仓的在审 PR 清单因匿名接口额度耗尽未能取到，相关结论只基于主分支提交与分支级原子源，已在信源核查表中注明。tilelang-hygon 的四笔提交位于未开 PR 的功能分支，属分支级动态，不代表主线已合入。Google News、Hacker News 与 arXiv 三条新闻通道在本窗口均无命中，报告主体因此全部来自代码侧。