调研窗口:2026-10-22 07:00 ~ 2026-10-23 07:00(北京时间,24 小时)。

信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon 分支 / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 三个文档站)+ 合入与在审真值检索(is:merged 与 is:open 两组,按仓库分组统计)+ 提交级 diff 与发布页 tag 核查 + 采用方仓核查(TileKernels / FlashQLA)+ 社区新仓检索 + Google News 中英文多组查询(走代理)+ Hacker News + arXiv + 产业媒体报道。


本期索引

  • 今日重点:35 笔合入里主仓把重心放到分析器生命周期,深寻 V3.2 的 top-k 选择器越界写被定位到 CUDA_ERROR_ILLEGAL_ADDRESS;昇腾 NZ 布局打通部分 UB→L1 拷贝(18 文件);CANN 社区新开 ops-tilelang 算子仓
  • 一、核心项目进展
    • 1.1 主仓:向量化重写前克隆 analyzer(#3612)(10-22)
    • 1.2 主仓:并行循环切分后旧 analyzer 绑定失效(#3226)(10-22)
    • 1.3 主仓:deepseek_v32 top-k 选择器候选缓冲越界(#3280)(10-22)
    • 1.4 主仓:CUDA 16 位反三角与二进制数学降级补齐(#3172)(10-22)
    • 1.5 主仓:sm8x 稀疏元数据非双射形状改为拒绝(#3125)(10-22)
    • 1.6 主仓:eager 前端保留 Python 源码位置(#3558)(10-22)
    • 1.7 在审:pass 能力检查、无符号 mulhi 与 assume 谓词(10-22)
  • 二、多后端适配(昇腾 / 海光 / 沐曦 / 摩尔线程)
    • 2.1 昇腾:NZ 缓冲支持部分 UB→L1 拷贝(#3530)(10-22)
    • 2.2 昇腾 PTO:MAD 方向谓词指数增长致编译停滞(#3581)(10-22)
    • 2.3 昇腾 PTO:SIMT 向量局部变量按 lane 分配(#3523)(10-22)
    • 2.4 昇腾仓:cann-bench 补 Dilation2D,同日回退 depthwise_conv_2d(#1916 / #1919)(10-22)
    • 2.5 昇腾仓在审:sparse_mla 前向与反传(#1929 / #1928)(10-22)
    • 2.6 海光:合并分支移除 enable_register_pipeline(10-22)
    • 2.7 沐曦 / 摩尔线程 / MLIR-Ascend:窗口内无推送(10-22)
    • 2.8 产业侧:SCMP 转载 SemiAnalysis 报告,TileLang 为六仓中唯一可用(10-22)
  • 三、生态与采用方(TileOPs / TileFoundry / 算子库)
    • 3.1 TileOPs:FP8 分页 warp 特化 kernel 共享内存别名致 NaN(#2708)(10-22)
    • 3.2 TileOPs:dense prefill 改为持久化遍历,D64 三消费者(#2706)(10-22)
    • 3.3 TileOPs:SM90 dense GQA decode 走 warp 特化(#2694)(10-22)
    • 3.4 TileOPs:因果滑窗用一个 KV head 的两个 key tile 服务整组 query head(#2696)(10-22)
    • 3.5 TileOPs:分页 decode 自动调优此前共用一张 block table(#2703)(10-22)
    • 3.6 TileOPs:masked softmax 的 FMA 溢出与 varlen RoPE 奇数半宽(#2699 / #2700)(10-22)
    • 3.7 TileOPs:FP8 grouped GEMM 分级暂存 scale 与 gate/up 乒乓 warp group(#2704 / #2693)(10-22)
    • 3.8 TileOPs:谓词广播与布尔操作数改 16 字节向量读(#2702 / #2697)(10-22)
    • 3.9 TileOPs:SM89 三条路径同日推进(#2705 / #2691 / #2690)(10-22)
    • 3.10 TileOPs:归约折叠单次发射与空 DSA 选择校验(#2689 / #2698)(10-22)
    • 3.11 TileOPs:基准参考改在 float32 上计算(#2692)(10-22)
    • 3.12 TileFoundry:安装 sm_89 架构与 RTX 6000 Ada 设备(#259)(10-22)
    • 3.13 TileFoundry:gmem 计价改按足迹与运行 mesh,并绑定被调方维度(#258)(10-22)
    • 3.14 TileFoundry:设备拒绝启动 kernel 时必须抛错(#256)(10-22)
    • 3.15 分发版:Sunrise 打出 0.1.15+sunrise.1.2.0 发布 PR(#9,在审)(10-22)
    • 3.16 采用方:TileKernels 与 FlashQLA 窗口内无推送(10-22)
  • 四、社区、教程与活动
    • 4.1 CANN 社区开设 ops-tilelang 算子仓(10-22)
    • 4.2 tilelang/agent:把算子优化写成 agent 的 SKILL 工作流(10-22)
    • 4.3 社区新仓:Sail 包装、远程自动调优 agent 与教程仓(10-22)
    • 4.4 文档与版本:文档站更新三次,VERSION 仍 0.1.15(10-22)
  • 五、趋势观察

今日重点

本窗口 tile-ai 组织合入 35 笔(TileOPs 20、tilelang 9、TileFoundry 3、tilelang-ascend 2、TileFoundry 文档站 1),另有 71 笔在审 PR 在窗口内更新。主仓九笔的重心从上周的”形状与几何校验”转向分析器生命周期:向量化重写前克隆 analyzer(#3612,单文件 2 行)、并行循环切分后旧绑定失效(#3226,6 文件)、eager 前端保留 Python 源码位置(#3558)——三笔都不会让 kernel 算错,而是让编译期在改写后读到过期状态。数值侧最值得记的是 #3280:deepseek_v32 的 top-k 选择器把 radix 候选暂存在固定 [2, 4096] 共享缓冲,却没有约束第一轮选出的阈值桶大小,低尾数相同的 FP32 分数或全等输入会越界写,崩在 CUDA_ERROR_ILLEGAL_ADDRESS。昇腾侧 #3530 是当日体量最大的一笔(18 文件 +993/-256),把压缩 NZ 填充保留到存储归一化,为部分 UB→L1 拷贝显式给出 C0 组行距;#3581 修掉 PTO 上一个会让编译停滞十分钟以上的谓词表达式指数增长。生态侧两条新线:CANN 社区 10-20 建了 ops-tilelang 算子仓、10-22 首次推送修复;Sunrise 分发版打出 0.1.15+sunrise.1.2.0 的发布 PR,把 TANG 后端 wheel 纳入分发。


一、核心项目进展

1.1 主仓:向量化重写前克隆 analyzer(#3612)

合于 10-22 23:43,只改向量化那一处源码,+2/-1。循环向量化在改写 IR 时复用传入的算术分析器,而分析器缓存着改写前的常量与区间约束;一旦向量化改变了循环边界或索引表达式,后续 pass 拿到的仍是旧结论。修法是在进入向量化重写前复制一份分析器,改写只作用于副本。这一笔没有随附回归测试,价值在于它从另一个入口堵住与 #3226 同源的问题——同一天那一笔处理的是并行循环切分后的失效绑定,两笔合看,分析上下文的生命周期被当成独立缺陷类在清。

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

1.2 主仓:并行循环切分后旧 analyzer 绑定失效(#3226)

6 文件 +55/-19,涉及 loop_partition、loop_vectorize、lower_tile_op 三处 pass 及其头文件。并行循环被切分后循环体的迭代变量被重绑定,但 pass 仍按切分前的变量查分析器,取到的是旧变量的区间与整除信息,简化与越界判定因此失真。修复把分析上下文按 pass 隔离,并简化了循环分析快照的初始化路径,配套加了一组最小化的并行循环分析回归用例。提交者中含 Codex 联署。

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

1.3 主仓:deepseek_v32 top-k 选择器候选缓冲越界(#3280)

2 文件 +91/-10,修 #1351。tl_topk_impl 把 radix 候选装进固定的 [2, 4096] 共享缓冲,但没有任何东西约束第一轮 radix 选出的阈值桶不超过 4096 个元素——FP32 分数低位尾数相同、或输入全等时,原子追加会写出缓冲末尾并崩在 CUDA_ERROR_ILLEGAL_ADDRESS。改法是在 tl_topk 包装层申请 (batch, 2, seq_len) 的全局 workspace,让两个 radix 缓冲都由 seq_len 兜住;公开的 tl_topk() 签名不变。随附 79 行复现用例,按 #1351 的输入构造。

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

1.4 主仓:CUDA 16 位反三角与二进制数学降级补齐(#3172)

3 文件 +463/-8。asin、acos、asinh、acosh、atanh、log1p、atan2、hypot、copysign、ldexp、nextafter 缺 CUDA 专用下降规则,通用规则只处理 float32/float64,float16 与 bfloat16 调用在 codegen 阶段无人接管。修复把它们注册进 CUDAMath 并补 half_t / bfloat16_t 重载:十个算子按 float32 求值再转回,nextafter 改为按 16 位编码做整数比较步进,以便在 fast math 下保住 bfloat16 次正规数。float32/float64 的函数名保持不变,common.h 侧新增 182 行模板实现。

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

1.5 主仓:sm8x 稀疏元数据非双射形状改为拒绝(#3125)

单文件 tilelang/layout/gemm_sp.py +6/-0,修 #2606。make_cutlass_metadata_layout_sm8x 的 interleave 只在 M 为组大小整数倍且 K 为偶数时才是双射;不满足时原先静默生成错误元数据,稀疏 GEMM 结果无提示地错。现在两个条件都改为显式拒绝。原先随 PR 加的稀疏元数据形状测试在评审中被要求移出,未进入主分支。

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

1.6 主仓:eager 前端保留 Python 源码位置(#3558)

4 文件。eager 前端把 Python 的 for / with 语句降成生成式调用后,报错与调试映射回的是生成代码的行号。修复把生成的循环调用映射回 for 语句、上下文包装映射回 with 语句,并补上可迭代参数解包的位置;与 #3125 一样,评审要求移出随附的源码位置回归测试。

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

1.7 在审:pass 能力检查、无符号 mulhi 与 assume 谓词(10-22)

主仓在审队列在窗口内更新 50 笔。值得跟的几条:[Transform] Locate VerifyBufferInit reads outside tile ops(#3620)、[Codegen] Add unsigned mulhi intrinsic(#3617)、[Transform] Keep assumption checks with local dependencies in scope(#3611)与 [Transform] Preserve tl.assume predicates that simplify to true(#3610)、[CPU] Support clz in C and C host code generation(#3609);昇腾侧 [BugFix][Ascend] Reject odd FP4 view geometry in GM->L1 copies(#3608)与 [Ascend] Support exact placement of versioned on-chip buffers(#3607)同日更新。另有 #3613 明确标注”请勿评审”的原生 argmin/argmax 归约草案。

来源:https://github.com/tile-ai/tilelang/pulls


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

2.1 昇腾:NZ 缓冲支持部分 UB→L1 拷贝(#3530)

当日体量最大的一笔:18 文件 +993/-256。压缩 NZ 布局的填充此前在逻辑布局阶段就被抹平,做部分 UB→L1 拷贝时无法表达源与目的各自的 C0 组行距。修复把紧凑 NZ 填充保留到存储归一化阶段、再按容量精确物化一次,拷贝区域下降时显式给出以 32 字节为块的源/目的 C0 组行距,并把行距贯穿到 scatter 模板、Ascend/PTO 的 DMA codegen 与访问足迹;同时加了几何、对齐与越界校验,保留物理视图所有者,调用方迁移到逻辑分配形状加显式行距。src/ascend/op/nd2nz.h 新增 83 行。

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

2.2 昇腾 PTO:MAD 方向谓词指数增长致编译停滞(#3581)

2 文件 +48/-5。带重复可变守卫的 PTO kernel 可以花十分钟以上做分析,即使里面没有矩阵乘——这是 NPU 执行之前的编译期停顿。根因是 PTOMmadDirectionAnalyzer 在守卫读取可变内存时保守保留进入状态,若两条分支都保持 MAD 方向,Merge 收到同样的谓词 S,而实现总是构造并化简 S || S;算术化简器在复现用例里消不掉这个重复,连续分支便让谓词表达式按倍数翻番。修复在两侧谓词相同时跳过重建,并补 20 行定向测试。

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

2.3 昇腾 PTO:SIMT 向量局部变量按 lane 分配(#3523)

4 文件 +273/-29。SIMT 向量局部变量被当标量 PTODSL 缓冲分配,没算 lane 数,却用 Python 下标访问;AllocatedBufferValue 不支持逐项赋值,循环索引写入在 PTODSL AST 重写阶段就更早失败。修复按每个向量槽分配足够标量元素并把槽下标翻译成标量偏移,复用既有的连续 load/store 路径,保持 AscendC 的向量局部语义,不改循环展开、不加新的下降 pass。补 float32x2 / float32x4 两组 codegen 测试,共 209 行。

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

2.4 昇腾仓:cann-bench 补 Dilation2D,同日回退 depthwise_conv_2d(#1916 / #1919)

#1916 合入 Dilation2D 前向(cann-bench level3 任务):NHWC、fp16/fp32、任意 kernelSize/stride/rate 与 SAME/VALID,越界 tap 的贡献取 -inf + filter 与 golden 逐位一致,NaN 输入传播、+inf filter tap 使越界贡献变 NaN;实现为纯 Vector kernel 加 Python wrapper,按昇腾拷贝引擎的 32B 行跨距对齐约束派发,grid 按几何自适应而不用写死核数。同一天 #1919 把 10-21 刚合入的 depthwise_conv_2d(#1915)整体回退,提交信息未给原因,示例仓处在”进一个、退一个”的状态。

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

2.5 昇腾仓在审:sparse_mla 前向与反传(#1929 / #1928)

#1929 新增 sparse_mla_fwd(DSA 稀疏 MLA 前向:varlen packed 随机行 gather、在线 softmax、QK/PV 双 GEMM 融合),落 examples/dsa_sparse_finetune/,双 kernel(fast 特化加 portable 通用)配合 host 分派;fast 路径的适用域写死为 H=128 且 G=1、dim/tail 为 512/64、topk 为 256 的整数倍、NI 在 1 到 6 之间,域外走通用路径。#1928 新增 sparse_mla_bwd(dQ GEMM 链、P/dP 向量链、dKV 随机行散射),落 examples_experiment/dsa/ 并登记进算子测试清单,前向结果 k1 已折进本算子。昇腾仓在审合计 11 笔。

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

2.6 海光:合并分支移除 enable_register_pipeline(10-22)

gcy_merge_main 于 10-22 11:11 推送一笔 6 文件 +10/-24 的清理,把 enable_register_pipeline 开关整体删掉,涉及 HCU 寄存器软流水的注入与流水规划两个 pass、pipeline_utils.h 以及语言层的 loop 扩展,示例与性能脚本同步改调用。删开关通常意味着寄存器流水路径已经默认生效、双轨维护不再划算;但 main 分支仍停在 10-20 的 preserve constant K origins in hoisted MLS windows(#25),这批改动还在合并分支上等落地。

来源:https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main

2.7 沐曦 / 摩尔线程 / MLIR-Ascend:窗口内无推送(10-22)

tilelang-metax 最后推送停在 10-16,tilelang-musa 停在 09-30,tilelang-mlir-ascend 停在 10-19,三条适配线在窗口内均无提交。昇腾、海光两条线承担了本窗口全部后端改动。

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

2.8 产业侧:SCMP 转载 SemiAnalysis 报告,TileLang 为六仓中唯一可用(10-22)

南华早报 10-22 发文《Jensen’s nightmare》,转述 SemiAnalysis 周四发布的华为-DeepSeek 报告:DeepSeek 最新软件发布加上对昇腾 950PR 的更多实验,说明与美国的差距在缩小,并判断 DeepSeek 到 2027 或 2028 年将”主要”依赖华为”已相当成熟”的技术栈。该报告原文即上一期已收录的那份,其中”六个获得昇腾支持的仓库中只有 TileLang 可用”的判断未变,另提到昇腾 950 后端是并入 TileLang 主仓而非停留在独立适配层。属同一信源的二次传播,不构成新的技术进展,但说明该判断正在进入主流媒体口径。

来源:
https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report


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

3.1 TileOPs:FP8 分页 warp 特化 kernel 共享内存别名致 NaN(#2708)

2 文件 +80/-30。GQAPagedFP8WSKernel 在”七个单 token 请求排在 300 token 请求之前、4099 行 FP8 cache、48 行页”的组合下返回 NaN。根因是共享内存规划器把 warp 角色各自的缓冲读成不同时段活跃,于是把 split handover 放到了仍被 TMA 填充的 key 或 value 环地址上,handover 的 running maxima 也被压到环标志位。修复把这几块缓冲在共享内存里显式隔开,并补 24 行回归用例(qwen3-fp8-ragged 口径)。

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

3.2 TileOPs:dense prefill 改为持久化遍历,D64 三消费者(#2706)

6 文件 +590/-5,新增 GQADensePersistentWSKernel(对外名 gqa_dense_persistent,优先于 gqa_dense):query tile 数超过 SM 数时,一个 CTA 占一个 SM 交替轮次走完自己的 tile,并在当前 tile 的 epilogue 期间装载下一个 tile 的 query。D64 上每个 tile 是 192 行 query 摊给三个 64 行消费者 warpgroup,K/V 重读减少三分之一;D128 保持 128 行。

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

3.3 TileOPs:SM90 dense GQA decode 走 warp 特化(#2694)

新增 GQADecodeWSKernel:TMA 生产者 warp 填两槽 K/V 环,一个 WGMMA warpgroup 维持 head 组的在线 softmax;split 数按一次填满常驻 CTA 槽来定,并保证每个 split 至少四个 KV tile,单 split 时在 CTA 内归一化并直接写输出、不再过 combine kernel;KV head 在 CTA 间变化最快。

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

3.4 TileOPs:因果滑窗用一个 KV head 的两个 key tile 服务整组 query head(#2696)

新增 GQADenseSlidingWindowBandKernel:GQADenseFwdOp 原先对每个 query head 都重载同样两个 key/value tile,一个 CTA 只做一个(query tile, head);新 kernel 一个 CTA 对(query tile, KV head)把 band 的两个 key/value tile 只装一次,依次对该组 query head 做归约,同时让下一个 head 的 query tile 在飞。针对的是滑窗行数不超过 128 个 key 的场景。

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

3.5 TileOPs:分页 decode 自动调优此前共用一张 block table(#2703)

4 文件 +24/-13。分页 decode 的 autotune 供给让每个请求拿到同一张 block table,等于只测了一个请求的页、且其页常驻 L2,候选排名反映的是被服务批次实际付出的流量。实测在 1500 MHz H200 上,MHA decode 调优器在 opt-1.3b-b16 上选了 block_N 128(冷 80.0 us)而放弃 block_N 64(72.5 us),掉到 flashinfer 的 0.994。修复让调优按请求各自读不同 cache 页。

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

3.6 TileOPs:masked softmax 的 FMA 溢出与 varlen RoPE 奇数半宽(#2699 / #2700)

#2699(4 文件 +54/-6)修 dense GQA 滑窗前向的 NaN 行:因果窗下的 128x64 调优瓦片让每个 query tile 从首行窗口开始走,最后几行会连续遇到两个全屏蔽 key tile;第一块之后两个 running maxima 都在 -1e38 钳位上,fast math 把 prev * scale - curr * scale 融合成 FMA,舍入误差(约 2 的 99 次方)溢出成 NaN,修法是在缩放重定标前先减 running maxima。#2700 修 varlen query RoPE 在奇数半旋转宽度(rotary_dim 10、14 等)下的循环边界:并行度 rows 乘 chunks 不能被旋转分支的线程数整除时 TileLang 丢掉了循环上界,warp 特化路径的后续 query 行算错,通用路径(dim 256)直接挂死。

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

3.7 TileOPs:FP8 grouped GEMM 分级暂存 scale 与 gate/up 乒乓 warp group(#2704 / #2693)

#2704 新增配置键 stage_scales:让生产者用 cp.async 把每一级的 A/B scale 拷进共享内存并在落地时到达该级 barrier;stage_scales=0 仍生成旧 kernel。背景是 masked FP8 grouped GEMM 在数学 warpgroup 里从全局读每步 K 的 A scale,环加深会压缩 L1 并让这些 load 停住,ds-v3-e8-up-masked 因此只有 DeepGEMM 的 0.979。#2693 给 GemmTemplate 加乒乓形态:两个数学 warpgroup 各跑完整的 64 行 tile、轮流取 CTA 的 tile,一个组的 epilogue 落在另一个组的 mainloop 下,靠轮转 barrier 排序。

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

3.8 TileOPs:谓词广播与布尔操作数改 16 字节向量读(#2702 / #2697)

#2702 让行广播谓词(Eq/Ne/Gt/Lt/Ge/Le 与 LogicalAnd/Or)把每 1024 列整块的 stride-1 操作数按 16 字节向量经 load16_evict_last 读,部分块与其他 staged body 保留 fragment 拷贝;linear-bias-float32(1024x4096 对 1x4096)六种比较从 0.99 提到 1.06,逻辑算子 1.04。#2697 把比较与逻辑的二值 kernel 中全尺寸操作数改为 16 字节 evict-first 向量读,float32 行此前只在 torch.compile 的 0.99 到 1.00。

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

3.9 TileOPs:SM89 三条路径同日推进(#2705 / #2691 / #2690)

#2705 把 SM89 的 SSD chunk scan 默认流水级数降到两级,减少 scan kernel 的缓冲需求,六个形状的 FP16/BF16 实测取三轮、每轮 200 次 CUPTI 采样的中位数。#2691 在现有规划器选择多分区时拆分 SM89 的 paged DSA key 区间,复用既有 FP32 部分输出与 log-sum-exp 的 merge;单分区时保持 BF16 输出与单 kernel,并把编译目标纳入 merge 缓存标识、为每个目标建独立 JIT builder。#2690 在 SM89 上对 query head 组 32/64、head dim 64/128 的 NSA top-k 选择改用 32 列 tile,保留 64 行 tile 与同一套打分排序实现,八个 FP16/BF16 用例的 GPU busy 中位数降 2.1% 到 15.0%。

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

3.10 TileOPs:归约折叠单次发射与空 DSA 选择校验(#2689 / #2698)

#2689 新增 ReduceFoldKernel,把每个保留索引的 lead 段连续后缀在一次发射里折叠(此前对轴前缀与后缀的归约要两次发射),并让 ReduceEdgeKernel 与 VectorNormEdgeKernel 把行能被 512 线程块在飞持有的调用都交给它;背景是 Sum 在 3d-multidim-float16(0.977)与 mlp-intermediate-float32(0.990)上落后 torch-compile。#2698 修分页 DSA 校验:query 行没有任何有效选中 key 时参考实现返回零输出与负无穷 LSE,而减参考 LSE 会得到 NaN 误差界、把正确结果判错;修复给空行精确的零 LSE 界,并补”全部行皆空”时校验器必须拒绝非零输出与有限 LSE 的控制用例。

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

3.11 TileOPs:基准参考改在 float32 上计算(#2692)

ref_program 此前既当正确性参考、又是 torch 基线计时的对象,许多参考直接以 16 位输入 dtype 调 torch,结果取决于设备侧 torch 每一步怎么舍入。CUDA 掩盖了这个差异,CPU torch 与 torch_musa 不掩盖——GroupNorm 与 InstanceNorm 在 MTT S5000 上失败,而同一 kernel 对 float64 参考是通过的。修复让 ref_program 在 float32 或更宽上计算,torch 基线则按输入 dtype 计时。这是少见的一笔由国产 GPU 后端暴露出、但要先在参考实现层修的问题。

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

3.12 TileFoundry:安装 sm_89 架构与 RTX 6000 Ada 设备(#259)

11 文件 +392/-22。新增 nvidia_sm89.toml 与 nvidia_rtx_6000_ada.toml 两个目标描述(分别 114 与 124 行),docs/spec/target.md 同步补 42 行。这一笔把 TileFoundry 的设备建模从 Hopper 扩到 Ada,与同日 #256 暴露的”按 sm_90 目标编译的测试在 sm_89 单卡上只错值”是同一件事的两面;配套还加了 sm_89 的 e4m3 warp mma 声明(#261,在审)。

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

3.13 TileFoundry:gmem 计价改按足迹与运行 mesh,并绑定被调方维度(#258)

当日改动面最广的一笔:74 文件 +1128/-553,修 #239、#252、#253。原先 132 个 CTA 只探 40 个时,计价按”每个窗口每个位置移动一次”算,under-fill 根本看不出来;只能由被调方声明的维度绑不上;analyze 的输出目标与逐函数内存行含义不清。修复把 check 与 analyze 的维度绑定集中到 instantiate,被调方或子模块声明的维度现可绑定。提交由社区贡献者发起、维护者接手改写。

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

3.14 TileFoundry:设备拒绝启动 kernel 时必须抛错(#256)

7 文件 +69/-11。此前启动器发出 kernel<<<grid, block, smem, nullptr>>>(...) 后从不读启动状态,设备拒绝启动时调用照常返回,输出保持旧内容,错误到不了调用方;在 sm_89 单卡上,按捆绑 sm_90 目标编译的测试只表现为值不对,compute-sanitizer 才显示 cudaLaunchKernel 返回 cudaErrorNoKernelImageForDevice。修复让 CUDA 启动器在读 cudaGetLastError() 后抛错,CPU 侧 module 同步改,规格文档与测试一并更新。

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

3.15 分发版:Sunrise 打出 0.1.15+sunrise.1.2.0 发布 PR(#9,在审)

10-22 16:48 建 candidate/sunrise-0.1.15-1.2.0-20261019 分支并开 PR,把 Sunrise 分发版更新到 0.1.15+sunrise.1.2.0,纳入 TileLang 0.1.15 及配套评审过的 TileOPs、TileKernels、TVM、tvm-ffi 变更;源导入与合并、分发适配分开提交,最终适配配置 TANG 后端的 Sunrise wheel 并修掉切片 GEMM 回归(两个累加器标志都对公开模板校验,保留 stride、origin 与 offset 断言)。自报验证:本机契约测试 55 通过、TANG 流水线 codegen 测试 54 通过、两个 lint job 与信任边界检查通过;但 S2 wheel 构建出包后,装后 import 冒烟检查在 tangGetDeviceCount 处失败。分支上三笔分别是 0.1.15 分发适配、CI 收窄到 TANG 后端与切片 GEMM 累加器标志测试。

来源:https://github.com/tile-ai/tilelang-sunrise/pull/9

3.16 采用方:TileKernels 与 FlashQLA 窗口内无推送(10-22)

deepseek-ai/TileKernels 最后推送停在 10-14,QwenLM/FlashQLA 同样停在 10-14,两家采用方在本窗口均无提交。

来源:https://github.com/deepseek-ai/TileKernels


四、社区、教程与活动

4.1 CANN 社区开设 ops-tilelang 算子仓(10-22)

Atlas-Ascend-Fork/ops-tilelang 于 10-20 建仓,定位是 CANN 社区面向昇腾 NPU 的 TileLang 高性能算子仓:统一管理 Kernel 实现、自动化测试用例、标准工作负载与性能基线,目标是形成可复用、可验证、可持续演进的算子资产。仓内已有 src、tests、scripts、docs 四层与 OAT.xml、SECURITY.md 等开放原子社区规范文件,10-22 推送一笔修复”合法上界种子与用例哈希组合后溢出导致 CPU 测试失败”;更早的提交里有”把 Ascend A5 引用替换为 Ascend 950”与”benchmark 失败或中断的会话仍覆写性能基线”。以单一 DSL 为唯一实现语言的社区级算子仓此前未见,这条线值得按周跟踪。

来源:https://github.com/Atlas-Ascend-Fork/ops-tilelang

4.2 tilelang/agent:把算子优化写成 agent 的 SKILL 工作流(10-22)

tilelang/agent(建仓 08-28,3 星)10-22 推送,README 本身是一段可执行的 prompt 模板:以 @tilelang-operator-dev/SKILL.md 注入领域知识,输入只有算子名、GPU/CPU 参考实现、shape 与”tilelang npu 交付、中文交付”约束;注释写明 agent 需在 A5 环境执行、加载后按 10 步闭环走,Skill 内含选型决策树、8-CSV 规则、L0 到 L4 策略与 11 条陷阱。把 4.1 的算子仓、3.15 的指令级调度与本仓合看,昇腾侧正在形成”人写规则、agent 调优、仓里沉淀基线”的分工。

来源:https://github.com/tilelang/agent

4.3 社区新仓:Sail 包装、远程自动调优 agent 与教程仓(10-22)

flytiger-eco/tilelang-wrapper-for-sail 于 10-21 建仓、10-22 推送(Python,无说明);cc-c122/TileLang-Remote-Autotuning-Agent(Python,10-22 推送)与 easy-tilelang/tilelang-devkit、easy-tilelang/tilelang-tutorials 同日推送。数量多、单仓星标低,属教程与工具层的自然生长,尚无可以判断的技术方向。

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

4.4 文档与版本:文档站更新三次,VERSION 仍 0.1.15(10-22)

tilelang.github.io 三次 Update docs,TileFoundry.github.io 一笔 render TileFoundry v0.0.3,TileOPs.github.io 一笔。主仓 VERSION 停在 0.1.15,最近发布仍是 09-30 的 v0.1.15;仓库 8742 星、915 fork、开放 issue 389 条。

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


五、趋势观察

一,主仓的缺陷类型换了层。 上一周主仓清的是”形状与几何校验”——cast 舍入注解被静默丢弃、无 arch 的 CUDA target 静默退到 sm_50、稀疏元数据形状不校验。这一窗口九笔里有四笔落在分析器与源码位置这类编译期状态上:analyzer 在改写后被复用、切分后变量重绑定、eager 前端丢失 Python 位置。这类缺陷的共性是”生成的 kernel 可能仍然正确,但编译器对你说的和它实际做的不一致”,而 TileLang 正在接入 agent 自动写 kernel(4.2),编译期状态一致性会直接决定 agent 迭代闭环能不能收敛。

二,昇腾侧的进展从算子数量转到底层数据通路上。 #3530 一笔 18 文件,做的是 NZ 压缩布局在部分 UB→L1 拷贝中的行距表达;#3581 与 #3523 分别修编译期谓词爆炸与 SIMT 向量局部变量的分配语义。三笔都不新增算子,都是让既有路径在更多形状上可编译、可正确执行。配合 CANN 社区另起 ops-tilelang 算子仓(4.1),可以读出分工:算子资产由社区仓承载,主仓专注后端语义与 codegen。

三,TileOPs 两天连续 20 笔,工作重心在 SM90 特化与 SM89 补位。 窗口内新增三个 SM90 warp 特化 kernel(dense decode、因果滑窗 band、FP8 varlen 重叠),同时把 SM89 的 SSD chunk scan、paged DSA、NSA top-k 三条小规模路径一并推进。值得注意的是两笔 NaN 修复(#2708、#2699)都出现在 warp 特化路径上:共享内存规划器的活跃期判断和 fast math 下的 FMA 融合,是这类手工编排 kernel 最容易出错的两种地方,而这恰好是 TileLang 承诺要替用户管住的部分。

四,一个反向指标是基准层。 #2692 把参考实现改到 float32 计算,起因是 MTT S5000 上 GroupNorm 与 InstanceNorm 失败——问题出在参考实现跟着输入 dtype 走,而不是 kernel 本身。同日 TileFoundry #256 让”设备拒绝启动”从静默变成抛错。两笔都在同一个方向:把”看起来通过了”变成”确实通过了”,对依赖这些工具做芯片适配的团队,这类修复的实际价值高于几条性能数字。

五,分发与产业两条线在加速。 Sunrise 分发版把 0.1.15 与 TANG 后端的 wheel 打包成发布 PR,说明 TileLang 已经出现不止一个下游分发形态;产业侧 SCMP 转载 SemiAnalysis 的判断,把”六仓中只有 TileLang 可用”推到主流媒体。技术节奏上这仍是平静的一天(主仓无发布、无 breaking change),但生态层的组织动作明显多于代码层。


附录:素材与核查说明

信源核查表

信源 核查结果
tile-ai 组织仓库 29 仓,窗口内推送 10 个;合入 35 笔,在审更新 71 笔
tilelang 主仓 9 笔合入,1 笔单行修复,1 笔 18 文件昇腾改动;VERSION 0.1.15,无新发布
TileOPs 20 笔合入,全部为 Attention / MoE / Elementwise / Reduction 的 kernel 与测试
TileFoundry 3 笔合入,另 4 笔在审;0.0.3 于 10-21 发布的说明已在上一期收录
tilelang-ascend 2 笔合入(1 新增示例、1 回退示例),11 笔在审
tilelang-hygon main 无提交;合并分支 gcy_merge_main 一笔开关清理
tilelang-sunrise 新建 candidate 分支并开 0.1.15+sunrise.1.2.0 发布 PR,未合入
tilelang-metax / musa / mlir-ascend 窗口内无推送,分别停在 10-16 / 09-30 / 10-19
TileRT / tilescale / DeepStack 窗口内无推送
采用方 TileKernels / FlashQLA 均停在 10-14,无推送
Google News(中英文多组,走代理) 命中 3 条,2 条为同一份 SemiAnalysis 报告的媒体转载,1 条为 DeepSeek 融资与主题无关已剔除
Hacker News 无 TileLang 相关讨论
arXiv 以 TileLang 为标题或摘要关键词检索,窗口内 0 条
社区新仓 5 个新仓或活跃仓,均已核对建仓日期与推送时间
会员或厂商官网 本主题不适用

完整信源清单