调研窗口:2026-10-16 07:00 ~ 2026-10-19 07:00(北京时间,72 小时,覆盖周末)。

信源:GitHub(tile-ai 组织 29 仓库推送全量核查,窗口内 8 个仓库有推送:TileOPs、TileOPs.github.io、TileOPs-nightly、TileFoundry、tilelang、tilelang-hygon、tilelang-metax、tilelang.github.io;组织窗口内合并 61 笔〔TileOPs 44、TileOPs.github.io 5、tilelang 5、tilelang-hygon 3、TileFoundry 2、tilelang-metax 2〕、在审 69 笔〔tilelang 51、TileOPs 8、tilelang-ascend 5、tvm 2、TileFoundry 1、tilelang-hygon 1、tilelang-mlir-ascend 1〕;非默认分支逐查〔ascendc_pto / dev / snapshots〕)+ 发布页与版本文件核查 + 采用方仓核查(TileKernels / FlashQLA)+ 社区仓核查(tilelang.github.io、TileOPs.github.io、easy-tilelang、AnacondaRecipes、learn-tilelang、BM1690 pipelines)+ 夜间流水线核查 + Google News 中英文多组查询(走代理)+ Hacker News + arXiv。


本期索引

  • 今日重点:TileOPs 单日 20 笔把八算子推入已实现,昇腾三笔与海光块缩放收口——PTO 向量化框架进入在审(10-16)
  • 一、核心项目进展
    • 1.1 主仓:纯 TMA 生产者纳入正向屏障,修复 warp-specialization 死锁(#3503)(10-16)
    • 1.2 主仓:兄弟流水线共享缓冲只展开一次,修 LayoutInference 崩溃(#2342)(10-16)
    • 1.3 主仓在审:Profiler 三连修复与 eager 前端源码定位(#3559 / #3560 / #3561 / #3558)(10-18)
  • 二、多后端适配(昇腾 / 沐曦 / 海光 / 摩尔线程)
    • 2.1 昇腾:SS 缓冲、RNG 状态重置与 SimdVF 归约三笔合入(#3514 / #3331 / #3437)(10-16)
    • 2.2 昇腾:T.Parallel 向量化到 PTO VMI 框架在审,44 文件、近两万行(#3512)(10-18)
    • 2.3 海光:块缩放 GEMM 的 scale copy 三笔——K-major 打包、FP32 MLS、流水化(#22 / #23 / #24)(10-17)
    • 2.4 沐曦与摩尔线程:metax dev 线同步主线并修 xfail;musa 窗口内无推送(10-16)
  • 三、生态与采用方(TileOPs / TileFoundry / 算子库)
    • 3.1 TileOPs:注意力侧两笔转正——MLAVarlen 转 implemented、DSAPaged 落地(#2629 / #2633)(10-18)
    • 3.2 TileOPs:CausalConv1d 双算子、Hadamard 变换与 SM90 W4A8 GEMM 补齐(#2627 / #2624 / #2590)(10-18)
    • 3.3 TileOPs:线性注意力四家族逐点位性能收口(GLA / GDN / KDA / DeltaNet)(10-18)
    • 3.4 TileOPs:Softmax 跨 16-CTA 集群分裂与 GLA decode 状态切片(#2632 / #2631)(10-18)
    • 3.5 TileFoundry:存储计划与发射计划分离,schedule 里程碑修正(#243 / #241)(10-18)
    • 3.6 采用方与夜间流水线:TileKernels / FlashQLA 无推送;夜间快照恢复至 4 次(10-19)
  • 四、社区、教程与活动
    • 4.1 文档:中文 trace 指南对齐,八个新算子写入 API 页(#76)(10-18)
    • 4.2 版本线:TileOPs 仍 pin TileLang 0.1.12,升级 0.1.15 的 PR 关闭未合并(#2554)(10-18)
    • 4.3 媒体与学术:Google News / Hacker News / arXiv 窗口内零新增(10-18)
  • 五、趋势观察
    • 5.1 从「算子补齐」转向「逐点位性能收口」
    • 5.2 注意力链路对齐 DeepSeek-V3.2 生产格式
    • 5.3 昇腾从「能跑」走向「标量并行自动向量化」
    • 5.4 编译器进入正确性硬化期
  • 附:素材与核查说明
  • 完整信源清单

今日重点:TileOPs 单日 20 笔把八算子推入已实现,昇腾三笔与海光块缩放收口——PTO 向量化框架进入在审(10-16)

周末三天的重心从主仓移到算子库:TileOPs 在 10-18 一天合入 20 笔,窗口内累计 44 笔,直接把 8 个此前仅有 spec 的算子推入 implemented 状态,覆盖 MLA varlen 前填、DSA 分页解码、因果卷积、Hadamard 变换、W4A8 GEMM、分页 KV 汇聚与注意力状态合并。主仓 tilelang 同窗口 5 笔合入,两笔是 CUDA warp-specialization 与流水线布局的正确性修复,三笔属昇腾;海光仓库把块缩放 GEMM 的 scale copy 从打包、原生化到流水化补齐。更值得关注的是在审队列:昇腾 PTO 向量化框架 #3512 以 44 文件、近两万行新增把 T.Parallel 编译到 PTO VMI,目标直指 Ascend 950;主仓 51 笔在审中 27 笔标 BugFix,说明项目正处在算子铺开之后的正确性硬化阶段。


一、核心项目进展

1.1 主仓:纯 TMA 生产者纳入正向屏障,修复 warp-specialization 死锁(#3503)

在纯 TMA 生产者配置下,ProducerConsumerWSRewriter 让生产者循环跑遍所有生产者 warp,但此前只有被选中的线程到达正向(full)屏障,1 至 3 号生产者 warp 只等背压(empty)屏障,与消费者不同步,可能落后两个相位并在复用的 parity 上挂死整块。改动让每个生产者线程都到达正向屏障(arrive_count 取 producer_extent),与 SIMT、cp.async 生产者的既有做法一致,修 #3470。

来源:tile-ai/tilelang#3503(10-16)

1.2 主仓:兄弟流水线共享缓冲只展开一次,修 LayoutInference 崩溃(#2342)

两个相邻的 T.Pipelined 循环共用 alloc_shared 缓冲但 num_stages 不同(例如 2 与 4)时,第一条流水线把缓冲从 (64,32) 展开为 (2,64,32),第二条的收集阶段把已展开的缓冲当成原始分配再插一维,得到四维 (4,2,64,32),LayoutInference 的 ICHECK 直接失败。改动按原始分配识别并只展开一次,替换此前的 #2337,修 #2309,新增 349 行、删除 6 行。

来源:tile-ai/tilelang#2342(10-16)

1.3 主仓在审:Profiler 三连修复与 eager 前端源码定位(#3559 / #3560 / #3561 / #3558)

Profiler 三笔针对基准误差与跨设备问题:#3559 修跨设备 CUDA Graph 捕获为空(复用上一张卡缓存的捕获流,replay 不执行,回归计数为 9 而非 27);#3560 在后续调用覆盖前快照输出;#3561 把 cache-flush 时间排除出早停判据,避免先试慢候选改变最终胜者。另有 #3558 把 eager 重写器的模板行号映射回用户语句首行,修 traceback 与覆盖率错位。

来源:tile-ai/tilelang#3559、#3561、#3558(10-18)


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

2.1 昇腾:SS 缓冲、RNG 状态重置与 SimdVF 归约三笔合入(#3514 / #3331 / #3437)

#3514 新增 T.alloc_ss:dav-3510 上一个 AIC 与其两个 AIV 共享的标量存储,每个 AIV 可发布自己的 uint32 字,AIC 经打包的 uint64 视图读取;组内 3 KiB SS 池、32 字节对齐基址、版本计数与同步走标准缓冲管线。#3331 在 AddFunction 中重置 RNG 状态,避免兄弟内核残留的 __asc_rng_state 导致 bisheng 报未声明标识符。#3437 让 SimdVF 内 dim=1 归约只存 lane 0。

来源:tile-ai/tilelang#3514、#3331、#3437(10-16)

2.2 昇腾:T.Parallel 向量化到 PTO VMI 框架在审,44 文件、近两万行(#3512)

该 PR 让用户在 T.SimdVF(lanes=L) 内书写标量并行 T.Parallel(逐元素、控制流、归约),由编译器向量化为 PTO VMI 指令,再经 PTOAS 编译到 Ascend 950 运行。分工上,TileLang 侧负责逻辑 lane 布局、向量化 pass 管线、VMI 指令生成与 Fragment 逻辑映射,工具链侧负责物理寄存器切分、mask 物化与向量寄存器分配。改动 44 个文件、新增 18988 行、删除 67 行,是 PTO 线从执行路由走向自动向量化的关键一步。

来源:tile-ai/tilelang#3512(10-18)

2.3 海光:块缩放 GEMM 的 scale copy 三笔——K-major 打包、FP32 MLS、流水化(#22 / #23 / #24)

#22/#23 支持 K-major LDS 的打包 scale copy,并支持 gfx946 原生 FP32 MLS GEMM(kPack=2),同时补块缩放、TF32、FP32 回归;#24 进一步支持块缩放 GEMM 中的流水化 scale copy,新增 357 行、删除 47 行、5 文件。三笔在 10-16 至 10-17 间全部合入,海光线的块缩放 GEMM 从「能算」推进到与主仓流水线对齐。

来源:tile-ai/tilelang-hygon#22、#24(10-17)

2.4 沐曦与摩尔线程:metax dev 线同步主线并修 xfail;musa 窗口内无推送(10-16)

tilelang-metax 默认分支为 dev,窗口内两笔合入:#160 修 MetaXGPU 的 xfail 用例(含 test_cuda_fma_fmul.py 等),#161 将 main 合入 dev,属跟随主仓基线的常规同步。tilelang-musa 默认分支 main 窗口内零提交,最新一次仍是 09-30 的「Prepare public TileLang MUSA 0.1.15」。tilelang-mlir-ascend 最新提交停在 10-15,无窗口内推送。

来源:tile-ai/tilelang-metax(10-16)


三、生态与采用方(TileOPs / TileFoundry / 算子库)

3.1 TileOPs:注意力侧两笔转正——MLAVarlen 转 implemented、DSAPaged 落地(#2629 / #2633)

#2629 把 MLAVarlenFwdOp 状态改为 implemented,此前打包 varlen 的 MLA 前填内核已合入,本笔补齐 manifest 要求,并让 MLAVarlenCall 按 DeepSeek-V3 8×4K、DeepSeek-V3 混合 512–16K、Kimi K2 4×8K 三种配置构建调用,BF16 与 FP16 双精度注册基准,对标 FA3 的 flash_attn_varlen_func 与 FlashInfer 的 BatchPrefillWithRaggedKVCacheWrapper。#2633 新增 DSAPagedFwdOp,在 FlashMLA 的 DeepSeek-V3.2 656 字节行格式上做分页 FP8 潜缓存的分块解码,返回 o 与 lse,SM90 内核以两个 64 槽键块缓冲在三个 warp 组间轮转,128 头时两段 64 头按 2-CTA 集群协作。

来源:tile-ai/TileOPs#2629、#2633(10-18)

3.2 TileOPs:CausalConv1d 双算子、Hadamard 变换与 SM90 W4A8 GEMM 补齐(#2627 / #2624 / #2590)

#2627 新增 CausalConv1dFwdOp 与 CausalConv1dUpdateFwdOp,宽度 2 至 4、可选 bias 与 SiLU、FP32 累加,前填走打包序列(cu_seqlens 与可选 has_initial_state),解码每序列推进一步,状态形状 [B, D, W-1] 与 vLLM 的 causal_conv1d_update 对齐。#2624 的 HadamardTransformFwdOp 支持 base_order 1/12/20/28/40、n 至 2^28、fp16/bf16 配 FP32 累加,用一次 16 位 MMA 对完成四级蝶形。#2590 的 GemmW4A8FwdOp 在 SM90 上做 FP8 e4m3 激活乘组缩放 INT4 权重、BF16 输出,decode 核服务 M≤32,更大 M 走 warp 特化前填核。

来源:tile-ai/TileOPs#2627、#2624、#2590(10-18)

3.3 TileOPs:线性注意力四家族逐点位性能收口(GLA / GDN / KDA / DeltaNet)(10-18)

窗口内 TileOPs 绝大多数合入标 [Perf],集中在四条线性注意力线。GLA:前填状态链去掉一次矩阵乘(#2626)、decode 状态按 value 切片一键一块(#2631)、SM89 BF16 的 gate 梯度归约展开(#2618)与键归约部分展开(#2617)。GDN:前填只在走查留下空闲 SM 时才分区(#2628),强 gate 衰减下保持前填有限(#2594)。KDA:chunk scan 拆成状态走查与块并行输出(#2623),状态尾条件改为布尔判定(#2630)。DeltaNet:反向状态梯度每块只做一次矩阵乘(#2625)。

来源:tile-ai/TileOPs#2626、#2631、#2628、#2623(10-18)

3.4 TileOPs:Softmax 跨 16-CTA 集群分裂与 GLA decode 状态切片(#2632 / #2631)

#2632 针对 lm-head 这类少行长行(4×102400,FP16/BF16)新增 SoftmaxClusterSplitKernel:旧实现每行 58 段经全局内存交换 (max, sum) 并跨 grid 屏障,新实现把一行摊到 16-CTA 集群的寄存器上,每个元素只取一次指数,跨 warp、跨 CTA、跨集群用 shuffle 树按 max-first 折叠,每 CTA 只向对端共享内存发一对值。#2631 则把 GLA decode 每头一块改为每块一个 value 切片,只读一次状态,键轴归约先用 16 字节向量做局部和再一次批量交换。

来源:tile-ai/TileOPs#2632、#2631(10-18)

3.5 TileFoundry:存储计划与发射计划分离,schedule 里程碑修正(#243 / #241)(10-18)

#243 把原始分配决策留在 analysis 层,移除可执行别名发布与 MemoryMetadata.alias,拆出只读的 StoragePlan / IssuePlan、TirBuilder、命名与单一 LoweringContext;访问结果做记忆化,转发复用其值,并修复在稀疏物理布局下仍按逻辑字节预留的共享内存放置问题,39 文件、新增 2543 行、删除 1889 行。#241 是一笔 314 文件的里程碑修正(新增 53049 行、删除 69314 行),覆盖 M1 至 M8:修正 Reshard 流量与别名、存储实例比较、规范布局与局部投影、窗口绑定与词法作用域校验,并迁移两个 Qwen 示例、重生成 31 个完整长组与全部 64 个短 TIR 文件。

来源:tile-ai/TileFoundry#243、#241(10-18)

3.6 采用方与夜间流水线:TileKernels / FlashQLA 无推送;夜间快照恢复至 4 次(10-19)

deepseek-ai/TileKernels 最新提交停在 10-14 的 v2.1.0-release 合并,QwenLM/FlashQLA 停在 10-14 的 fix-sm120,两者窗口内均无推送。TileOPs-nightly 的 snapshots 分支在窗口内产生 4 次快照:10-16 11:35、10-17 04:13、10-18 06:02、10-19 05:41,均为「bench: (run )」形式,说明夜间基准流水线已恢复按日推进。

来源:TileOPs-nightly snapshots(10-19)


四、社区、教程与活动

4.1 文档:中文 trace 指南对齐,八个新算子写入 API 页(#76)

TileOPs.github.io 窗口内 5 笔合入,其中 #76 让中文 trace 指南与英文页一致,跟随 TileOPs #2619 的 trace.run 变更(slots 常从被跟踪构建中拆出,仅在跟踪开启时解码并写 HTML),并把主仓已实现的八个算子写入对应家族页:GemmW4A8FwdOp、HadamardTransformFwdOp、MergeAttentionStatesFwdOp、PagedKVCacheGatherFwdOp、MLAVarlenFwdOp、DSAPagedFwdOp、CausalConv1dFwdOp 与 CausalConv1dUpdateFwdOp,Transform 页同时去掉「尚未实现」标注。tilelang.github.io 另有一笔由 GitHub Actions 自动生成的文档更新。

来源:tile-ai/TileOPs.github.io#76(10-18)

4.2 版本线:TileOPs 仍 pin TileLang 0.1.12,升级 0.1.15 的 PR 关闭未合并(#2554)

主仓最新发布仍是 09-30 的 v0.1.15,而 TileOPs 的 pyproject.toml 至今写死 tilelang==0.1.12。尝试升到 0.1.15 的 #2554 已关闭且未合并,PR 自述的阻塞是 29 个在 0.1.12 上通过的 smoke test 在 0.1.15 上失败:其中 28 个源于动态 extent 的 T.Pipelined 在 cp.async 尾声留下 ptx_wait_group(3 - t * 2)、CUDA codegen 报 Downcast 失败(与主仓 #2859 关闭 explicit_unroll 默认值有关),另 1 个是 LegalizeSafeMemoryAccess 在折叠为 False 的分支下判空失败。该 PR 还计划在 TileOPs 侧加 FIXME(staged-rollout) 绕过,再重建 runner 镜像。

来源:tile-ai/TileOPs#2554(10-18)

4.3 媒体与学术:Google News / Hacker News / arXiv 窗口内零新增(10-18)

Google News 中英文共 13 组查询(TileLang、tilelang 算子、tile-ai、TileOPs、TileFoundry、TileKernels、tilelang 昇腾、tilelang 国产芯片等)在窗口内无命中,最近一条仍停在 10-02 的 DeepSeek 开源昇腾组件系列报道;Hacker News 的 tilelang 查询零命中;arXiv 的 TileLang 与 tile-level 检索无窗口内新论文,最近相关条目为 07-24 的 TileSight。

来源:Google News TileLang 检索(10-18)


五、趋势观察

5.1 从「算子补齐」转向「逐点位性能收口」

窗口内 TileOPs 44 笔合入里,标 [Perf] 的占绝大多数,且几乎每笔都点名对标基线:torch-compile、quack、FA3、FlashInfer,并给出具体形状(serving-b1-h8、serving-b1-h16、lm-head 4×102400、8192 至 16384 的 RMSNorm 行)。这说明算子面铺开之后,工作重心已转为在真实服务形状上逐个消掉零点几个百分点的差距,而非继续堆算子数量。

5.2 注意力链路对齐 DeepSeek-V3.2 生产格式

DSAPagedFwdOp 直接实现 656 字节行格式的 FP8 潜缓存分页解码,MLAVarlenFwdOp 按 DeepSeek-V3 与 Kimi K2 的真实配置注册基准,PagedKVCacheGatherFwdOp 与 MergeAttentionStatesFwdOp 补齐 KV 汇聚与状态合并。这条链路加上线性注意力四家族,已经覆盖从长上下文前填到解码的完整推理路径,且以 vLLM 与 FlashMLA 的存储布局为准绳。

5.3 昇腾从「能跑」走向「标量并行自动向量化」

上期的主线是 PTO 执行路由合入,本期 #3512 把 T.Parallel 直接编译到 PTO VMI,并明确以 Ascend 950 为目标、与 PTOAS 分工物理寄存器与 mask 物化。同期三笔小修复(SS 缓冲、RNG 状态、SimdVF 归约)都是在补 dav-3510 上的存储与同步语义。这条线正在从「把 kernel 跑起来」转向「让用户按标量写法写、编译器负责向量化」。

5.4 编译器进入正确性硬化期

主仓 69 笔在审 PR 中 tilelang 占 51 笔,其中 27 笔标 [BugFix],集中在 ThreadSync(地址写入判定、谓词顺序)、Transform(向量化域规范化、跨线程 WAW、流水线传递依赖)与 codegen(cp_async 等待深度、ROCm FP4 常量、Metal hoisting)。这批问题多在特定组合下才暴露,属于典型的规模化使用后回灌的边界条件;TileFoundry 的 #243 拆出只读计划对象,也是同一方向的结构性治理。


附:素材与核查说明

  • 本窗口为周一三天窗口,覆盖 10-16 07:00 至 10-19 07:00,与上一期(10-15 07:00 至 10-16 07:00)不重叠。
  • 组织级仓库列表经 GitHub API 全量拉取,29 个仓库中 8 个在窗口内有推送;众仓库主分支提交经 Atom feed 逐条核对,非默认分支(tilelang-metax 的 dev、TileOPs-nightly 的 snapshots、tilelang-ascend 的 ascendc_pto)单独核查。
  • 合入与在审数量取自 GitHub 搜索接口的真值检索(is:merged 与 is:open 各一次),未使用逐仓 pulls 列表累加,避免重复计数。
  • tilelang-ascend 默认分支 ascendc_pto 最新提交停在 10-14,窗口内无主分支合入;其 5 笔在审 PR 于 10-16 有更新,属延续在审,未计入本期新增。
  • tile-ai/tvm 为 apache/tvm 的 fork(分支 tilelang_main),窗口内有 2 笔在审 PR 更新,最新推送为 10-07,与 TileLang 主线无直接交付关系,仅记录不展开。
  • 新闻侧零新增已用裸查询验证代理通路正常(同一查询可返回 37 条历史结果),确认是窗口内无新报道而非采集失败。

完整信源清单