调研窗口：2026-10-21 07:00 ~ 2026-10-22 07:00（北京时间，24 小时）。

信源：GitHub tile-ai 组织逐仓核查（Atom 提交流，覆盖 tilelang / TileOPs / TileFoundry / tilelang-ascend〔ascendc_pto〕/ tilelang-hygon〔main 与 feature 分支〕/ tilelang-metax〔dev〕/ tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack）+ 合入与在审真值检索（is:merged / is:open 两组，按仓库分组统计）+ 发布页 tag 核查 + 文档站提交核查 + 采用方仓核查（TileKernels / FlashQLA）+ 社区新仓检索 + Google News 中英文多组查询（走代理）+ Hacker News + arXiv + 第三方产业分析报告。

本期索引

  • 今日重点：TileFoundry 0.0.3 正式发布，把指令级调度交给 agent；TileOPs 一天合入 20 笔，MLA 反向与 SM89 小规模路径同日补齐；主仓八笔合入集中在”不报错但结果不可信”的编译器缺陷；昇腾 DeepGEMM 对齐由在审转合入（10-21）
  • 一、核心项目进展
    • 1.1 主仓：批式 finalize_reducer 此前只归约首个输出块（#3070）（10-22）
    • 1.2 主仓：BF16→FP32 向量化转换改走内联 PTX 打包 lane 转换（#3595）（10-21）
    • 1.3 主仓：展开 pass 保留空 stepped 循环，ROCm 的 Select 改逐 lane 下降（#2887）（10-21）
    • 1.4 主仓：软件流水保留传递性拷贝依赖，线程域参与区域冲突判定（#2822）（10-21）
    • 1.5 主仓：cast 的源/目标/舍入三组合前置校验，拒绝静默丢注解（#3350）（10-21）
    • 1.6 主仓：无 arch 的 CUDA target 拒绝而非静默填 sm_50（#2860）（10-21）
    • 1.7 在审：移位符号性与 TIRx 循环作用域修复在 tvm fork 与主仓成对提交（#3597 / tvm #87 / #86）（10-21）
  • 二、多后端适配（昇腾 / 沐曦 / 海光 / 摩尔线程等）
    • 2.1 昇腾评测：SemiAnalysis 深度报告设 TileLang Ascend 专章，称六仓中当时只有 TileLang 可用（10-21）
    • 2.2 主仓昇腾：DeepGEMM 启发式与 epilogue 对齐由在审转合入（#3577）（10-21）
    • 2.3 主仓昇腾：warp_reduce_max/min 在 PTO 上的整数符号性修复（#3582）（10-21）
    • 2.4 昇腾仓：亚 32B 紧凑行距源无法表达，keepdim reduce 输出错乱（#1795）（10-21）
    • 2.5 昇腾仓：JIT 缓存改原子发布，并在缓存被删除时恢复（#1903）（10-21）
    • 2.6 昇腾仓：cann-bench 补 depthwise_conv_2d 示例（#1915）（10-21）
    • 2.7 昇腾在审：indexer_topk_reducesum、GLM-5.3 kpool compress 与 WY fast 系列算子迁移（10-21）
    • 2.8 海光：feature 分支推进 FP8 解码 GEMM，定位到地址链占 kernel 13%（10-21）
    • 2.9 沐曦 / 摩尔线程 / MLIR-Ascend：窗口内无推送，三条线停在既有版本（10-21）
  • 三、生态与采用方（TileOPs / TileFoundry / 算子库）
    • 3.1 TileFoundry：0.0.3 打 tag 发布，指令级调度 + 重构分析 + H200 与 Apple AMX 示例（10-21）
    • 3.2 TileOPs：MLA 预填充反向补齐 op 层与 SM90 warp 特化 kernel（#2678 / #2680 / #2688 / #2686）（10-22）
    • 3.3 TileOPs：dense GQA 支持 FP32 attention sink，-inf 关闭、+inf 清零该 head（#2515）（10-21）
    • 3.4 TileOPs：SM89 小规模路径集中补位，FP8 BMM 默认瓦片降到 64×64（#2681 / #2672 / #2669 / #2654）（10-21）
    • 3.5 TileOPs：NSA top-k 在 SM89 上扩到 2048 压缩块并补齐输出槽（#2674 / #2675 / #2673）（10-21）
    • 3.6 TileOPs：GDN 前向在 walk 填满 SM 时单 kernel 完成，分量化归一化按 cluster 切宽行（#2682 / #2679）（10-21）
    • 3.7 TileOPs：GQA 反向在 warp 特化 kernel 里直接服务分组 KV head（#2685）（10-21）
    • 3.8 TileOPs 在审：折叠 edge 轴单次发射、NSA top-k 收窄列 tile（#2689 / #2690）（10-22）
    • 3.9 采用方：TileKernels 与 FlashQLA 窗口内无推送（10-21）
  • 四、社区、教程与活动
    • 4.1 社区教程：中文《TVM / TileLang Cookbook》更新，从 jit 讲到 cubin（10-21）
    • 4.2 社区新仓：tilelang-wrapper-for-sail 建仓，暂无说明（10-21）
    • 4.3 社区工具：远程自动调优 agent 仓继续推送（10-21）
    • 4.4 文档与版本：文档站七次自动更新，主仓 VERSION 仍 0.1.15（10-22）
  • 五、趋势观察

今日重点

本窗口合入 32 笔（TileOPs 20、tilelang 8、tilelang-ascend 3、TileOPs 文档站 1），另有一项正式发布：TileFoundry 于 10-21 10:26 打出 0.0.3。三条线的重心各不相同。主仓八笔合入几乎全是”不报错但结果不可信”的缺陷清理——批式归约只处理首个输出块、cast 舍入注解被静默丢弃、无 arch 的 CUDA target 静默退到 sm_50、流水规划丢失传递性依赖；TileOPs 一天把 MLA 预填充反向从 op 层做到 SM90 warp 特化 kernel，同时把 SM89 上的 FP8 BMM、W4A16、分页 DSA、NSA top-k 四条小规模路径一并推进，在审队列从 16 笔降到 9 笔；昇腾侧昨日仍在审的 DeepGEMM 启发式与 epilogue 对齐（#3577）今日合入，昇腾仓则修掉一个会造成 keepdim reduce 输出错乱的内存访问建模缺陷（#1795）。产业侧，SemiAnalysis 的华为-DeepSeek 深度报告专门写了一节 TileLang Ascend，给出的判断是六个获得昇腾支持的仓库中当时只有 TileLang 可用。

一、核心项目进展

1.1 主仓：批式 finalize_reducer 此前只归约首个输出块（#3070）

合于 10-22 01:16，2 文件 +29/-5，修 #2623。批式 finalize_reducer 的下降只在缓冲区开头调用一次 run_batch，其余连续输出块根本没有被归约：输出维度大于 batch 时，后面各段拿到的不是归约结果，而是未初始化或上一轮遗留的值。修复按 ceil(OutputDim / batch) 逐块发射带正确偏移的批式 all-reduce——128 个输出、batch=4 时发射 32 次调用，同时把重映射 reducer 缓冲改用扁平视图、每批以常量线性偏移寻址，去掉原先的形状与步长重建，并保留全部 chunk 调用。验证在 L40S 上完成：新增 sum/max/min 数值回归，覆盖 fp16 与 fp32、batch 4/8/16，23 项 finalize_reducer 相关用例通过，外部复现矩阵 16 例零失配。属典型的”小尺寸正确、大尺寸静默错”的一类缺陷。

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

1.2 主仓：BF16→FP32 向量化转换改走内联 PTX 打包 lane 转换（#3595）

合于 10-21 19:28，4 文件 +19/-7。
新增 tl::__tl_cvt_bf16x2_to_f32x2，用内联 PTX 把成对的 BF16 lane 一次转成 f32x2，替掉逐 lane 的标量提升路径；2、4、8 三种 lane 宽度均保持支持。同步更新了向量化 cast 与布局推断的期望值，并把 BF16 cast 的测试收敛到既有覆盖范围，不扩大断言面。改动不涉及语义，只动生成的指令序列，影响面是所有走向量化 cast 的 BF16 计算路径。

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

1.3 主仓：展开 pass 保留空 stepped 循环，ROCm 的 Select 改逐 lane 下降（#2887）

合于 10-21 19:20。循环展开 pass 此前会把 trip count 为 0 的 stepped 循环当成空壳消掉，改变后续 pass 看到的程序结构；修复改为在展开过程中保留这类空循环。同笔还改了 ROCm 侧的向量化 Select：改为逐 lane 下降，避免整组下降在不支持的目标上失败。两处改动一在 IR 结构、一在后端下降，合在同一笔里，说明它们由同一个复现场景引出。

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

1.4 主仓：软件流水保留传递性拷贝依赖，线程域参与区域冲突判定（#2822）

合于 10-21 18:13。流水规划此前只跟直接依赖，间接（传递性）的拷贝依赖会在规划阶段丢失，dependent copy 的等待随之被删；同时物理区域冲突检查在规划与下降两侧各有一套口径。修复用一个后向 worklist 同时跟随标量与 buffer 拷贝依赖，让规划与下降共享同一份物理区域冲突检查，保留重放绑定在作用域内，并补齐传递生产者、strided alias 与顺序 TMA 回落路径的覆盖。第二笔提交把线程域引入区域冲突判定：访问先展开到物化线程域再比较区域，使互不相交的共享内存切片仍可流水，同时保住跨线程依赖与 strided 缓冲的物理重叠；同一套域用于流水规划、依赖校验与 buffer 版本化。

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

1.5 主仓：cast 的源/目标/舍入三组合前置校验（#3350）

合于 10-21 13:49。round 只在显式列出的组合下可下降：PTX cvt.rs 要求 f32 源、32 位 rbits 操作数与特定的打包目标。其余组合此前在分支分派时被放过，舍入注解就此静默消失；标量循环本来就会拒绝这些组合，向量化循环不拒绝，于是同一个 cast 在两种循环形态下行为不同。典型例子是 f32→E8M0：它的 helper 向正无穷舍入、根本不接受舍入参数，而向量化分支从未读过 round。现在在分派之前统一拒绝非法组合，任何分支都无法再丢注解；配套测试只给真正接受 rbits 的 rs 模式传该参数。

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

1.6 主仓：无 arch 的 CUDA target 拒绝而非静默填 sm_50（#2860）

合于 10-21 13:32，1 文件，修 #1049。在有 CUDA toolkit 但没有可用 GPU 的机器上，未指定 arch 的 CUDA target 会退化成裸 cuda，TVM 随后静默填入 arch=sm_50，编译继续推进到一个错误的 SM，全程没有任何诊断——这正是 tilelang.compile(func) 配 target="auto" 的陷阱。修复落在 _cuda_target_from_arch：自动检测与显式 target="cuda" 共用的唯一入口，改为抛出可操作的错误，并在错误里给出显式 arch 的两种写法（{"kind":"cuda","arch":"sm_90"} 或 TILELANG_DEFAULT_TARGET）。headless 编译只要给了具体 arch 依然可用，因为 codegen 与 nvcc 只依赖 toolkit，设备访问推迟到 kernel 启动。PR 过程中一度加入三个无 GPU 单测，最终提交把它们撤下，改动只剩入口校验。

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

1.7 在审：移位符号性与 TIRx 循环作用域修复成对提交

主仓 #3597 修混合符号移位时丢失符号性的问题，同期在 tile-ai/tvm fork 提交 #87「TIRx 保留移位左操作数类型」，两者是同一处逻辑在语言层与 TIRx 层的对应改动。另有 tvm #86 把上游 apache/tvm#19927 的循环作用域修复回移。fork 在审 25 笔、主仓在审 163 笔，TIRx 侧的先行使 fork 承担了部分改动的前置验证，主仓再跟进，这条协作路径本窗口第二次出现。

来源：https://github.com/tile-ai/tilelang/pull/3597、https://github.com/tile-ai/tvm/pull/87、https://github.com/tile-ai/tvm/pull/86

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

2.1 昇腾评测：SemiAnalysis 深度报告设 TileLang Ascend 专章

SemiAnalysis 于 10-21 发布《The Huawei-DeepSeek Symbiosis》，副题直接列出「华为 Ascend 950PR 深挖与微基准、DSv4.1 Flash 分析、DeepSeek 的昇腾支持、AICPU、TileLang Ascend」。与 TileLang 直接相关的三点：其一，获得昇腾支持的六个仓库（TileLang、DeepGEMM、DeepEP、TileKernels、FlashMLA、DeepSelect）中，文章写作时只有 TileLang 可用，因为 CANN 9.2.0 尚未发布，而 DeepSeek 要求 CANN 9.2.0 与 C++20 编译器，9.2.0-beta.2 不提供后者；其二，昇腾 950 后端已并入 TileLang 主仓、而非只存在于单独的昇腾适配仓，提供原生编译、调度、同步与运行时集成，AutoSchedule 依据数据依赖与预估执行时间决定重叠与片上缓冲槽数，在 950PR 上对小尺寸单算子最多缩短 40%，编译器还会自动插同步、删掉多余等待，并检查条件分支；其三，与 CANN、Triton-ascend 对比时「某些尺寸可以打平或超过 CANN」，最优实现随尺寸切换（作者注明这是热缓存下的朴素基准）。文章同时点出生态缺口：多数优化面向超节点，x86 主机上 vLLM-Ascend 0.26.0rc1 仍有多处问题。

来源：https://newsletter.semianalysis.com/p/huawei-deepseek

2.2 主仓昇腾：DeepGEMM 启发式与 epilogue 对齐由在审转合入（#3577）

合于 10-21 10:20，29 文件 +1392/-1036。上一期记为在审，本期完成合入：补齐 DeepGEMM-Ascend 的配置选择与 epilogue，覆盖密集、批量与分组 GEMM 的融合 bias、Alpha、BF16 随机舍入与 FP8 量化，并在作者自己的提交里删掉冗余 setup、恢复测试门控。这笔与 2.1 的外部判断互相印证——昇腾侧反复打磨的正是 GEMM 路径的启发式与数值收尾，而这恰是「可用」与「跑得好」之间的差距所在。

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

2.3 主仓昇腾：warp_reduce_max/min 在 PTO 上的整数符号性修复（#3582）

合于 10-21 10:21。PTO 的 pto.redux_max/min 拒绝带类型的整数：PTOAS 要求无符号性载体加一个显式 signedness，于是 int32/uint32 的 T.warp_reduce_max/min 直接编译失败，报「requires signedness=’signed’ or ‘unsigned’ for integer values」。修复把 warp_reduce_sum/max/min 统一走 contrib.ptodsl 的 tl.redux_* 包装（现已导出），由包装提供载体与符号性，并按 AscendAllReduce 的方式对 value dtype 设门；问题单引用了主仓 #3541 与 PTOAS #1579。验证分两层：op × dtype 网格的 lower-only codegen 测试，以及稀疏的双目标（ascend/pto）执行测试——用 ≥2^31 的 uint32 lane 把符号性写错直接暴露出来。

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

2.4 昇腾仓：亚 32B 紧凑行距源无法表达，keepdim reduce 输出错乱（#1795）

合于 10-21 17:15，2 文件 +302/-3，修 #1682，是本窗口技术含量最高的一笔。根因在模板层：copy_ub_to_gm 无条件生成一条多 burst 的 DataCopyPad，而 MTE3 引擎实现的是系统级「UB 窄行按 32B 槽存放」约定——源地址每 burst 前进 ceil(blockLen/32) + srcStride 个 32B 块。对按 32B 边界写行的生产者这是精确匹配；但 T.reduce_max 的 (M,1) fp32 输出是 V-pipe 紧凑写入、行距只有 4B，引擎无法表达。实测第二个 burst 读到的是 element 8 而不是 element 1（[127, 1151, 71, 199] 里 1151 恰是 element 8 的值），写入侧没问题，错在 DMA 读源的地址序列。修复分两半：两侧皆 packed 且 total ≤ 65535 时折叠成一条 flat 单 burst（M×L 字节，利用「单条 burst 长度任意」的契约）；不可表达的非 32B 倍数行距改为 trap() 显式中止，用编译期 if constexpr 守卫，对 32B 倍数与亚 32B 槽布局不生成额外代码。为什么不降级：bisheng 的跨管道自动同步只覆盖直线 intrinsics，用条件分支包裹 DataCopyPad、或兄弟分支里放 GM 标量写，都会让 bisheng 生成坏的同步代码并导致设备死锁（实测逐行 DMA 被运行时分支包裹即触发 aicore 异常）。14 个用例覆盖 flat/multi 路径、keepdim reduce 复现、UB 循环复用、不可表达 pitch 的合法单 burst 形式与显式中止；测试环境为 Ascend910 A3（dav-c220）/ CANN 9.0.0 本地构建。PR 自述保留两个边界：亚 32B 紧凑行距且带行距目的（如把 (M,1) fp32 写进宽张量的某一列）仍走槽步进；blockCount > 4095（12 位字段）仍会静默截断（4200 行实测只拷前 104 行），拟另立 issue。

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

2.5 昇腾仓：JIT 缓存改原子发布，并在缓存被删除时恢复（#1903）

合于 10-21 16:20，7 文件 +1026/-86，关联 #1836。原实现把编译产物原地覆盖到 kernel_lib.so，其他进程若已加载该共享库，就会读到被改写中的文件。修复先把产物复制到同目录下的独立临时文件，再用 os.replace() 原子发布；加载前若缓存库已被删除，经过错误消息与文件状态确认后返回 cache miss，重新编译；保存阶段目录消失则跳过保存、保留已经得到的编译或调优结果；权限错误、依赖库缺失、无效 ELF 与序列化内部异常继续向上传播，不被吞成 miss。覆盖面包括 KernelCache、AutoTunerCache 与 AutotuneResult，不改 cache key、算子计算逻辑与 CI 并行度。41 项 host-only 用例通过（原子发布 9 项、缓存删除恢复 32 项），PR 自述未做 NPU 侧集成验证。原子发布部分沿用 #1853 的实现并保留原作者信息。

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

2.6 昇腾仓：cann-bench 补 depthwise_conv_2d 示例（#1915）

合于 10-21 16:19。examples/cann-bench/ 新增 depthwise 二维卷积基准。附带的冒烟测试细节值得记一笔：bench_test.sh 判定 PASSED 的标准是输出里出现 KERNEL OUTPUT MATCH 或 TEST PASSED!，脚本原先打印的是「smoke test passed」，退出码为 0 却被判失败；修复按 LSTM/GRU 的冒烟测试对齐，避免新增用例在 CI 里持续假红。

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

2.7 昇腾在审：indexer_topk_reducesum、GLM-5.3 kpool compress 与 WY fast 系列迁移

昇腾仓的在审队列本窗口有 13 笔更新，集中在三组。#1917 实现 indexer_topk_reducesum 算子，对应稀疏注意力的索引器取 top-k 后求和；#1910 增加 GLM-5.3 的 kpool 压缩算子，指向新一代模型的 KV 压缩路径；#1827 / #1828 / #1802 / #1803 / #1799 是一组 WY fast 与 chunk 系列算子迁移（WY fast 本体及其反向、chunk_o 反向、Chunk Scaled Dot KKT、Chunk Delta H），明显对着线性注意力/门控 delta 网络的训练链。另有 #1785 修多维标量 BufferLoad 在 T.tile 二元运算里的索引、#1854 校验自动跨核同步配置、#1731 修 gemm_v0 的 W4A4 int4×int4 结果错误（#438）。

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

2.8 海光：feature 分支推进 FP8 解码 GEMM，定位到地址链占 kernel 13%

tilelang-hygon 的 main 分支窗口内没有新提交（最新仍是 10-20 合入的 #25），但 feat/hcu-smls-b-vgpr 分支在 10-21 11:04 与 15:27 有三次推送，内容是可复用的性能定位与调参结论。11:04 那笔把解码 GEMM 的 k-loop 展开，并让 M>96 的 tile 回到 bn96；作者的关键判断是「每步固定开销来自地址链而不是 barrier」——符号化共享级索引 (k+d) % 3 会下降成 v_mul_lo_u16 / v_sub_u16 / v_readfirstlane_b32 并落在 matrix_load 的地址路径上，占掉 kernel 13%，而 T.Pipelined 本来就会发同样的两个 barrier。15:27 那笔把 M∈(64,128] 的 wide-N 中段优先按 block 数切分，gate_up 在 M∈(64,96] 取 bm96/bn128/bk128/ns2，并把 builder 移进包内的 decode_cfg；同时把解码 GEMM 的记录文档从 3825 行整理成 588 行的「结果 + 优化日志」两段式，另附已关闭杠杆索引与带宽模型。

来源：https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-smls-b-vgpr

2.9 沐曦 / 摩尔线程 / MLIR-Ascend：窗口内无推送

tilelang-metax（默认分支 dev，最新推送 10-16，发布仍停在 09-17 的 v0.1.14）、tilelang-musa（最新 09-30，v0.1.15+musa.1）、tilelang-mlir-ascend（最新 10-19，tag 停在 09-09 的 v0.1.2.020.sglang.poc）三仓本窗口均无提交，也无新 tag。

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

三、生态与采用方（TileOPs / TileFoundry / 算子库）

3.1 TileFoundry：0.0.3 打 tag 发布，指令级调度 + 重构分析 + H200 与 Apple AMX 示例

上一期记录「0.0.3 公告已进 README、尚待打 tag」，本期 10-21 10:26 正式发布，PyPI 为 tilefoundry==0.0.3，代号「Agent-driven scheduling」。核心变化是调度权归属：编译器列出每个 HIR 操作可由哪些指令实现、各自需要什么，agent 在候选中选定，编译器再把选定的调度下降为 TIR。配套 CLI 有 tilefoundry schedule candidates 与 schedule facts 展示候选与约束，tf.schedule 记录选择、tilefoundry schedule finalize 完成下降；指令以 AtomSched 操作声明并带显式 TIR 操作数契约，从 CUDA MMA atom、SM90 WGMMA 与异步拷贝起步。分析部分重构为内存（分段分配、footprint 与复用窗口、循环感知的 buffer 生命周期）、成本（按执行作用域计工作量与流量、每条访问关系声明读什么）、结构（定义事实、估值上下文与作用域快照统一走 visit context，投影复用使每个语料用例只分析一次）。新教程沿一个 8192×5120 乘 5120×17408 的 GEMM 一路走到 AtomSched 在 H200 上的最优 kernel。

性能面给出三组真实数字：Qwen3.8-27B 预填充在单张 H200 上，1024 / 4096 / 8192 token 分别 80.29 / 309.61 / 631.19 ms，SGLang 为 89.06 / 317.11 / 638.25 ms；短序列 128 / 256 / 512 / 1024 token 分别是 SGLang 的 1.88× / 1.41× / 1.10× / 1.03×；Apple M2 Pro 上的 Qwen 预填充示例把大投影做成 FP16 AMX 外积。两组 Qwen3.8-27B 运行在每个长度上都与 transformers 产生相同首 token。已知问题两项：wheel 尚未打包 TileFoundry 的 CUDA 运行时头文件，从安装包出发 tf.compile 编不出 CUDA kernel（需源码 checkout，下个版本修），以及性能时间线对 insert_slice / tf.zeros 跨 mesh 位置的定价有误且不显示 mesh lane 欠填（#239）。

来源：https://github.com/tile-ai/TileFoundry/releases/tag/v0.0.3

3.2 TileOPs：MLA 预填充反向补齐 op 层与 SM90 warp 特化 kernel（#2678 / #2680 / #2688 / #2686）

四笔连成一条线。#2678（10-21 11:34）加上 MLADenseBwdOp 与 MLAVarlenBwdOp 两个 op 层，对齐 FA3 的 _flash_attn_backward 接口；#2680（10-21 20:58）补 SM90 kernel，针对 D_qk=192、D_v=128 给出 MLADenseBwdWSKernel 与 MLAVarlenBwdWSKernel 两套 warp 特化实现；#2688（10-21 23:28）改 MLAVarlenPrefillWSFwdKernel 的常驻 CTA 取数顺序——此前先给每个 head 领一个 query tile，首个 tile 走完之后只能按 head 逐个排队，现在改成一次只走同一个 head 的 query tile；#2686（10-22 03:08）做收尾清理，把单读者常量与 helper 移到使用处，不改行为。文档站 #78 同步把两个新 op 写进 Attention 页，起因是 check_api_pages.py 报主仓导出 207 个 op 中有两页缺失。

来源：https://github.com/tile-ai/TileOPs/pull/2678、https://github.com/tile-ai/TileOPs/pull/2680、https://github.com/tile-ai/TileOPs/pull/2688

3.3 TileOPs：dense GQA 支持 FP32 attention sink（#2515）

合于 10-21 22:42。为 dense GQA 的预填充与解码增加可选的逐 head FP32 attention sink，覆盖 SM90 上的 FP16/BF16/FP8、RoPE、softcap 与滑动窗口；其他架构对 sink 调用直接拒绝。语义上每个 sink 是分母里一个不缩放、不加 softcap 的 logit：-inf 关闭该 sink、+inf 直接把该 head 输出清零，split decode 在大 logit 下也要保住 sink 权重。顺带修了一处基准失真——分块预填充的基准此前被当成单 query 解码在测；融合 RoPE 的尾部用例把 KV 长度调到 271/1025，以便覆盖拆分执行。

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

3.4 TileOPs：SM89 小规模路径集中补位（#2681 / #2672 / #2669 / #2654）

四笔合入都在补 L40S/4090 级的短板。#2681 把 FP8 BMM 在 SM89 上的默认瓦片改成 64×64——条件是结果 grid 每 SM 至多一块，大 grid 仍保留 128×128；作者给出 PyTorch 2.13.0 + CUDA 13.0 + TileLang 0.1.12 的单卡实测，1×16×64×128 从 5.161 降到 3.799 us、16×128×128×1024 从 12.943 降到 7.414 us，4096³ 反而从 512.573 变为 583.680 us（该行瓦片未变，属噪声）。#2672 把 64 列 W4A16 默认的 M 上界从 16 扩到 64，M=17 / 32 / 64 的重放时间分别下降 43.35% / 47.20% / 40.00%，同时保留显式配置覆盖与数值容差。#2669 用 MMA 实现 SM89 的分页 DSA，保留 656 字节 FP8 latent cache 布局、BF16 舍入、自然对数 LSE 与 dead position 行为：batch 4 / 128 head / 256 位置 0.269 ms 对 PyTorch FP32 参考 0.727 ms，48×64 head×2048 位置 2.543 对 2.535 ms，作者明确写「大用例基本打平，不构成普遍优势」。#2654 针对小 token 数只计算并搬运覆盖的 8-token 累加器前缀，在 132 SM 的 H200、N 可被 128 整除、133–225 输出 tile、K∈{7168, 7680, 8192} 的条件下把 40-token 解码打开到 M33–40。

来源：https://github.com/tile-ai/TileOPs/pull/2681、https://github.com/tile-ai/TileOPs/pull/2672、https://github.com/tile-ai/TileOPs/pull/2669、https://github.com/tile-ai/TileOPs/pull/2654

3.5 TileOPs：NSA top-k 在 SM89 上扩到 2048 压缩块并补齐输出槽（#2674 / #2675 / #2673）

三笔围绕同一处 SM89 路径。#2674 让 group48 复用既有 tiled top-k 程序：一个 token 一个 tile（48 个 query head 行）、列向 GEMM warp 布局；#2675 把 SM89 上的 top-k 选择从原来的上限扩到 1024–2048 个压缩块，覆盖 head dim 64 / 128 / 256；#2673 修候选不足时的输出填充——请求的每个 top-k 槽都要填 -1，此前 32 槽的排序带会让 32–63 号槽留着旧值。另有 #2684 给打包请求的 NSA 排序容量加上界，按 min(total_chunks, ceil(total_tokens / block_size)) 再对齐到打分 tile。

来源：https://github.com/tile-ai/TileOPs/pull/2674、https://github.com/tile-ai/TileOPs/pull/2675、https://github.com/tile-ai/TileOPs/pull/2673

3.6 TileOPs：GDN 单 kernel 预填充与分量化归一化切分（#2682 / #2679）

#2682 直指线性注意力的 GDN 前向：其预填充在夜间跑批的 79 行上慢于 FlashInfer（0.57×–0.98×），改法是在「walk 能填满 SM」的情形下把预填充收进单个 kernel。#2679 处理 FusedAddRMSNormPerTokenQuant：单元素约 25 条 SASS 指令花在逐元素数学上，改为把元素数学打包，并把宽行跨 cluster 切分。两笔都是把「单算子内串行开销」换成「更宽的并行与更少的指令」，与同日 SM89 那组小规模路径调优同一思路。

来源：https://github.com/tile-ai/TileOPs/pull/2682、https://github.com/tile-ai/TileOPs/pull/2679

3.7 TileOPs：GQA 反向在 warp 特化 kernel 里直接服务分组 KV head（#2685）

合于 10-21 22:58。SM90、head dim 128 时，GQABwdOp 此前把分组 KV head 交给流水化 kernel，按 query head 逐个 CTA 走；现在改由 warp 特化的 GQA 反向 kernel 直接吃分组 KV head，减少重复搬运与调度开销。同日 #2687 修 PrivateUse1 重命名回归测试里的设备后端自动加载问题，属于测试基础设施的清理。

来源：https://github.com/tile-ai/TileOPs/pull/2685、https://github.com/tile-ai/TileOPs/pull/2687

3.8 TileOPs 在审：折叠 edge 轴单次发射、NSA top-k 收窄列 tile（#2689 / #2690）

窗口末尾两笔仍在审。#2689 针对 reduction：把边缘轴折叠进一次发射，并让 fold 的尾部装载先发；#2690 在 SM89 上给 NSA top-k 用更窄的列 tile。合入 20 笔之后，TileOPs 的在审队列从上一期的 16 笔降到 9 笔，是本窗口组织内唯一在审数量明显下降的仓库。

来源：https://github.com/tile-ai/TileOPs/pull/2689、https://github.com/tile-ai/TileOPs/pull/2690

3.9 采用方：TileKernels 与 FlashQLA 窗口内无推送

deepseek-ai/TileKernels（1991★）与 QwenLM/FlashQLA（732★）本窗口均无提交，最新推送都停在 10-14；FlashQLA 的版本仍为 09-30 的 v0.1.3。采用方侧的平静已经持续一周，主线的动能仍在 tile-ai 组织内部。

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

四、社区、教程与活动

4.1 社区教程：中文《TVM / TileLang Cookbook》更新

RuneFang/tvm_tilelang_cookbook（34★，7-21 建仓）10-21 有推送。自我定位写得很清楚：面向「能写点 Python、听过 CUDA 但没读过编译器代码」的读者，手把手从 @tilelang.jit 一路挖到 .cubin。这类把 DSL 到编译产物全链路讲透的中文材料仍然稀缺，对昇腾、海光等新后端的入门者尤其有用。

来源：https://github.com/RuneFang/tvm_tilelang_cookbook

4.2 社区新仓：tilelang-wrapper-for-sail 建仓，暂无说明

flytiger-eco/tilelang-wrapper-for-sail 于 10-21 09:04 建仓，无描述、无 README，仓名指向为 SAIL 提供 TileLang 包装。这类「厂商内部运行时 + TileLang 外壳」的小仓通常是新硬件接入的前哨，但内容尚未公开，本期只作记录。

来源：https://github.com/flytiger-eco/tilelang-wrapper-for-sail

4.3 社区工具：远程自动调优 agent 仓继续推送

cc-c122/TileLang-Remote-Autotuning-Agent（1★，7-01 建仓）10-21 22:57 有推送。与 TileFoundry 把调度交给 agent 的方向一致，社区侧也在尝试把自动调优做成可远程驱动的 agent，而非本地脚本遍历。

来源：https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent

4.4 文档与版本：文档站七次自动更新，主仓 VERSION 仍 0.1.15

tilelang.github.io 在窗口内自动更新七次（10-21 10:18、11:26、11:33、13:54、18:26、19:23 与 10-22 01:23），均随主仓合入触发。主仓 VERSION 仍为 0.1.15（该版本 09-30 发布），README 的最新里程碑仍停在 09-30 的 Ascend 950 后端；TileOPs 无新 tag；TileFoundry 今日补上 v0.0.3 tag，是窗口内唯一的新版本。TileOPs-nightly 的 pushed_at 虽有变动，但其默认分支自 08-30 起无新提交，变动来自非默认分支，不计入开发活动。

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

五、趋势观察

编译器缺陷的性质从「算错数」扩到「选错目标」。 本期主仓八笔合入里，四笔改变的是「编译器以为自己在做什么」：无 arch 的 CUDA target 静默退到 sm_50（#2860）、cast 的舍入注解按分支静默丢弃（#3350）、批式归约只处理首个输出块（#3070）、流水规划丢失传递性依赖（#2822）。它们的共同点是不报错，只在某些尺寸或某些循环形态下给出不可信结果——比崩溃更难定位，也说明项目把静默正确性放在了功能扩张之前。

昇腾线的瓶颈已移到内存模型与运行时边界。 亚 32B 紧凑行距无法用 DMA 表达、只能显式中止（#1795）；缓存并发改为原子发布并补删除恢复（#1903）；PTO 的整数符号性要靠包装层补齐（#3582）。这三笔都不是「加算子」，而是「让已有算子在边界条件下不出错」，配合 DeepGEMM 启发式对齐合入，昇腾侧正从能力覆盖转向工程可靠性。

分析机构开始把 TileLang 当作昇腾可用性的判据。 SemiAnalysis 用「六个仓库中只有 TileLang 当时可用」来定义 DeepSeek 昇腾支持的完成度，这比任何单篇技术博客都更能说明 TileLang 在国产算力叙事里的位置：从实验性 DSL 变成默认承载层。同一篇报告也给出反向信号——优化集中在超节点、x86 主机支持仍不完整，通用性落后于性能。

注意力栈的重心从推理前向移到训练反向。 TileOPs 一天内补齐 MLA 预填充反向的 op 层、SM90 warp 特化 kernel 与取数顺序优化，再加上 GQA 反向的分组 KV head 处理，与昇腾侧上月以来的反向算子补位（indexer_bwd、SiTU-GLU 反向、WY fast 反向）完全同向。训练侧的 TileLang 覆盖度正在被系统性补齐。

SM89 从「能跑」走向「调优」。 FP8 BMM 瓦片、W4A16 的 M 上界、分页 DSA、NSA top-k 四条线同日动刀，且都给出单卡实测与「不构成普遍优势」这类限定语。消费级与单卡推理场景成了性能目标之一，同时作者对基准口径的自我约束明显变严。

调度与正确性的分工在 TileFoundry 与主仓之间变得清晰。 主仓继续收紧语义与正确性，TileFoundry 0.0.3 则把「选哪条指令」交给 agent、把「保证语义」留给编译器，并用 H200 与 Apple AMX 两个示例给出可复现数字。这条分工如果稳定下来，TileLang 生态会出现「语言层守正确、agent 层搏性能」的双层结构。

附录：素材与核查说明

信源核查表

信源 核查结果
tile-ai/tilelang（main） 窗口内 8 次提交、8 笔合入；VERSION 仍 0.1.15；在审 163 笔，其中 21 笔窗口内更新
tile-ai/TileOPs（main） 窗口内 20 次提交、20 笔合入；在审 9 笔，3 笔窗口内更新
tile-ai/tilelang-ascend（ascendc_pto） 窗口内 3 次提交、3 笔合入；在审 235 笔，13 笔窗口内更新
tile-ai/tilelang-hygon（main / feat-hcu-smls-b-vgpr） main 无提交；feature 分支 3 次推送（10-21 11:04、15:27）；在审 3 笔
tile-ai/TileFoundry（main） 无新提交；v0.0.3 于 10-21 10:26 打 tag 发布；在审 1 笔
tile-ai/tilelang-metax（dev）/ musa / mlir-ascend 三仓窗口内均无提交（最新分别 10-16、09-30、10-19），无新 tag
tile-ai/TileRT / tilescale / DeepStack 三仓窗口内均无提交（最新分别 09-24、08-25、10-05）
deepseek-ai/TileKernels / QwenLM/FlashQLA 均无提交，最新推送均停在 10-14
tilelang.github.io / TileOPs.github.io 分别 7 次 / 1 次自动更新，均随主仓合入触发
TileOPs-nightly 默认分支自 08-30 无新提交，pushed_at 变动来自非默认分支
Google News（中英多组）/ Hacker News / arXiv 窗口内三篇产业报道围绕华为-DeepSeek 协同，仅 SemiAnalysis 一文含 TileLang 专章；arXiv 最新相关论文仍为 07-24

补充说明：海光线的推送不在默认分支，判定方法是先取 feat/ 分支的提交流，再用仓库事件接口核对 PushEvent 的 ref 与 before→head；TileOPs-nightly 的 pushed_at 变动同样按此法排除。在审数量取各仓 is:open 的检索总数，窗口内更新数取 updated:>=窗口起点 的检索总数。本窗口起点（10-21 07:00）之前半日内落地的提交由上期报告覆盖，本期不重复计入。

完整信源清单