TileLang周报(2026-10-12 ~ 2026-10-16)
报告周期:2026 年 10 月 12 日(周一)至 10 月 16 日(周五),共 5 个工作日、5 期日报素材 素材来源:本刊每日 TileLang 动态日报 5 期(10-12、10-13、10-14、10-15、10-16);另经 tile-ai 组织 29 个仓库的 168 小时窗口补充核查、夜间流水线六份快照的测试结果逐份复算与各仓库发布元数据复核(详见附录)
一、本周要闻
-
主仓 ROCm 线连续四日收笔,DeepSeek 新架构内核成为第三收笔线(10-13 ~ 10-14):10-13 合入 DeepSeek NSA 前向/单 token 解码的 64 lane 单波实现(#3152,+111/-2)与 flash-decoding 的 LDS/VGPR 预算自适应瓦片(#3149,+82/-5);10-14 凌晨再合入 mHC(流形约束超连接)前/后处理内核(#3156,2 文件 +83/-6,HIP 侧改在完整 64 lane 波前上分工并把 FP32 混合权重转 BF16 供 MFMA)与 safe-gated KDA 解码示例(#3249,+360,含 GLM-5.3-Flash TP8 专用化配置);10-14 另有 cumsum/cummax 的 64 lane 段宽修复(#2326)。四笔共同指向「按设备现实自适应」:gfx942 波前尺寸、MFMA 输入类型、LDS/VGPR 预算逐项清偿,ROCm 支持从「跑通」转入「与新架构发布同步」。
-
昇腾 PTO 由「独立目标」升为「一等执行路径」,NSA 双迁移与回滚纪律同周发生(10-12 ~ 10-15):10-15 合入执行路由 #3316——
target="ascend"与execution_backend="pto"组合经执行感知的目标规范化后直连CodeGenTileLangPTO,auto路径维持 AscendC + TVM-FFI;前置的示例目标选择 #3474(10-14 合入,移除硬编码target="ascend",CI 对 ascend / pto 双目标各跑一遍)与本窗口在审的tvm_ffi后端、VMI 向量化、SS buffer、部分 L0 拷贝 pitch 等七笔构成家族。同周 10-12,昇腾 NSA Decode(#1701,+755,28 例分层测试)与 NSA Bwd(#1730,+1593,20 例)合入;10-14 三笔「文档事实核查」PR(#1612、#1676、#1681)因低优先级预期失败用例的异常正则不匹配被整批回滚(#1888 至 #1890),bitwise_rshift 以 重做版 #1893 删减预期失败断言后重新提交;10-15 开出 dav-3510 上 gemm_v0 的静默数值错误修复(#1900),暴露--npu-arch无条件取dav-2201导致 A5 内核瞄错架构。 -
TileOPs 一周 52 笔合入:Op 基类重构定稿、MoE 双算子落定、SM89 成为第二平台(10-12 ~ 10-16):10-11 合入的 Op 基类重构(#2520,363 文件、+3362/-5227,自称「发布前」收口)把带与不带编译边界的算子统一经
_run_call执行、select_kernel_key/select_implementation合并为key_for、编译边界改从 manifest 推导,文档站次日跟进《Op 基类指南》;10-13 MoE 双算子同日落定——masked FP8 分组 GEMM(#2521,20 文件 +1788/-37,对齐deep_gemmmasked 语义)与 DeepSeek 分组 top-k(#2535,15 文件 +1193/-69,对齐 vLLMgrouped_topk并修 T=512 时 32/132 SM 的调度问题);SM89 从 10-12 起成批收口(GQA 预填/解码、GDN / DeltaNet 预填、GLA / DeltaNet 反向、KDA 预填、Norm 全谱),RTX 40 系与 L40 / 6000 Ada 级工作站在算子库侧与 Hopper 并列。 -
夜间底盘六份快照:正确性 1032 → 1134(+102)、基准 1947 → 1989(+42),全程零失败(10-12 ~ 10-16):本周六份快照的正确性用例由 1032 项升至 1134 项、基准由 1947 项(52 套件)升至 1989 项(55 套件),六份均为零失败零错误、各 1 项跳过。单日最大增幅出现在 10-15(+41,与 TileOPs 16 笔合入同拍);10-16 出现两次快照(凌晨与夜间各一),为本周首次单日双快照。评测环境为 H200、CUDA 13.2、torch 2.13.0+cu132、tilelang 0.1.12(详见第四章表格)。
-
多后端同版本收口:海光 v0.1.15 升级落地并开出分布式 IPC 线,沐曦追平上游,摩尔线程全周静默(10-13 ~ 10-16):海光 10-14 合入 v0.1.15 全量升级(#20,保留 HCU 后端并新增块缩放 GEMM 的 scale 缓冲布局规划)并建立
v0.1.15-release分支,同日以 #19(+2577/-28,44 文件)开出 Distributed 线(统一进程启动器、IPC allreduce 测试、IPC GEMM allreduce 示例等);10-15 以 #21 把块缩放异步拷贝布局当日开、当日合进v0.1.12-release维护线;10-16 再合入 #22(K-major 打包 scale 拷贝与原生 FP32 MLS GEMM)。沐曦 10-13 将 main 推至上游 v0.1.15 节点(镜像 #3309),10-16 dev 线合入上游 main 并修一处 xfail 用例。摩尔线程自 09-30 后全周无活动,是四家中唯一未对齐 v0.1.15 的后端。 -
TileFoundry 一周十笔合入:稀疏解码贯通 TIR、Qwen3.8-27B 预填首次全长度超越 SGLang、核心重构与外部审计同期(10-12 ~ 10-16):10-12 合入稀疏解码贯通 TIR 的 #223(84 文件、+4424/-1543,修丢失循环工作量、查找计费过高与分阶段更新错误);10-13 合入 Qwen3.8-27B 预填示例 #229(89 文件、+38069),为该仓首个在 1024 / 4096 / 8192 三个长度上全部快于 SGLang 的示例(8192 时打分只占 FLOP 的 3.1%,负载超出 H200 拐点 36 倍);10-15 同窗落地实例槽重构(#232,源码全套 1674 通过)与 lowering / printer 契约(#233);10-16 补上 analysis 访问上下文(#238)与 issue 237 的分配/lowering 修正(#241)。重构落地的同一窗口,第三方技能仓即把 TileFoundry vendor 进工作流并列出 4 个 emitter bug 与 2 个运行时 bug——进入外部审计期。
-
采用方首次出现版本迭代:TileKernels v2.1.0 昇腾侧重写、FlashQLA 修复 Blackwell sm120(10-14):DeepSeek TileKernels v2.1.0(16 文件、+352/-382)是该仓 09-30 加入昇腾支持后的第一次版本迭代,改动集中在昇腾侧(
norm_forward_asc大幅重写、SwiGLU 前后向、Engram 门控),README 要求升至 TileLang 0.1.15、PyTorch 2.13、Ascend 950 NPU 与 CANN 9.2.0;Qwen FlashQLA #48(5 文件、+82/-5)修复自 0.1.13 起 Blackwell sm120 前向/反向不可用(block_DV被误改为 128、缺aggregate_card_state),RTX 5090 + tilelang 0.1.15 + CUDA 13.0 全部测试通过。 -
硬件矩阵同周扩张:SM121 点亮、SM120 块缩放双线在审、D=256 解码扩至 SM80/SM90、CPU OpenMP 网格并行落地(10-12 ~ 10-15):SM121(GB10 / DGX Spark)块缩放 MMA 于 10-11 点亮(#3430,DGX Spark 驱动 580.178.04 / CUDA 13.0 实测),此前
sm_121a下所有块缩放 MMA 会编译成BPT.TRAP;SM120 侧的 MXFP8/MXFP4(#3495)与 mxf8f6f4 全操作数家族(#3099)全周在审。消费级与工作站侧,D=256 FP16/BF16 稠密解码由 SM89 扩至 SM80/SM90(#2563,附 Gemma-2B 8K cache 负载,同时修复稠密 RoPE split 忽略 softcap);主仓 10-14 合入 CPU 后端的网格级 OpenMP 并行(#3109)与 Windows NVCC 改走cl.exe(#3510)。
二、核心项目进展(主仓 tilelang)
版本与发布
- 主仓最新标签与发布仍为 v0.1.15(09-30 发布),全周无新标签、无新发布页条目,至本周五已 16 天。发布页最近三条依次为 v0.1.15(09-30)、v0.1.14(09-02)、v0.1.13(08-03),此前的月度节奏在本周被打破。
- 适配线标签全周未动:昇腾
v0.1.2.000、MLIR 昇腾v0.1.2.020.sglang.poc、沐曦v0.1.14(10-13 仅 main 镜像到 v0.1.15 节点、未打新标签)、摩尔线程v0.1.14+musa.1;海光以 v0.1.15 主线加v0.1.12-release维护线双线并行。 - 发布静默的对照面是下游约束上移:TileKernels v2.1.0 明确要求 TileLang 0.1.15 起,FlashQLA 的 sm120 修复亦在 0.1.15 上验证——本周起「0.1.15 基线」成为采用方事实上的最低要求。
编译器、执行路径与语言层
- PTO 执行路由(#3316,10-15 合入):新增执行感知的目标规范化钩子(挂在 BackendContext 创建处),昇腾专属规范化器在解析前把原生 Ascend 目标专门化为 PTO 目标变体并保留
arch等属性,设备代码生成改走PTO_CODEGEN与target.build.tilelang_pto_without_compile;文档给出TILELANG_DEFAULT_TARGET=ascend加TILELANG_EXECUTION_BACKEND=pto的等价配置。测试覆盖上下文解析、目标属性保留、缓存分发与环境变量选择。 - PTO DeepGEMM codegen(#3330,10-14 合入):跨控制流跟踪
tl.ascend_set_mmad_direction并把生效方向编码进 MAD 指令、拒绝冲突方向下可达的 MAD;支持 SIMD 后增量加载、混合 AIC/AIV 内核的 UB 地址到 PTODSL 指针转换、动态 L0C→UB 维度与步长、int16 有符号打包 MX scale。自报 15 个 DeepGEMM PTO 用例 14 个通过,test_bf16_m_grouped_padding仍受 PTOAS VPTO 后端无法合法化 FP32→BF16 转换阻塞。 - 独立 AIV 执行(#3485,10-14 合入):扩展 Ascend AutoSchedule,使 AIV0 与 AIV1 可在
sid依赖下执行不同工作(交替收尾、反量化与收尾分工、不相交输入分区、循环内角色切换);AssignCore提前到延迟估计与调度之前,并为 AIV 局部 UB 缓冲与标量状态建立独立身份。 - C 系 codegen 共享基类(#3289,10-14 合入):
CodeGenTileLangCBase统一 CUDA、HIP、CPU C、Ascend C++、Metal 与 C 主机代码生成;修复T.import_source、T.Kernel(prelude=...)与pragma_import_c在源文本缺少末尾换行或以反斜杠结尾时吞掉下一生成行的问题。验证:聚焦回归 83 通过、codegen 与 launch 套件 405 通过。 - CPU 网格并行(#3109,10-14 合入):新增
MaterializeCPUParallelGrid把带注解的 CPU grid 循环转为 OpenMP 并行循环,配套循环 collapse、线程数配置、最小 trip 数门限与分配下沉,补齐 Linux / macOS / Windows 的编译标志——CPU 后端首次具备网格级并行。 - Host IR 线全周在审:统一 CUDA 与 CuTeDSL 执行路径的 #3473 与直连 Host IR JIT 加原生缓存的 #3481 自 10-10 开出后未见合入,仍是主仓最大的在审方向。
- PTO 家族七笔(10-15 起在审):
tvm_ffi执行后端(#3516)、T.Parallel到 PTO VMI 向量化(#3512)、SS buffer(#3514)、部分 L0 拷贝显式 pitch(#3511)、NZ 缓冲 UB→L1 部分拷贝(#3530)、PTO Ramp 值按类型向量发射(#3535)、SIMT VF 内向量局部访问修复(#3523)。其中 SS buffer 于 10-16 上午合入(#3514,日报截点之后经补充核查确认)。
正确性与缺陷修复
- 10-12(周末窗口收口):GEMM dtype 白名单(#3245,不支持组合在 CUDA codegen 前拒绝,缺陷单 #3004)、For 循环步长(#3427)、JIT 输出索引归一化污染(#3318)、popcount 窄整数下沉(#3300,
int8(-1)计 8 位)、位归约放行浮点源(#3332)、CuTeDSL 栅栏选举(#3400,2D/3D 线程块下多线程重复发布到达)、NVRTC / CuTeDSL 参数身份(#3414)。 - 10-13:
thread_storage_sync精确线程边界(#3291,4 文件 +292/-96,修 #2954);昇腾 PTODSL 归约助手 vendor(#3469,+555/-115)修复 CANN Weekly 下ModuleNotFoundError。 - 10-14:SimdVF 多缓冲
T.fill(#3455,多缓冲版本维导致整缓冲填充检查拒绝)、bisheng 临时文件清理开关(#3461)、跨 warp allreduce 分组错发(#3494)。 - 10-15:JIT 拒绝不可哈希编译期参数(#3417,防止可变默认值静默复用陈旧内核,91 通过)、测试编译时间消除(#3517)、治理配置(#3513,CodeRabbit 自动标签)。
- 10-16(含日报截点后经补充核查确认):Ascend SS buffer(#3514)、SimdVF dim=1 归约逐行单元素存储(#3437)、昇腾 RNG 跨内核状态复位(#3331)、兄弟流水线共享缓冲防双扩(#2342)、TMA-only 生产者限定 warp 0(#3503)。
- 10-16 凌晨新开待审:cumsum 共享缓冲在 warp 特化下的修复(#3536)——消费者线程从 128 起编号而扫描模板未收到区间下界,导致
CumSum1D/CumSum2D在消费者线程上提前返回、缓冲保持未写;修后 #3528 复现用例 131072 项全对(main 上 129019 项错)。另有 profiler L2 只读刷写(#3534)与共享扫描线程数诊断(#3532)。
性能与硬件矩阵
- SM121(GB10 / DGX Spark):#3430 让
tl::sm120_mma_sync_blockscaled按 arch 与家族同时接受 SM120 / SM121,NVF4 测试按 SM12x 家族门控,CuTeDSL TMA 架构表放行sm_121/sm_121a。 - SM120 块缩放:MXFP8(
kind::mxf8f6f4、1X、UE8M0)与 MXFP4(mxf4nvf4、2X、UE8M0)扩展(#3495,+884/-51,FP32 累加、每 32 个 K 元素一标度)与全 f8f6f4 操作数家族(#3099)持续在审。 - ROCm:除要闻四笔外,10-15 的 scan 波前修复(#2326)让
cumsum/cummax按目标波前宽度选择段宽,32 线程块在编译期排除 64 lane 模板;B300 上 CUDA scan 套件 17 通过。 - SM89 成批性能政策:MLA 解码按批与头块数选 2/4/8/16 KV 分裂(#2511:单请求 bf16 4k 0.158→0.034 ms 约 4.6 倍、32k 1.170→0.162 ms 约 7.2 倍)、KDA 预填按 16 列值瓦片切分(#2510:4k/32k/64k 降 22.9%/19.8%/18.5%)、扫描加宽(#2574:4096~16384 token 由 8 值头扩到 16 值头,19 行提速 2.5%~19.7%)。
- Windows / Metal:Windows NVCC 改走 MSVC(#3510,CUDA 13.3 下复现用例修复);Metal 批十笔(#3194 至 #3210)在审,覆盖段归一累加器、simdgroup GEMM 片段、BF16 向量类型、短常量循环展开等。
评审队列观察
- 主仓开放 PR:170(10-12)→ 168(10-13)→ 161(10-14)→ 160(10-15)→ 165(10-16),周内在 160 至 170 区间震荡;组织口径的开放 PR 合计由 498(10-12)升至 507(10-13),至本周编制时为 532 笔。
- 组织口径的合入与产出:本周 PR 合入 109 笔、新开 158 笔(主仓 30、TileOPs 52、昇腾适配仓 8、TileFoundry 10、海光 3、MLIR 昇腾 2、沐曦 2、文档站 2;10-12 至 10-16 北京时间窗口,GitHub 检索实测)。10-16 单日在审更新达 172 笔,为本周最高。
- 结构性判断延续并加强:主仓瘦身(170 → 165)、TileOPs 变厚(9 → 16 后回落)、昇腾适配仓回升(230 → 237)——「合入闸门收紧、在审队列变厚」的格局贯穿全周,发布静默期的准备动作集中在评审侧。
三、多后端适配(昇腾 / 沐曦 / 海光 / 摩尔线程等)
昇腾(Ascend):迁移合入、执行路径收敛与回滚纪律
- 迁移合入:NSA Decode(#1701,10-12,+755)以 Developer mode hybrid 实现单核解码(seq_q=1),显式 L1/L0C/UB 分配、5 段 GM workspace 经
workspace_idx=[6..10]自动分配,编译器自动切 Cube/Vector 作用域并插跨核 flag 同步;分层测试 28 例(L0 阈值、L1 功能含 DeepSeek-V3 典型配置 D=128/S=16/BS=64、L2 异常与边界特值),910B3 fp16 实测。NSA Bwd(#1730,+1593)提供nsa_fwd与nsa_bwd_single两枚 JIT 核,反向按 flashattn_bwd 范式以片上直连T.copy(L0C→UB、UB→L1)省去 9 枚 GM workspace,分层测试 20 例。 - 回滚与重做:10-14 三笔文档事实核查 PR(round 语义 #1612、bitwise_rshift #1676、bitwise_xor #1681)被整批回滚(#1888 至 #1890)。原因由重做版 #1893 说明:每日全量套件执行低优先级用例时,预期失败用例的异常正则不匹配——实际抛
AscendVectorInstructionSelection的InternalError而非断言的Compilation Failed,而这些用例标记low_priority、PR 阶段 CI 不执行,由每日全量运行兜底发现。重做版删除全部预期失败断言(支持表外 dtype 改由文档 dtype 支持表覆盖),保留 docstring 重写与真机测试。 - 在审与队列:dav-3510 gemm_v0 静默数值错误(#1900,10-15 开出)为本周最值得跟踪的一笔——无异常、无 NaN/Inf、无零值,多重集检查全过,只在远端表现为精度失败;两笔可分提交(
--npu-arch取自平台、gemm_v0 staging 限制)以待验证依赖顺序。另有 KDA chunk_delta 反向/前向迁移(#1895 / #1896)、bitwise_xor 文档重做(#1899)、situ_glu(#1891)、cann-bench/lstm 示例(#1892)、indexer_bwd(#1887)。开放 PR 由 230 升至 237。 - 主仓侧:除要闻的 PTO 路由与独立 AIV 外,示例目标选择 #3474 让
TILELANG_DEFAULT_TARGET决定目标、DeepGEMM JIT 入口支持execution_backend="auto";跨 warp allreduce 分组修复 #3494 是对 #3469 合入后社区正确性回馈的响应,回路在两日内闭合。
MLIR 昇腾(tilelang-mlir-ascend):显式内存控制落地,知识库修正
- NPUIR 显式内存与流水线控制 #190(10-13 合入,12 文件、+1092/-102):支持带符号 pitch 的双行 GM→UB 拷贝与连续 GM/L1 ND 拷贝、物理缓冲视图与可变标量存储、GEMM L0 ping-pong 事件与 MMAD / Fixpipe 单元标志(含显式同步要求)、运行时尺寸的
T.vbrc区域;修复 Expert 有符号 floor 除/模、无符号算术与惰性条件降低。 - 10-12 的全量性能报告(#197,7 算子 112 例全部通过,CANN 8.5.0 / 910B2C / msprof)与 10-15 的知识库修正(#202,ArgMax N-split 的 partial / merge 描述统一为单次 device launch 内完成)分别合入;
T.Pipelined()参数化(#199)、FloorDiv / FloorMod codegen(#200)、Cube 内存分配 API(#201)在审。开放 PR 稳定在 32 笔。
海光(Hygon):三线结构成形——老线交付、新线追上游、新分支做分布式
- 寄存器流水线(#18):全周在审,49 文件、+4144/-318,为后端加入寄存器流水线(与共享内存流水线协同、按流水级插入提交与屏障)及 warp 数超 4 时的 warp-diverge,作者自述 example_gemm 已达 rocblas 水准。
- FP8 特性分支:
feat/hcu-mqa-fp8上 MLS channelwise FP8 GEMM 自报 432 TFLOPS,10-12 补「M≤32 时 fp8 解码 GEMM 带宽超越 hipblas」并收敛命名(均为分支内作者自测口径,未关联 PR)。 - 主线与发布线:10-14 合入 #20(整体升级到上游 v0.1.15,保留 HCU 后端与
build.sh工作流,新增 HCU 块缩放 GEMM 的 scale 缓冲布局规划与配套 LDS 布局),并建v0.1.15-release分支;10-15 #21 把块缩放异步拷贝布局当日开、当日合进v0.1.12-release维护线;10-16 合入 #22(K-major 打包 scale 拷贝与原生 FP32 MLS GEMM,日报截点之后)。 - 分布式线:#19(+2577/-28,44 文件)与
distributed分支全周两次推进,覆盖统一分布式进程启动器、IPC 元数据初始化校验、BF16 / 全对等 IPC allreduce 测试、IPC GEMM allreduce 示例、IPC 关闭与初始化失败的传播、无效 peer 守卫、STRIDE 远程块拷贝,并有一次「move distributed language APIs」重构。多卡集合通信成为海光下一阶段主线。
沐曦(MetaX):追平上游后转入 dev 线合流
- 10-13 main 分支推至上游 v0.1.15 版本节点(镜像 #3309),镜像线追平发布;dev(适配主工作线)此前停在 09-24。10-16 dev 合入上游 main(
Merge branch 'main' ... into dev)并修一处[fix bug][MetaXGPU]的 xfail 用例——该笔(10-16 12:35)在 10-16 日报截点之后,经本周补充核查补录。开放 PR 稳定在 16 笔。
摩尔线程(MUSA)与 Sunrise:全周静默与依赖升级
- 摩尔线程自 09-30 后全周无推送、开放 PR 6 笔,是四家国产后端中唯一未对齐 v0.1.15 的一家;
v0.1.15+musa.1仍是最新发布(09-30)。 - Sunrise(S2/S3 的 TANG 后端仓)10-12 仅两笔 Dependabot CI 依赖升级(#7、#8)在审,开放 PR 3 笔,无代码面动态。
四、生态与采用方
TileOPs:52 笔合入的一周——重构定稿、MoE 收口、全算子逐日刷榜
- 治理:Op 基类重构 #2520(10-11 合入,363 文件、+3362/-5227)以「先行为、后机械改名」的十个提交完成:带与不带编译边界的算子统一走
_run_call、失败时按是否由本调用选定目标决定回滚、子算子不归任何delegate_for阶段则复合调用整体失败、_validate_dtypes移除、key_for合并两个选择入口、编译边界改从 manifest 推导。文档站次日合入《Op 基类指南》(#69),代码与文档同拍。 - MoE 双落定:masked FP8 分组 GEMM #2521(20 文件、+1788/-37,关闭 #2388)以 FP8 e4m3、1x128 / 128x128 块标度、bf16 输出对齐
deep_gemmmasked 语义,内核为 TMA producer warp 加两个 WGMMA consumer warpgroup 的持久网格,空瓦片不占 CTA、每份权重瓦片只从 HBM 读一次;清单收窄到 masked 布局并配 14 行负载(E=8 两行加 EP=8 解码十二行,覆盖 DeepSeek-V3、Qwen3-235B-A22B、Kimi-K2,batch 64 / 256),新增moe.masked_routed_metadata生成器模拟专家不均衡(DeepSeek-V3 b64 空 4/32、Kimi-K2 b64 空 13/48)。DeepSeek 分组 top-k #2535(15 文件、+1193/-69,关闭 #2268)补齐num_expert_group/topk_group/routed_scaling_factor并对齐 vLLMgrouped_topk,同时修复 T=512 时默认TOKENS_PER_BLOCK只占 32/132 SM 的调度问题。收尾重构 #2551 删除专用 TightRow 内核,布局覆盖扩至 SM80 / SM86。 - SM89 成为第二平台:10-12 三笔(D256 dense GQA 解码 #2529、dense GQA 预填 #2528、GDN / DeltaNet 预填 #2527,21 例「失败不可分发」转通过),10-13 四笔(GLA 反向 #2546:共享内存 114688→81920 字节;DeltaNet 反向 #2547:147968→98816;KDA 融合预填 #2548:256 线程分配 182528→101120;紧凑逐行分组 GEMM #2549),加 GQA split log-sum-exp 改 FP32 累加(#2550,bf16 大分数下 split 混合最多偏差 0.23,autotune 改按实际启动程序计时)与 D=256 解码扩至 SM80/SM90(#2563)。
- 注意力家族收口:10-15 三笔同日合入——varlen GQA 汇聚与流水线滑窗(#2537,FP32 逐头 sinks 覆盖 varlen 部分)、分页 GQA 的只读 RoPE / 双边滑窗 / 逐头 sinks(#2516)、分页 KV 写入算子 PagedKVCacheWriteFwdOp(#2539,按调用方槽位原位散射、跳过 -1 槽、FP8 独立 K/V 标度、保留 NaN 与溢出饱和语义,与 FlashInfer 及 vLLM 对照验收)。此前 10-12 已合入 MLA 解码 SM89(4.6 倍 / 7.2 倍)、MLA 调优分裂扩展(候选 24→40,冷调优 348.6 s→151.6 s)、WGMMA 描述符 lane 0 广播(26 行基准快 0.7%~4.1%)、KDA 预填、NSA top-k 位排序、分页 GQA TMA 特化核(H200 上被接管四行降至原 0.44~0.88 倍)。
- foundry 性能批进入日更节奏:10-13 八笔合入加八笔新开、10-14 十一笔收口、10-15 三笔,代表性数据——softmax 寄存器驻留
_held_row加 warp shuffle 折叠(#2570:39/42 行领先 1.1%~78.1%,仅 lm-head 两行落后 0.909x/0.934x);TMA 无法寻址 N 非 8 倍数输出时逐行逐 warp 存储(#2566:294.7→23.2 微秒,12.7 倍,快于 cublaslt-best 7.5 倍);GELU erf 链缩短五条指令(#2557:1.125x)、SiluAndMul 16 位门快速除法(#2565:1.074x,torch-compile 比从 0.945x 翻正)、分组 GEMM 分块收尾交替 staging(#2582:四行 MoE 负载全部领先 quack,fwd 1.015x / dgrad 1.012x / wgrad 1.033x)、W4A16 权重先行加载(#2586:decode-long-k 达 marlin-fp16 的 1.047 倍,8 行瓦片各提速 2.5%~7.6%)、sigmoid 免 IEEE 除法(#2588:1.020x~1.042x,测试节点 +2)。池化、归一化、归约、RoPE 一并铺开(MaxPool1d/3d、AvgPool3d、RMSNorm 按宽度分行、InstanceNorm 至多 2048 元素整向量行、any/all 逐字测试、组瓦片计数 warp 扫描:4096 组在 32 线程下 62 寄存器而非 152)。 - 线性注意力修复:KDA 预填求逆稳定化(#2584,相关键下 main 在 6 个 chunk 与 2 个 fused 用例失败含 NaN,改为 16 行对角块 FP32 求解并复用共享缓冲)、GLA 反向显存复用(#2569,131072→98304 字节落入 SM89 上限,PR 明示不主张加速)、GLA 预填状态跳过(#2540,省去整张 FP32 初态张量的分配与读取)、单瓦片解码态加载规避 TMA 挂起(#2538)、MoE 专家搜索跳过(#2575)、permute-align 收界(#2559)。
- 队列:开放 PR 9(10-12)→ 19(10-13)→ 26(10-14)→ 17(10-15)→ 16(10-16),一周内先增后减,合入速度与在审厚度交替主导。
TileOPs 夜间流水线:六份快照,正确性 1032 → 1134
- 10-12 03:40 — 快照 fc2a2a91(基准 70cfdf84):正确性 1032 项、基准 1947 项(52 套件),零失败零错误、1 项跳过
- 10-13 03:53 — 快照 4cd9f968(基准 6ec0d086):正确性 1042 项、基准 1947 项(52 套件),零失败零错误、1 项跳过
- 10-14 04:18 — 快照 76e8cac2(基准 a79a21e4):正确性 1064 项、基准 1966 项(53 套件),零失败零错误、1 项跳过
- 10-15 05:15 — 快照 56c2ab6b(基准 35a25677):正确性 1105 项、基准 1970 项(53 套件),零失败零错误、1 项跳过
- 10-16 11:35 — 快照 d6ab1d14(基准 cb15bce7):正确性 1129 项、基准 1982 项(54 套件),零失败零错误、1 项跳过
-
10-17 04:13 — 快照 d44ab6d3(基准 67250e81):正确性 1134 项、基准 1989 项(55 套件),零失败零错误、1 项跳过
- 周净增:正确性 +102 项、基准 +42 项、套件 +3;六份快照各有一项跳过(H200 上不适用的 SM89 调优用例、不支持的 KDA 配置等),无失败无错误。评测环境 H200、CUDA 13.2、torch 2.13.0+cu132、tilelang 0.1.12,最新快照内 tileops 版本
0.0.1.dev1338+g67250e817。 - 与合入的同拍关系:10-13 的 +10 对应 SM89 系列、10-14 的 +22 对应 MoE 与 SM89 反向/预填、10-15 的 +41 对应 TileOPs 16 笔合入(性能批普遍带 test-node 增项)——「合入一笔、底盘爬一截」的滚动记账在整周延续。
- 需要说明的窗口差:10-16 日报生成时未见 10-15 夜间快照,实际该份于 10-16 11:35 落地(北京时间,晚于日报 07:00 截点),本周已补录;表末的 10-17 04:13 一份为 10-16 夜间运行的产物,属 10-16 当日的第二次快照。
TileFoundry:十笔合入——从稀疏解码到可交付工程形态
- 稀疏解码贯通 TIR(#223,10-12 合入,84 文件、+4424/-1543):承接 #218 的 DSA 解码遗留问题——修复丢失的循环工作量、查找计费过高与分阶段更新错误,落地索引式异步收集、Bitcast 与 Where 下沉;分配约束改为单一访问器(返回 Storage),捕获更新直接写入 tile(16B 共享内存),参数初始化的携带值获得独立存储与前置拷贝;16 组 FixtureCase 守卫各含 CLI 调用与输出检查。
- Qwen3.8-27B 预填(#229,10-13 合入,89 文件、+38069):64 层 gated-delta-net / attention 混合的文本预填,在 1024 / 4096 / 8192 三个长度上全部快于 SGLang;8192 时打分只占 FLOP 的 3.1%、其余是 GEMM,属 FLOP 受限、无流量可省,示例把
analyze --roofline作为决策输入。边界:SGLang 基线由人工给定,其两列口径在 8192 上相差 9.1%,示例需论证取列;此前该仓示例均为解码且仅及 SGLang 零头。 - 重构与契约:实例槽重构(#232,10-15 合入)把模块子节点改为显式共享实例槽(重复层共享定义、保留独立资源与实例身份),放置地址按 SSA 叶子共享、活跃区间按定义缓存并在调用点翻译,删除旧内联机制与
Program,源码全套 1674 通过 / 11 跳过;lowering 与 printer 契约(#233,10-15)保留作者区域、存储定义、更新窗口与分阶段所有权,指令准入在推断、候选生成、lowering 与 TIR 验证之间共享。10-16 追加 analysis 访问上下文(#238)与 issue 237 的分配/lowering 修正(#241)。 - 工程化入口:scaffold 独立工程(#227,10-13 合入,47 文件、+1356/-368)按 Target codegen 生成可编辑的独立 CMake 工程(参数视图、特化分发、启动器、主机注册与占位 CUDA 体),codegen 产物与链接拆为六个模块文件;schedule CLI 加固(#231,10-14 合入)统一指令匹配器、新增
tf.empty、加schedule finalize --dim,冻结全套源码 1617 通过 / 11 跳过。 - 发布与在审:发布页仍为 v0.0.2(09-10),全周无新版本;开放 PR 1→2→1→3→2。外部审计侧,第三方
tilelang-cuda-skills于 10-15 把 TileFoundry vendor 为子模块(固定 main 与 cutlass 提交),公开标注compile()路径为实验性并列出 4 个 emitter bug(Cast 发射、fill 张量后缀、tensor_view 坐标等)与 2 个运行时 bug(reduce-workspace 收尾、warps_per_group 分派),同时指出缺失 host TMA tensor-map emitter,并划出受支持路径(HIR 生成、check / analyze、schedule finalize、TileLang 内核)。
TileSight、TileRT 与采用方
- TileSight:10-13 建立 dev 分支并出现分支级推送,公开分支含 main 与 tilelang(后者含 TileLang 前端提交),无发布内容、开放 PR 0。作为 09-23 开源的瓦片中心性能模型,本周处于分支整理期。
- TileRT:全周无推送(最近一次 09-24),发布页仍为 v0.1.5.post2(08-06);本周无发布 PR、无 benchmark 文档更新。
- TileKernels:v2.1.0 于 10-14 合入(详见要闻),10-14 之后至本报告编制无新推送,开放 PR 计数回落。版本要求(TileLang 0.1.15 / PyTorch 2.13 / Ascend 950 / CANN 9.2.0)使它成为「0.1.15 基线」的第一份下游硬约束。
- FlashQLA:sm120 修复(#48,10-14)恢复
block_DV=64并补做aggregate_card_state,RTX 5090 + tilelang 0.1.15 + CUDA 13.0 全部测试通过;PR 提出后续拟探索节点内 / 跨节点上下文并行。
五、社区、教程与活动
PTO-ISA:tilelangGym 新仓与 SuperNPUBench 的发布节奏
- tilelangGym 于 10-14 建立,定位「测试不同 benchmark / microbenchmark 的 PTO lowering 与性能敏感度」;10-15 七笔(06:47 至 07:59)把独立 SimtVF 标准测试套件成型:SV / CF / SP 加 PTO-DSL 的种子提交与版本 pins、回归(#1),脚本与布局整理、SV1 SIMD keynote 与单张 PASS 表(#3),remat 顶点全 PASS 表与逐行 Simt 对 SIMD-DSL 对照(#4),并把 SIMD 约 85% 与 Simt IPC / bubble / ILP 分析链入 ANALYSIS——为 PTO lowering 的性能敏感度分析建立数据台。
- SuperNPUBench 全周高频:10-12 四笔提交(rms_norm 16 KiB 瓦片分解 R=8192 归约、gfsim RES_CHECK 文件 IO 锁步安全、matmul_test 简化 4-PE、#225 合并)与两个发布(ops-20261012、ops-20261013,后者按 suite 拆 2 小包、含全部编译产出 ELF);10-13 发布 ops-20261013.2 并推进 matmul_test 直连 CUBE TCVT、topk 可移植 gfsim fixture;10-14 三笔归档(rms_norm_dynamic_simt_32k、group_norm_grad 静态内核进 V0、matmul reuseB 内核);10-15 组归一化反向单趟归约与 gamma/beta 均衡(#245)。
- tilelang-puzzles-ascend 的
simt_compare分支全周推进:T.reduce_max对串行扫描实测快 8.94 倍,12 个内核改用T.reduce_absmax折叠并重测。
第三方工具与资料:从列表仓到外部审计
- tilelang-debugger:10-10 两笔(统一源码访问捕获 + 带链接证据的离线 HTML 调试报告),在 TileLang 0.1.12 与 0.1.15 上各跑 117 项 CPU / TIR 测试、H200 上 42 次采集与 66 对基线 / 插桩发射;本周(10-12 起)无新推送。
- easy-tilelang 教程仓:10-13 推送「adapt tilelang new ascend」,含 Ascend A5 的 lower-trace 编译观察等条目;10-14 之后静默。
- TileMC(面向 manycore 架构的第三方目标):10-14 一笔,扩展 FP32 支持、panel 映射与工作负载缩放。
- BM1690 流水线仓(算能 TPU):10-14 十八笔、10-15 六笔板级日志——Matmul / RMSNorm 板级冒烟、Split-K 板级结果、RoPE 联调与启动超时诊断、驱动重初始化与退出语义、7.0.0-34 内核主机构建等,属个人实验仓的联调记录。
- SwHwCodesign 资料线未成形:10-11 建立的 TileLang-Awesome 于 10-12 已不可访问(404,搜索无同名新仓),同账号另建六个 AI-* 占位仓,至 10-14 公开仓仅剩 2 个(AI-Compiler、AI-Infra)且无推送。
- tilelang-cuda-skills:10-15 对 TileFoundry 的首批外部审计(内容见第四章),是「工具链被真实使用」的第一个公开样本。
媒体与学术:报道潮之后的持续静默
- 10-10 至 10-11 的中文深度解读(知危,风闻与知乎两版)以 TileLang 为叙事核心,讲 Tile 抽象、三档编程接口与自动调优缓存机制,并给出昇腾适配对标 Ascend C「性能几乎无差别(A2 / A3)」的对比;10-11 境外聚合稿(tech-insider)属转引性质,非独立信源。
- 10-12 起连续四个窗口 Google News 中英文多组查询零新增,Hacker News 无相关条目,arXiv 无新预印本(最近一篇仍为 07-24 的 TileSight 性能模型论文)。媒体侧在 09-29 至 10-11 的报道潮之后进入平静期,本周主体内容全部来自 GitHub 组织动态。
六、趋势观察
- 发布静默期的「静态准备」在四个方向同时到位:TileOPs 的 Op 基类重构自称发布前收口、夜间底盘六份快照零失败增长、TileFoundry 的示例从解码跨到预填全长度超越、多后端 v0.1.15 对齐——主仓自 09-30 已 16 天无标签,但下一版本的静态准备在本周基本就位。是否发布,取决于评审闸门而非内容。
- 多后端从「版本对齐」走向「同语义收口」:海光完成 v0.1.15 全量升级并另立 v0.1.12 维护线,沐曦镜像追平后转入 dev 合流,昇腾把重心放在文档事实核查、缓存语义与真机验证(#1900),摩尔线程是唯一缺口。四家的差距已从「有没有后端」变为「语义是否一致」。
- ROCm 成为第二条与新模型发布同步的线:连续四日 DeepSeek NSA、mHC、KDA 与 GLM-5.3 相关内核落地,技术债按设备预算(波前 64 lane、MFMA 输入类型、LDS/VGPR 上限)逐笔清偿;与 SM121 / SM89 的多代并修合流后,后端拓扑已从「英伟达优先」转为「多设备同时维护」。
- TileOPs 的性能工程从「重点算子攻坚」转为「全算子逐日刷榜」:单日 8 至 11 笔、每笔带外部对照(torch-compile / cuBLASLt / quack / FlashInfer / FLA),残余落后被公开写进 PR(lm-head 两行 0.909x、GLA 反向明确不主张加速)——把「落后」量化并留档,比单点提速更能说明工程纪律。
- 昇腾的治理口径从「全部通过」转向「断言精确」:三笔文档核查因低优先级用例的异常正则不匹配被整批回滚再删减重做,暴露 PR 阶段不跑
low_priority、每日全量兜底的两级 CI 结构;这类回滚不是退步,而是把「预期失败」当成技术债清算的信号。 - TileFoundry 从编译器工具转向可交付产品形态:scaffold 独立工程、实例槽与 lowering 契约、跨脊线的预填示例论证,重构落地的同一窗口即被第三方 vendor 并列出缺陷台账——产品化的典型轨迹是「底盘重构之后,接口契约收敛与缺陷公开化」。下一期需跟踪这些 emitter / 运行时缺陷是否被吸收。
- 下一期关注点:主仓 v0.1.16 是否发布(已静默 16 天);SM120 块缩放 MMA 双线(#3495 / #3099)能否合入;cumsum warp 特化修复(#3536)与 PTO 家族七笔的推进;昇腾 #1900 的两笔提交与 A5 真机验证排期;海光分布式 IPC 是否出现首个合入的多卡 PR;TileOPs 已出现发布流水线整理类提交(周末),其是否引出首个正式版本;TileFoundry 的外部缺陷清单是否被吸收。
附录:素材与核查说明
- 素材:本刊每日 TileLang 动态日报 5 期(10-12 至 10-16),对应采集窗口 10-10 07:00 至 10-16 07:00(北京时间);周期内 10-16 白天与 10-16 夜间(北京时间 10-17 凌晨)的动态,经 168 小时窗口补充核查后纳入本报告;10-17(周六)当日的组织动态并入下周一日报。
- 补充核查(168 小时窗口,截至本报告编制):tile-ai 组织 29 个仓库中窗口内 12 个有推送(tilelang、TileOPs、TileFoundry、tilelang-ascend、tilelang-hygon、tilelang-metax、tilelang-mlir-ascend、TileSight、TileOPs-nightly、tilelang.github.io、TileOPs.github.io、tilelang-sunrise);PR 合入与新建数按 10-12 至 10-16 北京时间窗口检索实测;夜间流水线六份快照的
test_results.xml与bench_results.xml逐份复算(用例、失败、错误、跳过、套件数)。 - 口径与边界:性能数据均为 PR 说明或发布说明的自述(H100、H200、B300、RTX 5090、RTX 6000 Ada 等),本刊无对应硬件、未做独立复现;夜间快照为单次流水线记录,非横向评测;发布时间以各仓库标签与发布页复核为准,适配仓的发布页与分支推进可能存在滞后;TileKernels 的昇腾性能由其自身 README 声明,非第三方评测。
- 下一期周期建议:2026-10-19 ~ 2026-10-23。
信源核查表
| 信源 | 核查结果 |
|---|---|
| tile-ai 组织(29 仓库) | 168 小时窗口内 12 个仓库有推送;本周 PR 合入 109 笔、新开 158 笔;当前开放 PR 532 笔 |
| 主仓 tilelang | 本周合入 30 笔;开放 PR 170→168→161→160→165;发布页与标签全周无新增(最新 v0.1.15,09-30) |
| TileOPs | 本周合入 52 笔;开放 PR 9→19→26→17→16;Op 基类重构、MoE 双算子、注意力三连、foundry 性能批同周落地 |
| TileOPs-nightly | 六份快照,正确性 1032→1134、基准 1947→1989(52→55 套件),全部零失败零错误 |
| TileFoundry | 本周合入 10 笔(#223、#221、#227、#229、#231、#232、#233、#236、#238、#241);开放 PR 1→2→1→3→2;发布仍为 v0.0.2(09-10) |
| 昇腾(ascend / mlir-ascend) | ascend:合入 8 笔(#1701、#1730、#1612、#1676、#1681、#1888 至 #1890);开放 PR 230→237。mlir-ascend:合入 2 笔(#190、#202);开放 PR 32 笔 |
| 海光 / 沐曦 / 摩尔线程 | 海光:合入 3 笔(#20、#21、#22)、#19 与 distributed 分支在审/推进;沐曦:合入 2 笔(main 镜像与 dev 合流);摩尔线程:09-30 后无活动、开放 PR 6 笔 |
| Sunrise | 仅 Dependabot 两笔 CI 依赖升级在审(#7、#8);开放 PR 3 笔 |
| 采用方(TileKernels / FlashQLA) | TileKernels v2.1.0(#38)与 FlashQLA sm120 修复(#48)均于 10-14 合入;10-14 之后无新推送 |
| TileSight / TileRT | TileSight 10-13 建 dev 分支并推送、无发布;TileRT 全周无推送(最近 09-24) |
| PTO-ISA 与社区第三方 | tilelangGym 新仓(10-14)与七笔 SimtVF 套件;SuperNPUBench 两笔发布加多笔归档;BM1690 24 笔板级日志;tilelang-cuda-skills 1 笔外部审计;TileMC 1 笔;easy-tilelang 1 笔;tilelang-debugger 与 SwHwCodesign 无窗口内推送 |
| 媒体与学术 | Google News 自 10-12 起连续四个窗口零新增;Hacker News 无相关条目;arXiv 无新预印本 |
主要引用
- 主仓 SM121 块缩放 MMA(#3430) — https://github.com/tile-ai/tilelang/pull/3430
- 主仓 GEMM dtype 白名单(#3245) — https://github.com/tile-ai/tilelang/pull/3245
- 主仓 ROCm NSA(#3152) — https://github.com/tile-ai/tilelang/pull/3152
- 主仓 ROCm flash-decoding 瓦片(#3149) — https://github.com/tile-ai/tilelang/pull/3149
- 主仓 DeepSeek mHC 前/后内核(#3156) — https://github.com/tile-ai/tilelang/pull/3156
- 主仓 safe-gated KDA 解码(#3249) — https://github.com/tile-ai/tilelang/pull/3249
- 主仓 ROCm scan 波前修复(#2326) — https://github.com/tile-ai/tilelang/pull/2326
- 主仓 PTO 执行路由(#3316) — https://github.com/tile-ai/tilelang/pull/3316
- 主仓昇腾示例目标选择(#3474) — https://github.com/tile-ai/tilelang/pull/3474
- 主仓独立 AIV 执行(#3485) — https://github.com/tile-ai/tilelang/pull/3485
- 主仓 PTO DeepGEMM codegen(#3330) — https://github.com/tile-ai/tilelang/pull/3330
- 主仓跨 warp allreduce 修复(#3494) — https://github.com/tile-ai/tilelang/pull/3494
- 主仓 CPU OpenMP 网格并行(#3109) — https://github.com/tile-ai/tilelang/pull/3109
- 主仓 C 系 codegen 共享基类(#3289) — https://github.com/tile-ai/tilelang/pull/3289
- 主仓 Windows NVCC 走 MSVC(#3510) — https://github.com/tile-ai/tilelang/pull/3510
- 主仓 JIT 不可哈希参数(#3417) — https://github.com/tile-ai/tilelang/pull/3417
- 主仓 cumsum warp 特化修复(#3536) — https://github.com/tile-ai/tilelang/pull/3536
- 主仓 SimdVF 多缓冲 T.fill(#3455) — https://github.com/tile-ai/tilelang/pull/3455
- 主仓发布页 — https://github.com/tile-ai/tilelang/releases
- TileOPs Op 基类重构(#2520) — https://github.com/tile-ai/TileOPs/pull/2520
- TileOPs masked FP8 分组 GEMM(#2521) — https://github.com/tile-ai/TileOPs/pull/2521
- TileOPs DeepSeek 分组 top-k(#2535) — https://github.com/tile-ai/TileOPs/pull/2535
- TileOPs varlen GQA 汇聚(#2537) — https://github.com/tile-ai/TileOPs/pull/2537
- TileOPs 分页 GQA 全特性(#2516) — https://github.com/tile-ai/TileOPs/pull/2516
- TileOPs PagedKVCacheWrite(#2539) — https://github.com/tile-ai/TileOPs/pull/2539
- TileOPs MLA 解码 SM89(#2511) — https://github.com/tile-ai/TileOPs/pull/2511
- TileOPs softmax 寄存器行(#2570) — https://github.com/tile-ai/TileOPs/pull/2570
- TileOPs coop2 逐行存储(#2566) — https://github.com/tile-ai/TileOPs/pull/2566
- TileOPs 分组 GEMM 存储重叠(#2582) — https://github.com/tile-ai/TileOPs/pull/2582
- TileOPs 夜间快照分支 — https://github.com/tile-ai/TileOPs-nightly/commits/snapshots
- TileFoundry 稀疏解码贯通 TIR(#223) — https://github.com/tile-ai/TileFoundry/pull/223
- TileFoundry Qwen3.8-27B 预填(#229) — https://github.com/tile-ai/TileFoundry/pull/229
- TileFoundry 实例槽重构(#232) — https://github.com/tile-ai/TileFoundry/pull/232
- TileFoundry lowering 契约(#233) — https://github.com/tile-ai/TileFoundry/pull/233
- 昇腾 NSA Decode(#1701) — https://github.com/tile-ai/tilelang-ascend/pull/1701
- 昇腾 NSA Bwd(#1730) — https://github.com/tile-ai/tilelang-ascend/pull/1730
- 昇腾 bitwise_rshift 重做(#1893) — https://github.com/tile-ai/tilelang-ascend/pull/1893
- 昇腾 dav-3510 gemm_v0(#1900) — https://github.com/tile-ai/tilelang-ascend/pull/1900
- MLIR 昇腾 NPUIR 控制(#190) — https://github.com/tile-ai/tilelang-mlir-ascend/pull/190
- 海光 v0.1.15 升级(#20) — https://github.com/tile-ai/tilelang-hygon/pull/20
- 海光 Distributed(#19) — https://github.com/tile-ai/tilelang-hygon/pull/19
- 海光块缩放异步拷贝布局(#21) — https://github.com/tile-ai/tilelang-hygon/pull/21
- 沐曦仓库 — https://github.com/tile-ai/tilelang-metax
- 摩尔线程仓库 — https://github.com/tile-ai/tilelang-musa
- TileKernels v2.1.0(#38) — https://github.com/deepseek-ai/TileKernels/pull/38
- FlashQLA sm120 修复(#48) — https://github.com/QwenLM/FlashQLA/pull/48
- PTO-ISA/tilelangGym — https://github.com/PTO-ISA/tilelangGym
- SuperNPUBench — https://github.com/PTO-ISA/SuperNPUBench
- tilelang-cuda-skills(TileFoundry 外部审计) — https://github.com/sablin39/tilelang-cuda-skills
- BM1690 pipelines — https://github.com/arcflute/tilelang-tpu-bm1690-pipelines
- tilelang-debugger — https://github.com/superAngGao/tilelang-debugger
数据来源:每日 TileLang 动态日报(2026-10-12 至 2026-10-16 共五期)及本周期实测核查;报告周期为 2026 年 10 月 12 日至 10 月 16 日。