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

信源：GitHub tile-ai 组织逐仓核查（29 仓库，窗口内推送 9 个）+ 合入与在审真值检索（合入 18 笔、在审更新 44 笔）+ 新增 issue 全量核对（8 条）+ 主分支逐笔提交核查（tilelang 5、TileOPs 7、TileFoundry 1、tilelang-ascend 4、tilelang-hygon 1、tilelang-lsp 1、文档站 5）+ 发布页与 tag 核查（5 仓）+ 采用方两仓 + Google News 中英文查询（走代理）+ Hacker News + arXiv。

本期索引

  • 今日重点：昇腾 950 的持久化 FlashMLA 内核落地（#3678，新增 1997 行、42 次提交）；向量化规划一处真实回归修复（#3734）；#line 指令改为逐行重锚定（#3703）；TileOPs 单日 7 笔性能与测试基建合入；TileFoundry 把共享内存排序做成可证明的分析层（#283），并测出 main 的 lowering 在 sm_89 上 7 个程序错 4 个；24 小时合入 18 笔、在审更新 44 笔、新增 issue 8 条
  • 一、核心项目进展
    • 1.1 主线：昇腾 950 的持久化 FlashMLA 内核，BF16 prefill 与紧凑 FP8 decode 双路径（#3678）（11-02）
    • 1.2 主线：cast 解耦被跳过后残留的 512 位向量，SM100 上 ViT 前向编译失败（#3734）（11-02）
    • 1.3 主线：#line 指令改为逐行重锚定，调试器映射可选密集模式（#3703）（11-02）
    • 1.4 主线：注意力示例按设备共享内存上限推 tile，RTX PRO 6000 上由 2 过 7 败变为 9 全过（#3702）（11-02）
    • 1.5 主线：CUDA float32x8 位转换改用内建 bit cast，SASS 指令从 19 条降到 18 条（#3748）（11-02）
  • 二、多后端适配
    • 2.1 昇腾：WY Fast 前向算子新增、wy_fast_bwd_split 反向分段迁移（#1827 / #1828）（11-02）
    • 2.2 昇腾：T.tile.fill 与 T.tile.clear 文档事实核查合入（#1634 / #1628）（11-02）
    • 2.3 昇腾在审：cython wrapper 缓存符号形状解析，承接宿主开销这条线（#1967）（11-02）
    • 2.4 昇腾在审：z3 版同步从主线迁入，整仓迁主线同步推进（#1965 / #1975）（11-02）
    • 2.5 昇腾在审：归约算子两个缺陷，首 Parallel 轴广播改标量读（#1978 / #1968 / #1979）（11-02）
    • 2.6 海光：gfx948 按 gfx946 兼容处理，先让既有指令可测（#32）（11-02）
    • 2.7 摩尔线程、沐曦、mlir-ascend 窗口内无推送（11-03）
  • 三、生态与采用方
    • 3.1 TileOPs：单日 7 笔合入，6 笔性能、1 笔测试基建（11-02）
    • 3.2 TileOPs：warp 专用 decode 的 softmax 与 QK/PV 重叠，478 次扫描调用几何均值 0.952 倍（#2817）（11-02）
    • 3.3 TileOPs：MoE 按 token 收集预置换行，H200 上 387.8 TFLOPS（#2815）（11-02）
    • 3.4 TileOPs：16 位逐元素去掉带范围检查的除法，8 字节操作数一线程一向量（#2814 / #2813）（11-02）
    • 3.5 TileOPs：Norm 统计量改为减行内元素，消掉靠近 1e36 的溢出（#2811）（11-02）
    • 3.6 TileOPs：tileops_test 目标让 tests/ops 脱离 CUDA 跑（#2816）（11-02）
    • 3.7 TileFoundry：共享内存排序做成可证明的分析层，83 个夹具计划 52 拒绝 31（#283）（11-02）
    • 3.8 TileFoundry 在审：CUDA lowering 发出同步计划并反向校验，main 上 7 个程序错 4 个（#285）（11-02）
    • 3.9 采用方：FlashQLA、TileKernels 窗口内无推送（11-03）
  • 四、社区、教程与活动
    • 4.1 昇腾同步正确性成为集中报障面：单日 8 条新 issue，两名报告人分别系统性提交（11-02）
    • 4.2 跨行与转置视图的 SIMD 加载丢掉 DMA 依赖，根因落在逻辑区域重建（#3749）（11-02）
    • 4.3 新闻侧：Google News 中英文、Hacker News、arXiv 三通道均无窗口内命中（11-03）
  • 五、趋势观察
  • 附、素材与核查说明

今日重点

这一天的主线只有一件事：TileLang 把 DeepSeek 的 MLA 解码内核搬到了昇腾 950 上，而且是按编译器可调度的写法搬的。#3678 一次性加入 1997 行、20 个文件、42 次提交，覆盖 BF16 prefill 与紧凑 FP8 decode 两条路径，还带可选的 FP8 与 FP4 额外 KV 缓存、attention sink、非法索引与逐 query 独立长度。实现方式值得注意：注意力循环、在线 softmax、反量化与调度全部写在 TileLang 里，靠 T.Stage 让编译器管理缓冲版本与同步，QK 与 PV 用同一块 L0B 存储环的两个不同 layout 视图，L1 的落位用 T.annotate_buffer_addresses 直接指定；prefill 在一个累加器分配里分出四个不相交的输出切片，归一化算完一片就排空一片。为了让这套写法成立，编译器侧同时补了三处能力——L0C 到 UB 的拷贝保留 NZ 输出、每个 AIV 流在自身流水线之前选定、显式调度的阶段在物理发射相位相同时保持源码顺序。目标形状写得很窄：头数 64、Dqk 与 Dv 都是 512、topk 为 64 的倍数。

第二条线是修一处真实的编译器回归。#3734 处理的是 SM100 上的 ViT 前向：Cannot convert type float32x16 to CUDA type。路径很长但每一步都清楚——向量化规划器为一个 256 位的 BF16 存储选了 16 个 lane，并推迟 cast 约束，指望 DecoupleTypeCast 把计算和访存拆开；而 #3579 之前刚把「读内存的守卫」从解耦里排除（重复这些守卫会改变语义），于是最终合法化那一步仍在按推迟策略选宽度，残留的 FP32 表达式被撑成 512 位向量，CUDA 后端表达不了。修法是把推迟约束做成规划器的参数而非 IR 属性，只有最后一步合法化传 false，守卫的 FP32 转 BF16 降到 8 个 lane，成功解耦的 BF16 拷贝仍保留 16 个 lane；地址 cast 也从数据向量类型约束里摘出来，靠既有的访问分析判合法性。

第三条线是调试体验。#3703 把 #line 指令从「每条 TIR 语句只锚定首行」改成可选的分级发射：级别 0 关闭、1 保持原状、2 以上对每一行生成的非空行重新锚定，包括被展开成多行的语句和收尾大括号。实现放在 CUDA、HIP、CPU C、C host、Metal、Ascend 共用的 C 系输出流里，不动 vendored 的 TVM；没有源码跨度的语句和函数显式切到 <tilelang-generated>，用户原文逐字保留以免破坏宏续行与字符串字面量。新增环境变量 TILELANG_EMIT_LINE_DIRECTIVES 作为 Python 侧的默认值，在缓存查找之前折进 pass 配置。本地 42 项通过、6 项跳过，CUDA 与 HIP 的用例在作者机器上被跳过，cuda-gdb 的实际行为尚未直接验证——作者自己写明了这一点。同日 LSP 仓跟进一笔，把 CUDA 调试对齐到新的行映射。

算子侧的节奏没有停，但方向变了。TileOPs 单日 7 笔合入里有 6 笔性能、1 笔测试基建，全部来自同一作者；每笔仍然附 H200 实测表与清单回归门槛。其中最有价值的是 #2816：新增一个只在测试时存在的 tileops_test 目标，在 PrivateUse1 上为每个算子算出共享的 workload 参考，让 tests/ops 的冒烟测试不再必须站在 CUDA 上。改动 1281 行、59 个文件，作者同时把边界写清楚了——比对共享参考只能证明后端无关的契约与可移植性，不能证明真实后端的数值正确性，那条线仍挂在 #2611。

最后是 TileFoundry 的一条暗线。#283 把「共享内存访问应当如何排序」做成分析层：一个同步族按 lowering 的发射顺序走每个函数，对每次共享内存访问取指令、发射 mesh、操作数与字节范围，在冲突访问前规划 commit/wait 与作用域内栅栏。83 个调度夹具里它规划 52 个、拒绝 31 个（24 个用 TMA 拷贝、6 个用 wgmma，两者的完成语义尚未声明；1 个需要在跨线程边界不等的循环里插栅栏）。配套的 #285 在审，把计划变成 CUDA lowering 并反向校验：sm_89 单卡上能构建的 7 个程序里 6 个与参考完全一致、synccheck 与 memcheck 全干净，而按 main 的方式 lower，同样是这 7 个里有 4 个结果错误；#283 的描述里给出的直接症状是 sm_89 的 E4M3 GEMM 最大偏 659。

一、核心项目进展

1.1 主线：昇腾 950 的持久化 FlashMLA 内核，BF16 prefill 与紧凑 FP8 decode 双路径（#3678）

合于 11-02 09:26，新增 1997 行、删除 17 行，20 个文件、42 次提交，是本窗口最重的一笔。内核面向昇腾 950，prefill 走 BF16、decode 走紧凑 FP8，可选 FP8 与 FP4 的额外 KV 缓存，并支持 attention sink、非法索引与逐 query 的独立长度。注意力循环、在线 softmax、反量化与调度全部用 TileLang 表达，靠 T.Stage 与编译器管理的缓冲版本和同步驱动；QK 与 PV 共用一块 L0B 存储环的两个不同 layout 视图，L1 落位由 T.annotate_buffer_addresses 指定；prefill 在既有累加器分配内分出四个不相交输出切片，并配独立同步事件，归一化完成一片即排空一片；decode 用整行输出拷贝。编译器侧补三处支持：L0C 到 UB 的拷贝保留 NZ 输出、每个 AIV 流在其流水线之前选定、显式调度阶段在物理发射相位打平时保持源码顺序。目标为头数 64、Dqk 与 Dv 均 512、topk 为 64 的倍数，pair gather 沿用已记录的零 burst 代码生成绕法；随附正确性与基准运行器、缓冲复用与 query 边界回归、以及带 FlashMLA MIT 声明的 README。

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

1.2 主线：cast 解耦被跳过后残留的 512 位向量，SM100 上 ViT 前向编译失败（#3734）

合于 11-02 09:05。在 e8225d7f5 上，面向 SM100 的 ViT 前向在受守卫的 FP32 转 BF16 输出存储处报 Cannot convert type float32x16 to CUDA type：规划器为 256 位 BF16 存储选 16 个 lane 并推迟 cast 约束，期待 DecoupleTypeCast 把计算与访存分离；#3579 让解耦跳过那些会读内存的守卫之后，后续 LegalizeVectorizedLoop 仍沿用同一策略继续推迟，残留的 FP32 表达式于是变成 512 位向量，CUDA 侧无法表示。修法是把 defer_cast_constraints 做成规划器的 C++ 参数（不是 IR 属性），由调用方按阶段选择：LayoutInference 与 LowerTileOp 维持 true，最终合法化一律传 false，并把算出的宽度作为显式 vectorize_hint 前传，避免重写器用默认策略重新规划。守卫的 FP32 转 BF16 降到 8 个 lane，成功解耦的 BF16 拷贝仍保留 16 个 lane，守卫安全检查保留；地址 cast 与访问指针参数从数据向量类型约束中排除，连续地址降为标量基址加步长，由既有访问分析判定合法性，int64 索引的对齐向量访问与四 lane 原子用例不受影响。

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

1.3 主线：#line 指令改为逐行重锚定，调试器映射可选密集模式（#3703）

合于 11-02 04:06，682 行新增、11 个文件。此前的 #line 发射只锚定一条 TIR 语句生成的第一行，后续行会被编译器逻辑行号推进，cuda-gdb 于是把生成的加载、运算与存储归到无关的 Python 行；合成语句与收尾大括号可能继承已推进的映射，没有源码跨度的函数还会继承前一个函数的源文件。本笔把 tl.emit_line_directives 扩成分级：0 或 False 关闭并保持默认，1 或 True 维持语句级，2 以上对每一行非空生成行重新锚定（含多行展开的语句与收尾大括号）。实现落在 C 系输出流，覆盖 CUDA、HIP、CPU C、C host、Metal 与 Ascend，未改动 vendored TVM；同时记录语句与函数的源码跨度，无源区域显式切到 <tilelang-generated>，用户原文逐字保留以避免破坏宏续行与字符串字面量。新增 TILELANG_EMIT_LINE_DIRECTIVES 作为 Python 侧默认值，在缓存查找前折进 pass 配置，显式配置优先（包括默认值非法时）。本地用 Metal 加 LLVM 构建验证：行指令与 compile-only 两组 42 项通过、6 项跳过；去掉指令后，向量化加法与条件选择两个夹具在级别 0、1、2 下产物完全一致。作者注明 CUDA 与 HIP 用例在本机被跳过，cuda-gdb 行为未直接测试。

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

1.4 主线：注意力示例按设备共享内存上限推 tile，RTX PRO 6000 上由 2 过 7 败变为 9 全过（#3702）

合于 11-02 06:59，345 行新增、13 个文件。多个注意力示例的非调优路径把 Hopper 尺寸的 tile 写死，在 sm_86、sm_89、sm_120 上每块动态共享内存上限约 99 KiB，启动直接失败：报错原文是要求把允许的动态共享内存设为 163840 字节，而设备 0 每块最多 101376 字节。现在每个示例从 shared_memory_per_block_optin 推 tile，做法沿用既有的 GQA decode 示例（example_gqa_decode.py）；数据中心配置不变，因为每条分支只在低于上一档 tile 需求时才触发。约 99 KiB 设备上的落点：attention sink 四个示例分别是 64×64 两阶段 128 线程（前向）与 64×32 一阶段 128 线程（反向），example_gqa_bwd.py 前向 64×64、反向 64×32，example_mla_decode.py 取 BLOCK_N 32、BLOCK_H 16。后两者新增 threads 参数：64 行或 16 行 tile 下 T.gemm 需要四个 warp 才能让每个 warp 独占整块 16 行，256 线程会直接 lower 失败（分别报 acc_s 与 acc_s_cast 的 layout 冲突、以及 M=16 N=32 无合法 warp 划分）。另外 flash-attn 2.x 没有 sm_12x 内核，四个 varlen 示例原本在参考步骤就抛错，现改为探测一次 flash_attn_varlen_supported，不支持则改用纯 PyTorch 的 varlen_attention_ref（覆盖 GQA、右下对齐的 causal 与 window_size）。验证在 RTX PRO 6000 Blackwell（sm_120）上：CUDA 9 项全过（原为 2 过 7 败），CuTeDSL 同样 9 项全过；H100 有 227 KiB 可选上限，仍走原 tile，同样 9 项通过。

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

1.5 主线：CUDA float32x8 位转换改用内建 bit cast，SASS 指令从 19 条降到 18 条（#3748）

合于 11-02 09:58，20 行新增、3 行删除，修 #3679 并取代 #3681。改动只用 CUDA：把 float32x8 的位重解释从位运算序列换成 __builtin_bit_cast(uint64_t, make_float2(v, v))，生成的 SASS 指令数从 19 条降到 18 条。幅度不大，但属于编译器后端里典型的「语义等价、产物更短」的一类收尾。

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

二、多后端适配（昇腾 / 海光 / 摩尔线程 / 沐曦）

2.1 昇腾：WY Fast 前向算子新增、wy_fast_bwd_split 反向分段迁移（#1827 / #1828）

两笔同在 11-02 02:33 合入，前后相差不到一分钟。#1827 以「迁移并新增」的方式把 WY Fast 算子搬到 Ascend NPU 侧，#1828 迁移并补上 wy_fast_bwd_split 的分段反向。WY 表示法是线性注意力一类模型里把递推矩阵乘化简成低秩形式的常用手段，前向与反向分段配合，等于把这条算子链在昇腾上补齐到可训练。两笔的 PR 描述只有模板问候语，改动范围以标题与提交为准，本报告不对内部实现细节作推断。

来源：https://github.com/tile-ai/tilelang-ascend/pull/1827、https://github.com/tile-ai/tilelang-ascend/pull/1828

2.2 昇腾：T.tile.fill 与 T.tile.clear 文档事实核查合入（#1634 / #1628）

同批合入的还有两笔文档侧的事实核查，分别针对 T.tile.fill 与 T.tile.clear 的 API 行为。#1628 的讨论里留了一句维护约定：这两个接口比较接近，一边有修改时请同步看看另一边是否需要对应调整。这属于持续进行的 API 事实核查批次，同期仍在审的还有 T.tile.bitwise_not（#1656）、T.tile.bitwise_and（#1651）、T.tile.bitwise_lshift（#1659）与 T.tile.bitwise_or（#1655）。

来源：https://github.com/tile-ai/tilelang-ascend/pull/1634、https://github.com/tile-ai/tilelang-ascend/pull/1628

2.3 昇腾在审：cython wrapper 缓存符号形状解析，承接宿主开销这条线（#1967）

11-02 更新。#1967 在 cython wrapper 里缓存符号形状的解析结果，目标是把每次调用的 dispatch 开销继续压下来，对应缺陷 #1898。这条线与上一期合入的 #1948 是同一目标的两个落点：#1948 把每次调用无条件构造 TVM Analyzer 的固定开销改成惰性构造（910B3 上从约 450 微秒降到约 106 微秒），#1967 处理的是符号路径上重复解析本身的成本。同族还有在审的 #3757，主线侧对齐动态 ABI 与共享符号校验：按 Var 对象而不是打印名做键，避免同名不同变量被当成同一符号，并把单例维度上的 stride 只当兜底、不让它抢先决定取值。

来源：https://github.com/tile-ai/tilelang-ascend/pull/1967、https://github.com/tile-ai/tilelang/pull/3757

2.4 昇腾在审：z3 版同步从主线迁入，整仓迁主线同步推进（#1965 / #1975）

#1965 于 11-02 12:18 更新，把主线的 z3 版 insert sync 迁到 Ascend 仓。#1975 同日 03:28 更新，仍是把 TileLang-Ascend 整体迁到主线多后端架构上。两条是同一收敛路径的两个层面：一条换同步算法的实现，一条换仓的骨架。与之配套的主线侧动作是 #3752（z3 rlimit 配置）在审，以及 #3735 提出在 thread sync 里证明模块化索引的不相交性——上一期报告里那笔 ThreadSync 回退留下的证明预算问题，正在从「配置求解资源」与「减少需要证明的形状」两头同时处理。

来源：https://github.com/tile-ai/tilelang-ascend/pull/1965、https://github.com/tile-ai/tilelang-ascend/pull/1975

2.5 昇腾在审：归约算子两个缺陷，首 Parallel 轴广播改标量读（#1978 / #1968 / #1979）

三笔都在 11-02 更新。#1978 修 AscendVidReduction 在行切分循环里对整块 [1,N] UB 错误地把 tile 算子尺寸折半（对应 #1800）；#1968 把 threads=2 时 sigmoid 与 round 的计数折半（对应 #1897）；#1979 让首 Parallel 轴上的广播操作数改走标量读，绕开 VEC 的 UB 对齐陷阱（对应 #1885）。三者都属同一类问题：昇腾侧向量化路径对形状与对齐的假设比 CUDA 侧更硬，边界条件一破就是编译期或运行期的硬错。

来源：https://github.com/tile-ai/tilelang-ascend/pull/1978、https://github.com/tile-ai/tilelang-ascend/pull/1968

2.6 海光：gfx948 按 gfx946 兼容处理，先让既有指令可测（#32）

合于 11-02 10:46。改动把 gfx948 当作 gfx946 兼容，复用既有的 MLS atoms、MMAC、barrier 与编译选项，目的是在接入 gather load 之前先让现有 builtins 能在新型号上跑通。这是典型的两步走：先把新架构号挂进现有路径拿到可测基线，再补真正需要新指令的部分。该仓上一笔常规拷贝的 LDS swizzle 修正在上一窗口合入。

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

2.7 摩尔线程、沐曦、mlir-ascend 窗口内无推送（11-03）

tilelang-musa 最后推送停在 09-30，tilelang-metax 停在 10-26，tilelang-mlir-ascend 停在 10-31 08:21（早于本窗口起点）。本期的后端动态因此集中在昇腾与海光两侧，其中昇腾占了合入 18 笔中的 4 笔、在审更新 44 笔中的多数。

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

三、生态与采用方（TileOPs / TileFoundry / TileKernels / FlashQLA）

3.1 TileOPs：单日 7 笔合入，6 笔性能、1 笔测试基建（11-02）

7 笔全部在 11-02 05:10 至 05:22 之间合入，作者同为 lcy-seso，覆盖注意力、MoE、逐元素、归一化、测试目标与 CI 契约六个方向。与上一期「一天 13 笔」相比数量减半，但结构变了：上一期是逐 shape 的调优扫尾，这一期出现了测试基建（#2816）与 CI 契约（#2818）两类非性能改动，说明算子库开始为多后端与可维护性做结构性准备。每笔仍然附 H200 实测表，并声明 manifest 里没有已测行回退。

来源：https://github.com/tile-ai/TileOPs/commits/main

3.2 TileOPs：warp 专用 decode 的 softmax 与 QK/PV 重叠，478 次扫描调用几何均值 0.952 倍（#2817）

合于 11-02 05:21。改动把 warp 专用 GQA decode 的 softmax 与 QK、PV 矩阵乘重叠起来：概率留在寄存器里，K 与 V 使用独立 barrier，共享一个三槽环形缓冲。选择条件是 D 为 128、64 个 key 的 tile、每个 KV 头一行 tile、无 softcap，且「一 CTA 一 SM」的栅格能满足占用率与流长度界；不满足的调用继续走串行程序。九个规划器边界都在宿主侧用固定设备事实校验，不需要真实分配注意力张量、也不要求 132 SM 的卡。实测环境 H200、CUDA 13.2、PyTorch 2.13.0+cu132、TileLang 0.1.12，用 CUPTI 设备忙时：形状 (16, 8, 4096, 64, 8, 128) 下 float16 为 0.074610 毫秒对 fa3 的 0.077490（1.039 倍，3.654 TB/s），bfloat16 为 0.074260 对 0.077220（1.040 倍，3.671 TB/s）；受影响的 478 次扫描调用记录的几何均值延迟为 main 的 0.952 倍。

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

3.3 TileOPs：MoE 按 token 收集预置换行，H200 上 387.8 TFLOPS（#2815）

合于 11-02 05:19。tight 布局的 MoE 输入改为按 token 收集，同一条 hidden 行在它的多条路由之间复用；token 填不满 GPU 时把路由拆到不同块上。同时删掉 tight 扫描里没人用的专家序源索引分配与写入。调度上新增一条规则：未融合、K 主序的 tight 组若平均行数大于 128 且不超过 192，选 H200 上的 ping-pong 64×256 tile；显式给出组大小的调用仍走成本模型。实测 (4096, 256, 8, 7168, 2048) 上 bfloat16 为 7.4417 毫秒对 flashinfer-cutlass 的 7.6982（1.034 倍，387.844 TFLOPS），float16 为 7.5040 对 7.9869（1.064 倍，384.624 TFLOPS）；已记录的路由与偶数组扫描覆盖两种 dtype 与偏置模式，manifest 无超噪声回退。

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

3.4 TileOPs：16 位逐元素去掉带范围检查的除法，8 字节操作数一线程一向量（#2814 / #2813）

#2814 合于 11-02 05:18，针对 bfloat16 的 Sigmoid 与 16 位 SiLUAndMul 改用四分之一缩放的倒数，绕开带范围检查的除法，并把该辅助函数与 SiLU 共用；同时省掉指数运算里用不上的次正规数守卫，融合门控输出跨线程按相邻 16 字节向量存放，并保留 gate 接近负 88 时乘上无穷值的未舍入 SiLU 激活。实测 bfloat16 Sigmoid 在 16777216 元素上 0.017504 毫秒对 torch-compile 的 0.018656（1.066 倍，3.834 TB/s），268435456 元素上 0.248545 对 0.262529（1.056 倍，4.320 TB/s）；SiLUAndMul 在 (2048, 14336) 上 0.041856 对 0.045440（1.086 倍）。#2813 合于 05:17，让 8 字节操作数在长启动下按一线程一向量处理，改动前先用 TileFoundry 的 roofline 分析确认这类按位或算子全部是访存受限：按 4.8 TB/s 的 HBM 带宽给出各 workload 的下界，hidden-state-prefill 的 int64 为 41.9 微秒、int32 为 21.0 微秒，llama-7b-ffn 的 int64 为 56.4、int32 为 28.2，cnn-feat-broadcast 的 int64 为 42.8，linear-bias 的 int64 为 14.0、int32 为 7.0。

来源：https://github.com/tile-ai/TileOPs/pull/2814、https://github.com/tile-ai/TileOPs/pull/2813

3.5 TileOPs：Norm 统计量改为减行内元素，消掉靠近 1e36 的溢出（#2811）

合于 11-02 05:10。根因写得很直接：所有均值与方差内核都对原始行做 fp32 求和，和随行的偏移量增长（约 N 乘 x ），靠近 1e36 的行直接溢出，全部输出变成 NaN；即使没溢出，均值也会在偏移量量级上舍入，x − mean 把这个舍入带进每个元素，于是常数行算出的方差等于舍入的平方，归一到正负 weight 而不是 bias。修法是给每个由行构成的 fp32 量加一条不变式：它只被行的扩散度约束，不被偏移量约束——实现上取「行内某个元素」作基准再相减，warp 主导的核用 warp 首元素、分片核用行的首元素。这条与上一期报告里 #2812 列出的四个 Norm 算子缺陷是同一类问题的正面修复。

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

3.6 TileOPs：tileops_test 目标让 tests/ops 脱离 CUDA 跑（#2816）

合于 11-02 05:20，是本窗口单笔改动最大的一笔：1281 行新增、524 行删除、59 个文件，关闭 #2803。新增一个只在测试时存在的 tileops_test 目标，在 PrivateUse1 上为每个算子计算共享的 workload 参考，tests/ops 的冒烟测试经由它运行；不支持的参考调用被显式列成 unvalidated 跳过，只在 CUDA 上成立的排除项也在同一目标下校验。算子到 workload 的注册表通过 workloads/calls 共享，基准专用的绑定、拷贝记账与输入覆盖留在 benchmarks/_cases；另补测试设备生成器、卷积支持、修正测试能力标记，并让参考实现遵守 GQA 反向的 scale 与空行语义。作者把局限写在显眼处：这个目标测的是后端无关的契约与测试可移植性，比对共享参考并不能证明真实后端的数值正确性，#2611 仍然开着；没有可用参考的调用继续以 unvalidated 形式可见。同批的 #2818 把 dispatch 契约从 workload 数据生成里拆开，为契约用例只实例化形状、不生成也不校验用不到的张量内容，并把五个宿主结构测试套件在 compile-contract-gate 里只跑一遍。

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

3.7 TileFoundry：共享内存排序做成可证明的分析层，83 个夹具计划 52 拒绝 31（#283）

合于 11-02 02:09。问题是实打实的：finalize 在共享内存访问之间不发射 commit、wait 或 barrier，于是通过共享内存暂存的调度程序会竞争，卡片上 sm_89 的 E4M3 GEMM 相对 torch 最大偏 659。#267 定下的方向是「在 HIR 上排访问顺序的分析」加「按目标发射顺序的 lowering」，并且只发最小正确形式、证不出来就拒绝；本笔交付的是分析部分。新增的同步族按 lowering 的发射顺序走每个函数：对每次共享内存访问，从 analysis/emission.py 取指令、发射 mesh 与操作数，从 IterationScope 取作用域，从内存族记录的落位取字节数；在冲突访问之前规划一次 commit 加一次 wait（针对未完成的 cp.async 写）与一次覆盖「同时持有两者的最内层作用域」的栅栏，按它先运行的子节点的活跃位置编号。算子通过新注册表 visitor_registry/completion.py 声明效果如何完成，什么都没声明的按 UNKNOWN 处理；该族拒绝 UNKNOWN 之后的访问、返回时仍在飞行中的 UNKNOWN 访问、在跨线程边界不等的循环里插栅栏、以及被保留调用且被调方持有共享内存的情形。默认没有调用方请求该族，因此 75 个会 finalize 的调度夹具字节与 main 完全相同；83 个夹具里计划 52、拒绝 31（24 个用 TMA 拷贝、6 个用 wgmma，尚未声明完成语义；1 个需要在线程相关边界的循环里插栅栏），语料用例逐条断言了这些拒绝。

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

3.8 TileFoundry 在审：CUDA lowering 发出同步计划并反向校验，main 上 7 个程序错 4 个（#285）

11-02 12:40 更新。这笔是 #283 的执行端：CUDA lowering 按同步计划发射步骤，并逐条检查实际发射与计划一致，不一致抛 LoweringError。语料上的结果：默认流程为 14 个程序加上同步、拒绝 31 个；这 14 个的 TIR 与 main 的差别只有新增语句，另外 30 个与 main 完全相同；经核心 visitor 后，main 会 finalize 的 75 个程序字节一致。单张 sm_89 上跑了 7 个能构建且带同步的程序：6 个与参考完全一致，且在 synccheck 与 memcheck 下全部干净；同样这 7 个按 main 的方式 lower，有 4 个结果错误。作者同时报出 main 已有的两个问题：GEMM 与 loop_outputs 上 racecheck 报同一 cp.async 的多个 lane 写同一地址（HIR 里的 copy_async 不接受拆分目标，于是每个线程拷整块 tile），loop_outputs 还会把一个输出写到错误的元素上。

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

3.9 采用方：FlashQLA、TileKernels 窗口内无推送（11-03）

QwenLM/FlashQLA 最后推送停在 10-30 03:02，deepseek-ai/TileKernels 停在 10-14 05:49，两仓在本窗口都没有提交。TileKernels 自 10-14 起已连续多期无推送，采用方的可见活动本期全部落在 TileOPs 与 TileFoundry 两个自家仓里。

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

四、社区、教程与活动

4.1 昇腾同步正确性成为集中报障面：单日 8 条新 issue，两名报告人分别系统性提交（11-02）

窗口内主仓新增 8 条 issue，编号集中在 #3740 到 #3749，全部指向昇腾 AscendC 与 PTO 两个后端的代码生成与同步插入。其中 7 条由同一报告人提交，环境为 TileLang 0.1.15、PyTorch 2.10.0+cpu、torch_npu 2.10.0，主题是同一段程序在两个后端上的行为分歧：生成代码在启动哪些核上不一致（AscendC 出 mix(1,2)、PTO 出 vector）、enable_auto_schedule=True 时 AscendC 不插同步、关闭自动调度时生成成功而开启后失败且报错信息不清、同一份代码 AscendC 能生成而 PTO 失败。这一组报告的共同点是都附了可复现的 prim_func 与逐行对照的生成代码，属于可直接转成回归用例的形态。加上在审的 #3735（thread sync 里证明模块化索引不相交）与 #3580（谓词与提升索引同步），昇腾侧的同步正确性已经成为当前最集中的一条线。

来源：https://github.com/tile-ai/tilelang/issues/3747、https://github.com/tile-ai/tilelang/issues/3745

4.2 跨行与转置视图的 SIMD 加载丢掉 DMA 依赖，根因落在逻辑区域重建（#3749）

同批里技术含量最高的一条，报告人环境为 main 分支 e8225d7f5、Ascend950PR、CANN 9.2.0。问题是两处连续 SIMD 加载丢掉了对更早一次 GM 到 UB 拷贝的依赖：一是行主序加载从行中间开始并跨到下一行，二是访问转置视图且物理连续维不是最后一个逻辑维。根因路径写得很细：生成的 tl.access_ptr 覆盖物理加载范围，而 MemoryAccessDetector::LogicalPointerRegion 重建出的逻辑区域漏掉了其中一部分，依赖分析于是认为输入 DMA 与加载无关，AutoSchedule 把输入拷贝挪到加载之后，InsertSync 也不发 MTE2 到 V 的 set/wait。最小复现里 work.shape 为 (2, 128)，先 T.copy(A, work[1, 0:64])，再在 T.SimdVF() 里 T.simd.vld(work[0, 64]) 读 128 个连续半精度元素：生成的读描述符是 tl.access_ptr(work[0, 64], 128, 1)，物理覆盖 [64, 192)，其中 [128, 192) 正是刚被 DMA 写过的部分，而重建出的区域是 work[0:1, 0:128] 即物理 [0, 128)，完全错过被写区间；行数按 ceil(128/128) 算成 1 行，起始列 64 没被计入。这条与本窗口主力在补的 FlashMLA、PTO 向量化是同一块地基，一旦命中就是静默错结果。

来源：https://github.com/tile-ai/tilelang/issues/3749

4.3 新闻侧：Google News 中英文、Hacker News、arXiv 三通道均无窗口内命中（11-03）

六组 Google News 查询（英文 TileLang、TileLang GPU kernel、tile-ai TileLang；中文 TileLang 算子、TileLang 昇腾、沐曦或海光或摩尔线程）返回的最新相关报道是 10-21 的 SemiAnalysis 华为与 DeepSeek 一文，其余集中在 09-29 至 10-01 的 DeepSeek 开源昇腾组件那一批，全部早于本窗口。arXiv 上 all:"TileLang" 的最新命中为 2026-07-24 的 TileSight，TileLang 与 Ascend 的联合查询无结果；Hacker News 的候选条目里与 TileLang 相关的只有仓库推送条目。候选清单 14 条中 12 条来自 GitHub 仓库发现，2 条是无关的 Show HN 与 LSP 仓，本期信息面完全由仓库侧构成。

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

五、趋势观察

这一天最值得记的判断是：昇腾不再是「适配目标」，而是「首发目标」。FlashMLA 这种 DeepSeek 自家解码路径上的核心内核，直接以昇腾 950 为目标写在主线里，用的还是 T.Stage、annotate_buffer_addresses 这类为显式调度准备的底层能力，而不是另开一个仓做移植。配合编译器侧那三处支持（L0C 到 UB 保留 NZ、AIV 流在流水线前选定、显式阶段按源码顺序发射），可以看出昇腾路径正在从「能跑」转向「能按调度意图跑」。

第二条线索是正确性的重心在下移。上一期主线为 TIRX 迁移回退了一整套 ThreadSync 正确性改进，这一期立刻出现两条补位：主线 #3752 调 z3 求解资源上限，#3735 提出用模块化索引的不相交性减少需要证明的形状；同时 #3749 把一处「依赖分析看不见物理加载范围」的缺陷挖到了 MemoryAccessDetector::LogicalPointerRegion 这一层。把这些和 TileFoundry 那条暗线放在一起看——分析层能证明 83 个夹具里的 52 个共享内存排序，而 lowering 按 main 的方式做会让 7 个程序里的 4 个算错——同一个信号出现了三次：TileLang 现在的主要风险不是功能不够，而是「生成的代码看起来对、实际不对」，而团队选择的做法是把正确性做成可拒绝的证明，而不是靠补丁逐个堵。

第三条是测试与契约开始独立成层。TileOPs 的 tileops_test 把冒烟测试从 CUDA 上摘下来，compile-contract-gate 把 dispatch 契约与 workload 数据生成拆开，两者都在同一小时内合入；作者同时明确写了「共享参考不能证明真实后端正确」。这种把「能测什么」和「证明了什么」分开表述的习惯，比多合几笔性能补丁更能说明这个项目正在走出原型期。

最后注意一个反差：算子侧单日 7 笔性能合入、6 笔都带 H200 实测表格，而新闻侧三通道 24 小时零命中。TileLang 的对外声量仍然集中在「DeepSeek 开源昇腾组件」这类事件上，日常进展几乎只发生在仓库里。对跟踪这个项目的人来说，这意味着日报的主体只能是提交与 PR，而不是报道。

附、素材与核查说明

信源核查表

信源 核查结果
tile-ai 组织（29 仓库） 窗口内推送 9 个仓库
主仓 tilelang 主分支 5 笔合入；新开 issue 8 条；文档站 5 笔自动构建
TileOPs 主分支 7 笔合入；站点 1 笔
TileFoundry 主分支 1 笔合入；1 笔在审新增
tilelang-ascend 主分支 4 笔合入；6 笔在审更新
tilelang-hygon 1 笔合入
tilelang-lsp 1 笔合入（对齐 CUDA 调试行映射）
TileOPs-nightly 1 次自动推送
tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack 窗口内无推送
发布页（tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon） 窗口内无新版本
FlashQLA / TileKernels 窗口内无推送
Google News（中英文六组查询）/ Hacker News / arXiv 无窗口内实质命中

完整信源清单