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

信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileOPs-nightly / TileFoundry / tilelang-hygon / tilelang-metax / tilelang-ascend / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 文档站与夜间基准仓)+ 合入与在审真值检索(合入 20 笔、在审更新 45 笔)+ 新增 issue 全量核对(7 条)+ 提交级 diff 核查 + 发布页与 PyPI 制品核查 + 外部后端仓核查(Hexagon / CANN 算子仓 / PPU 包装器 / 模型片段仓)+ Google News 中英文多组查询(走代理)+ Hacker News + arXiv。


本期索引

  • 今日重点:TileOPs 0.1.0 正式发布,206 个算子、2265 行基准里 2262 行不慢于最快实现;24 小时合入 20 笔、在审更新 45 笔、新开 issue 7 条;海光 scale 缓冲双 bank 合入 1389 行;第三方 Hexagon 后端单日 15 笔
  • 一、核心项目进展
    • 1.1 主仓:ROCm 的 exp/sin/cos 不再默认降到快速内建,sin/cos 误差从 7.8e-04 降到 6.6e-08(#3621)(10-27)
    • 1.2 主仓:自动目标改读当前 HIP 设备,多卡机上不再固定取 0 号卡架构(#3594)(10-27)
    • 1.3 主仓:TMA gather/scatter 的常量行索引转为 int32,字面量写法不再编译崩(#3365)(10-27)
    • 1.4 主仓:NVRTC 与 CuTeDSL 按首个输入张量所在设备取当前流(#3359)(10-27)
    • 1.5 主仓:自动调优缓存键纳入 compile_flags,同一内核换编译选项不再复用旧二进制(#3338)(10-27)
    • 1.6 主仓:ThreadSync 把 address_of 记为写,不透明调用后的栅栏补上(#3346)(10-27)
    • 1.7 主仓:运行时行程数的流水线 epilogue 直接展开,分页解码恢复可编译(#3356)(10-27)
    • 1.8 主仓:切片填充与负索引的静态边界两修,越界写与静默负地址关闭(#3352 / #3347)(10-27)
    • 1.9 主仓:T.ws(-1) 区域在所有线程上运行,落单写法不再丢存(#3354)(10-27)
    • 1.10 主仓:profiler 未指定设备时用当前设备,缓存冲刷不再打到 0 号卡(#3357)(10-27)
    • 1.11 新增缺陷:内核缓存命中破坏六个假定冷编译的测试;CuTeDSL 的 CUDA graph 捕获为空图(#3694 / #3693)(10-27)
  • 二、多后端适配(昇腾 / 海光 / ROCm / Hexagon)
    • 2.1 海光:HCU scale 缓冲按双 bank 分区,blockscaled GEMM 在手动 warp 特化下共用一套行映射(#29)(10-27)
    • 2.2 昇腾主仓:PerCoreTask 候选体允许内部 if/else,条件 L1 填充不再被整体拒绝(#3686)(10-27)
    • 2.3 昇腾主仓:CI 把 AscendC 与 PTO 拆成两个矩阵作业,可分别重跑(#3687)(10-27)
    • 2.4 昇腾仓在审:Cython 包装器惰性构造分析器,直指昨日记录的每次调用重建问题(#1948)(10-27)
    • 2.5 昇腾仓在审:GLM-5.3 的 kpool 压缩算子、PTO 上的张量移位量(#1910 / #1949)(10-27)
    • 2.6 昇腾主仓在审:PTO 的 T.Parallel 向量化、AutoSimVF 通道、批量三维 ND2NZ、写计数多缓冲、Bisheng 构建身份进缓存键(#3512 / #3320 / #3689 / #3629 / #3691)(10-27)
    • 2.7 ROCm 在审:pip 版 ROCm 10 的架构探测与镜像、FP4 转换与 E8M0 类型(#3549 / #3665)(10-27)
    • 2.8 昇腾仓日测:2903 项全通过,连续两日无失败(#1950)(10-28)
    • 2.9 沐曦:组织视图有推送时间戳,六个分支均无窗口内新提交(10-27)
    • 2.10 第三方:Hexagon DSP 后端单日 15 笔,HVX 向量化、VTCM 竞技场、Crouton 布局与 DMA 内建齐备(10-27)
  • 三、生态与采用方(TileOPs / TileFoundry)
    • 3.1 TileOPs 0.1.0 发布:206 个算子的首个正式版,公开 API 就此定形(10-27)
    • 3.2 TileOPs 0.1.0 基准:H200 上 2265 行里 2262 行不慢于最快实现,几何平均 1.37 倍(10-27)
    • 3.3 TileOPs 0.1.0 的已知缺口与下一步:融合算子、整模型、计算通信融合(10-27)
    • 3.4 TileOPs 合入四条:where 与 masked_fill 可调优、D128 因果对角按 64 键打分、稀疏长 GDN 预填充拆分、变更日志按发布日定日期(#2782 / #2783 / #2779 / #2781)(10-27)
    • 3.5 TileFoundry 合入两条:整窗口按操作数本身下发、动态共享内存按内核实际用量定尺寸(#273 / #271)(10-27)
    • 3.6 外部生态:CANN 社区的 TileLang 算子仓、PPU 版零改动包装器、DeepSeek V4.1 Flash 片段示例(10-27)
  • 四、社区、教程与活动
    • 4.1 文档站单日六笔自动更新(10-27)
    • 4.2 昇腾仓的文档事实核查四条,单条评论最多 46 轮(#1628 / #1655 / #1656 / #1659)(10-27)
    • 4.3 在审:把 lower trace 抽成独立的开发工具包(#3583)(10-27)
    • 4.4 新闻侧:Google News 与 HN 24 小时内无本次窗口的新报道(10-27)
  • 五、趋势观察
  • 附、素材与核查说明

今日重点

TileOPs 0.1.0 在今天窗口内落地,结束了连续多日的「卡在最后一格」状态。PyPI 上的 0.1.0 制品上传于 10-27 07:45(北京时间),比上一期日报的截点晚 45 分钟;GitHub 发布页在 13:15 挂出,仓库首次出现正式 tag。发布说明的硬指标是三组:206 个算子全部有 manifest 条目(签名、测试与基准负载、roofline 公式,CI 里逐条校验);H200 上 2265 行夜间基准中 2262 行不慢于最快对照实现,几何平均 1.37 倍,对照方包括 torch.compile、FlashAttention 3、FlashInfer、vLLM、DeepGEMM 与 FlagGems;206 个算子里有 195 笔带 foundry 标签的合入是由 TileFoundry 生成内核并通过同一套 CI 进库的。发布同时把公开 API 冻结,之后 0.x 只增不改,并新增 API 兼容性检查脚本,任何按上一版写法调用的 PR 会直接失败。

24 小时内 tile-ai 组织合入 20 笔(tilelang 13、TileOPs 4、TileFoundry 2、tilelang-hygon 1),在审更新 45 笔,新开 issue 7 条(其中 2 条是昇腾仓的每日测试机器人)。合入重心分两条线:一条是浮点语义与执行上下文,ROCm 把七个默认数学函数从快速内建改回精确实现(MI300X 上 sin/cos 误差从 7.8e-04 降到 6.6e-08),自动目标改读当前设备而不是固定 0 号卡,NVRTC 与 CuTeDSL 改为按输入张量所在设备取流,自动调优的缓存键补进 compile_flags;另一条是静态边界与线程域,切片填充按区域末端而不是按长度校验、超范围负索引在编译期拒绝、T.ws(-1) 的区域改为在所有线程上运行。海光合入 12 文件 1389 行,把 scale 缓冲拆成独立双 bank 并与 LDS bank 分开。第三方侧 Hexagon DSP 后端单日 15 笔,把 HVX 向量化、VTCM 竞技场、Crouton 布局与 DMA 内建一路铺到 Crouton 上的映射与归约。


一、核心项目进展

1.1 主仓:ROCm 的 exp/sin/cos 不再默认降到快速内建,sin/cos 误差从 7.8e-04 降到 6.6e-08(#3621)

合于 10-27 22:39,5 文件 +107/-17,修 #3618。在 HIP 上,T.exp、T.exp10、T.log、T.log2、T.log10、T.sin、T.cos 七个默认函数无论 tl.enable_fast_math 如何设置都会降到 __expf 一类快速内建,而 CUDA 侧按 #875 的约定走精确实现;反过来,显式的 T.__exp 等快速算子在 HIP 上根本不下沉,在 tilelang.rocm.language 里直接抛 AttributeError。本笔让 HIP 与 CUDA 采用同一策略:七个默认算子改派 HIPMath,八个带下划线的快速算子新增最低下沉规则指向 HIPFastMath,并统一挪进 FAST_MATH_INTRINSICS 列表由 ROCm 语言层导出。精度对比取 #3618 的复现:MI300X 上 exp 与 exp10 看相对误差、sin 与 cos 看绝对误差(x 到 1e4),四列依次为 T.exp、T.exp10、T.sin、T.cos,主干为 3.7e-06 / 4.1e-06 / 7.8e-04 / 7.8e-04,本笔为 8.3e-08 / 8.1e-08 / 6.6e-08 / 6.4e-08。

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

1.2 主仓:自动目标改读当前 HIP 设备,多卡机上不再固定取 0 号卡架构(#3594)

合于 10-27 22:37,2 文件 +24/-1。determine_target(“auto”) 在 ROCm 上落到注册的 HIP 探测器,而该探测器无条件读 torch.cuda.get_device_properties(0),于是当前 HIP 设备不是 0 号卡时,自动探测给出的是 0 号卡的架构;CUDA 侧早就在用 torch.cuda.current_device(),ROCm 没有跟上。修法是先取当前设备索引再查属性,并用 normalize_rocm_arch 去掉特性后缀。新增测试把 current_device() 打桩返回 1,同时记录所有被探测过的索引,断言探测到 gfx942 且只碰过 1 号卡;修复前该断言得到 probed == [0]。作者注明本机无 ROCm 设备,测试只用打桩 torch 验证探测助手,未跑真实 HIP 运行时;与本日另一笔 #3549 改动的是不同函数(#3549 在 tilelang/contrib/rocm.py,本笔在 tilelang/rocm/target.py)。

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

1.3 主仓:TMA gather/scatter 的常量行索引转为 int32,字面量写法不再编译崩(#3365)

合于 10-27 09:15,2 文件 +44/-2,修 #2601。tma_gather4 与 tma_scatter4 两个前端助手在收到 Python 整数形式的行索引时直接透传,落到 TIR 里没有确定的整型宽度,于是在 Blackwell 上编译崩溃;其他行表达式保持原样。本笔只把 Python 整数转成 int32 的 TIR 常量,并补一个受 Blackwell 门控的编译回归,用字面量行索引分别编译 gather 与 scatter,再检查生成的 CUDA 源码里确实出现对应指令。作者说明该测试只验证能编译,不验证运行时数值正确;且提交者本机的 Windows Python 3.13 环境里 Home Assistant 的 pytest 插件依赖 fcntl,pytest 收集阶段就失败,本笔的 TileLang 测试在本地并未真正跑起来。

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

1.4 主仓:NVRTC 与 CuTeDSL 按首个输入张量所在设备取当前流(#3359)

合于 10-27 08:11,2 文件 +5/-4,接 #3301 的 F149 与 #3408。两个适配器此前用字符串前缀判断目标类型,且默认取调用方当前设备的当前流,于是调用方当前 GPU 与张量所在 GPU 不一致时选错流。本笔改用目标种类判断,并按首个输入张量的设备查当前流,显式传入的流句柄仍然最高优先,无张量的 NVRTC 调用维持原语义。这是继昨日 out_idx 输出分配流(#3649)之后,流语义清理的第二笔。作者划清边界:本笔只修流选择,跨设备的默认流启动上下文仍是未解项,这类调用方必须自己在启动前把张量所在设备设为当前设备,补丁里没有逐调用的设备守卫;验证时用真实 Torch 流与 GPU 启动覆盖了同设备与跨设备的非默认流,但未建立跨设备默认流上下文的正确性。

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

1.5 主仓:自动调优缓存键纳入 compile_flags,同一内核换编译选项不再复用旧二进制(#3338)

合于 10-27 16:57,3 文件 +28/-1,接 #3301 的 F131。自动调优器把编译结果持久化缓存,但缓存标识里没有 compile_flags,于是同一个内核换一组设备编译器选项后仍会命中旧二进制,测出来的性能与真实编译结果对不上。本笔让 CompileArgs 真正存下并按原有顺序参与哈希,并把 flag 从 JIT 实现经 AutoTuner.set_compile_args() 一路透传到候选编译、逐配置的 pass 合并、分组设备编译与缓存内核加载。验证覆盖三条既有缓存键测试,以及用原生运行时对 CUDA 做的定点对比:列表与字符串两种 flag 形式都能在逐配置覆盖中存活;作者指出原先的合并逻辑会把编译器宏丢掉,正是这次要堵的口子。

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

1.6 主仓:ThreadSync 把 address_of 记为写,不透明调用后的栅栏补上(#3346)

合于 10-27 05:49,2 文件 +62/-1,接 #3301 的 F305,延续昨日的栅栏语义主线。address_of 本身不携带访问信息,它只把指针交给下游,但访问分类器把它的类型硬编码成读,于是一块共享缓冲经不透明调用写入后不会进 writes 集合,后续的全块读找不到冲突写、也不再规划栅栏,线程就这样抢过写入读到写前内容。实测的 CUDA 输出很直白:T.address_of(B[0]) 形式在 extern 之后没有栅栏,而 T.access_ptr(B, “w”) 形式有;后者显式声明了访问类型,所以这个洞一直没被发现。loop_unswitching.cc 早已按写保守处理 address_of,ThreadSync 是唯一的例外。本笔把 address_of 记为写,方向是保守的:指针用途未知就可能被写,而栅栏规划只需在安全侧错一次;新增两个用例,断言该内核下沉后 extern 之后跟着 __syncthreads()。

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

1.7 主仓:运行时行程数的流水线 epilogue 直接展开,分页解码恢复可编译(#3356)

合于 10-27 07:02,1 文件 +5/-1。T.Pipelined 循环的行程数是运行期值时,只要 epilogue 展开两步以上且带异步拷贝,CUDA codegen 就报 TypeError: Downcast from tirx.Sub to ir.IntImm failed。根因在 EmitImpl:prologue 与 epilogue 的 extent 是常量时才展开,以保证每次迭代的等待计数是常量;而 epilogue 的 extent 是 end - start,即 (n + max_stage) - n,n 为常量时在构造期折叠,n 为运行期时保持为 Sub,as_const_int 返回空,于是 epilogue 被当成 T.unroll 循环发射,其等待写成 T.ptx_wait_group(1 - k) 这样读循环变量的形式,而 ptx_wait_group 只收立即数。修法是判断前先看化简后的 extent,常量边界的情形化简回自身,与原来一致。该缺陷由下游 TileOPs#1852 报出并二分到 afcebed1 之后,破坏对象是分页解码——其行程数来自 KV 长度。作者在 sm_89、CUDA 13.0 上双版编译对比,新增测试在主干失败、在本笔通过。

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

1.8 主仓:切片填充与负索引的静态边界两修,越界写与静默负地址关闭(#3352 / #3347)

合于 10-27 05:37 与 05:14,合计 4 文件 +100/-3。前者针对 Fill 构造器的静态边界检查:原来逐维只校验 min >= 0 与 extent <= shape,而 extent <= shape 只对从缓冲起点开始的区域成立;一个起点靠里的切片两项都通过,却被下沉成越界写——T.fill(As[6:12], 7) 在长度 8 的缓冲上编译通过,发射出 As[threadIdx.x + 6] = 7,写到 6 到 11。修法是把区域末端也对照目标形状校验,且新条件要求起点与长度都是静态,锚定在零的区域与符号起点区域都不受影响;T.clear 共用同一节点,缺陷相同也一并关闭。后者针对 LegalizeNegativeIndex:原来把负索引加一次缓冲长度就了事,没有范围检查,长度 8 的缓冲上 A[-9] 得到 -1、A[-10] 得到 -2,跳出分配起点且无任何诊断,而 NumPy 与 PyTorch 对同样的索引抛 IndexError。修法是仅在分析器能证明回绕后仍为负时拒绝,并报出索引与长度,所有在范围内的常量与全部动态索引维持原行为。

来源:https://github.com/tile-ai/tilelang/pull/3352、https://github.com/tile-ai/tilelang/pull/3347

1.9 主仓:T.ws(-1) 区域在所有线程上运行,落单写法不再丢存(#3354)

合于 10-27 04:52,2 文件 +40/-3,接 #2978。WarpSpecialize 的文档把 -1 定义为「跳过 threadIdx.x 绑定」,也就是该区域不受 threadIdx 约束;实现里却把 -1 当普通 warp 组索引,构造出 [-warp_group_size, 0) 这个对任何线程都不成立的区间,于是区域在零个线程上运行,区间内所有存全部被丢弃且无诊断。最小复现很干净:256 线程里先全写 0,再在 T.ws(-1) 里写 5,结果 256 个线程写出的仍是 0;同样的存不加 ws 包装则 256 个线程都写 5。修法是当 id 里含 -1 时不发射任何范围约束,真实组索引仍按自身线程约束,并在文档串里补上 -1 这一例——原文档的示例恰好缺它。作者核查全仓未出现 ws(-1) 用法,因此没有既有代码依赖那个空区域。

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

1.10 主仓:profiler 未指定设备时用当前设备,缓存冲刷不再打到 0 号卡(#3357)

合于 10-27 07:26,1 文件 +1/-1。do_bench 在未显式指定设备而当前 CUDA 设备非零时,_cache_device(None) 落到模块级默认的 cuda:0,于是为冲刷缓存而做的分配可能落在与实测工作不同的 GPU 上,测出来的数字与真实负载不在同一张卡。修法是默认串改成不带索引的 “cuda”,由 PyTorch 在分配缓冲时解析当前设备;显式指定设备与 NPU、MPS 的默认值保持不变。上游 CI(run 37915443170)在 Quick Lint、CUDA、ROCm、Metal、Ascend 与 CuTeDSL 示例上通过;作者注明后续的 PR 触发式运行需维护者批准,尚未执行测试,本机也没有 CUDA 硬件或计时基准。

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

1.11 新增缺陷:内核缓存命中破坏六个假定冷编译的测试;CuTeDSL 的 CUDA graph 捕获为空图(#3694 / #3693)

两条 issue 均由 lucifer1004 在 10-27 开出。前者指出有六个测试在冷缓存下通过、用同一缓存目录再跑一次即失败,原因是它们假定编译一定会发生:例如 test_tilelang_kernel_gemm_sm75.py 里的 sm75 边界用例在四个参数下报 AttributeError: ‘NoneType’ object …;后者更硬:从 CuTeDSL 后端发出的内核启动没有被记进 CUDA graph,torch.cuda.CUDAGraph 捕获得到空图,replay() 什么都不跑,而同一内核用 cuda 后端编译则捕获回放正常,报告附了最小复现。两条都是「缓存与执行上下文」这一族缺陷的延伸,与 1.4、1.5 两笔合入指向同一片区域。

来源:https://github.com/tile-ai/tilelang/issues/3694、https://github.com/tile-ai/tilelang/issues/3693


二、多后端适配(昇腾 / 海光 / ROCm / Hexagon)

2.1 海光:HCU scale 缓冲按双 bank 分区,blockscaled GEMM 在手动 warp 特化下共用一套行映射(#29)

合于 10-27 08:26,12 文件 +1389/-314,是本窗口软件侧最大的一笔。改动把 scale 缓冲的行规划成独立的双 bank 分区,与 LDS bank 分开,并把同一套行映射同时交给分配、scale 拷贝与 MMAC 下沉三处使用,避免三处各自推一遍布局。细节包括:保留物理区间化的生命周期与 m0 对齐分页;压缩 scale 拷贝的下发条数;把空闲 wave 的工作摊到两个 scale bank 上;该布局用于手动 warp 特化下的 blockscaled GEMM。这一笔接上昨日合入的原生 B16 存储与动态 WASP GEMM(#28),海光在两天里完成了三笔主干合入。

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

2.2 昇腾主仓:PerCoreTask 候选体允许内部 if/else,条件 L1 填充不再被整体拒绝(#3686)

合于 10-27 07:24,2 文件 +85/-24。T.PerCoreTask 里的 T.Task 只要体内含 if/else 就被拒绝,哪怕分支只控制内部工作——例如一次拷贝之后按条件补 L1 填充。复现于 d85a70a9,物化阶段报「选择候选的控制流必须位于 task 作用域之外」。本笔删掉这条一刀切的拒绝,改为对 task 体做统一校验:同步栅栏包住整个候选,同时每个分支仍逐项检查管线与核的兼容性、嵌套 task、以及不受支持的跨核同步。回归测试接受内部 if 与 if/else、保留候选体、并保留原有的拒绝用例,另检查完成标志在外层守卫下跟随整个候选。task、调度单元与延迟三套共 113 例全通过,带条件 L1 填充的内核也能完成完整昇腾源码下沉;设备执行未测。

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

2.3 昇腾主仓:CI 把 AscendC 与 PTO 拆成两个矩阵作业,可分别重跑(#3687)

合于 10-27 07:05,7 文件 +56/-27。此前 AscendC 与 PTO 是同一个 Test for Ascend 作业里的两个步骤,任一后端专属失败都会让整条检查变红,且重跑必然连带另一套示例。本笔把示例与单元测试拆成两个矩阵作业:Test for AscendC 跑 examples/ascend(target 为 ascend,4 并发)与 testing/ascend -m “not pto”(32 并发),Test for PTO 跑 examples/ascend(target 为 pto)与 testing/ascend -m pto(32 并发)。划分靠既有的 pto 标记,并给只跑 PTO 的测试与双后端测试里的 PTO 参数补上标记;两个标记表达式互斥且合起来覆盖全套,不重复执行。fail-fast 置为 false,一个作业失败不取消另一个;两边都保留「示例失败但构建成功时仍跑单测」的行为,各自发布构建与测试摘要。共用测试(含同一前端 IR 分别过两个代码生成器的契约检查)只在 AscendC 一侧跑一次,约定写进 testing/ascend/README.md。

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

2.4 昇腾仓在审:Cython 包装器惰性构造分析器,直指昨日记录的每次调用重建问题(#1948)

更新于 10-27 07:42,作者 zwh1025。昨日日报记录的新缺陷 #1945 指出,昇腾仓的 Cython 包装器每次调用都新建一个 TVM 分析器,实测约 0.28 毫秒的固定开销;本笔以「惰性构造 TVM 分析器」为题,正对这个开销。这类小开销在大模型推理的逐层调用里会累积成可观测的延迟,属于昇腾侧性能清扫的低垂果实。目前仍在审。

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

2.5 昇腾仓在审:GLM-5.3 的 kpool 压缩算子、PTO 上的张量移位量(#1910 / #1949)

两条均在 10-27 更新。#1910 给昇腾 NPU 加 GLM-5.3 的 kpool 压缩算子示例——这是把具体模型的结构件搬进示例仓的做法,与 TileOPs 侧「算子按模型需求长」的思路一致;#1949 让 T.tile.bitwise_lshift 在 PTO 上支持张量形式的移位量,此前只接受常量。PTO 是昇腾侧相对独立的一条代码生成路径,移位量张量化意味着该路径在掩码与位打包类算子上向通用表达靠了一步。

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

2.6 昇腾主仓在审:PTO 的 T.Parallel 向量化、AutoSimVF 通道、批量三维 ND2NZ、写计数多缓冲、Bisheng 构建身份进缓存键(#3512 / #3320 / #3689 / #3629 / #3691)

五条在审构成昇腾侧本窗口的排队队列,均于 10-27 有更新。#3512 是 PTO 上把 T.Parallel 向量化的流水线框架与逐元素支持,属结构性改动;#3320 新增 AutoSimtVF 通道并注册进昇腾流水线,方向上是在 SIMT 与向量之间自动择路;#3689 支持批量三维 GM 到 L1 的 ND2NZ 拷贝,直接对应批量注意力与分块 GEMM 的搬运需求;#3629 增加写计数式多缓冲并重排 DeepGEMM 的加载;#3691 把 Bisheng 编译器构建身份纳入二进制缓存键,与 1.5 的 compile_flags 进缓存键是同一类思路——昇腾侧同样存在换编译器后误命中旧二进制的问题。

来源:https://github.com/tile-ai/tilelang/pull/3512、https://github.com/tile-ai/tilelang/pull/3320、https://github.com/tile-ai/tilelang/pull/3689、https://github.com/tile-ai/tilelang/pull/3629、https://github.com/tile-ai/tilelang/pull/3691

2.7 ROCm 在审:pip 版 ROCm 10 的架构探测与镜像、FP4 转换与 E8M0 类型(#3549 / #3665)

#3549 支持 pip 安装的 ROCm 10 SDK,改动含架构探测与 Dockerfile;这与 1.2 的 auto 目标修正相关但落在不同函数——本笔动的是基于 /opt/rocm 的 get_rocm_arch 探测,1.2 动的是 tilelang/rocm/target.py 里的探测助手。#3665 修 HIP 模板里的浮点到 FP4 转换,并补 E8M0 类型支持,指向分块缩放与低精度量化在 ROCm 上的可用性。两条均在审。

来源:https://github.com/tile-ai/tilelang/pull/3549、https://github.com/tile-ai/tilelang/pull/3665

2.8 昇腾仓日测:2903 项全通过,连续两日无失败(#1950)

昇腾仓的每日测试机器人 10-28 06:49 的报告显示总计 2903 项、通过 2903、通过率 100%,工作流运行号 40590261193;前一日的报告(#1947)同样是 2903 项全通过。对一条仍在快速演进的 NPU 后端来说,两日无失败的日测说明主干可用性稳定,这也为 2.3 的 CI 拆分提供了基线——拆分后两个作业各自仍应维持这个数字。

来源:https://github.com/tile-ai/tilelang-ascend/issues/1950

2.9 沐曦:组织视图有推送时间戳,六个分支均无窗口内新提交(10-27)

tilelang-metax 仓的 pushed_at 显示 10-26 07:08,落在本窗口起点之前;逐分支核查六个分支,窗口内没有新提交,按无实质变更处理,不列条目级内容。

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

2.10 第三方:Hexagon DSP 后端单日 15 笔,HVX 向量化、VTCM 竞技场、Crouton 布局与 DMA 内建齐备(10-27)

xwhzz/tilelang-hexagon 的 hexagon-backend 分支在 10-27 一天落了 15 笔提交,是本窗口后端侧最密集的外部动作。当日主线可以分成三段:底层从 CodeGenTileLangCBase 重建代码生成器,把 DMA 做成带提交组的 IR 内建,把启动网格与工作线程池下沉进 TIR,并在 TIR 里规划 VTCM 竞技场;中层铺 HVX 的单条向量化通路(算术映射、向量数学加一条逐元素路径)、用 HVX 转置树做多行归约、软件流水化 HVX 映射,并支持片段与 T.reduce_*;上层引入 Crouton 布局与区域拷贝、让 HVX 映射与归约直接在 Crouton 瓦片上跑、每个会话对 DSP 频率投票,最后补构建缓存与并行设备测试。提交里带一个破坏性标记(feat(hexagon)!),说明 Crouton 布局与区域拷贝改动了既有接口。Hexagon 是手机与嵌入式侧的 DSP,这条线若走通,TileLang 的硬件覆盖面就从 GPU、NPU 延伸到 DSP。

来源:https://github.com/xwhzz/tilelang-hexagon


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

3.1 TileOPs 0.1.0 发布:206 个算子的首个正式版,公开 API 就此定形(10-27)

10-27 13:15 挂出 v0.1.0,是仓库首个 GitHub release;PyPI 上的 0.1.0 制品上传更早,为 10-27 07:45(北京时间,即 10-26 23:45 UTC),比上一期日报的截点晚 45 分钟——这正是昨日日报写「首发仍在最后一格」之后立刻发生的变化。发布说明定位很明确:一个为 agent 时代重做的算子库,每个算子先写 YAML 规格(声明签名、测试与基准用的负载、roofline 公式),由编码 agent 开发与调优内核,再由规格派生出的测试与基准决定谁能进库;195 笔带 foundry 标签的合入就是这样进来的。本版固定公开 API:算子按 tileops.. 导入,docs/stability.md 列出公开模块,scripts/ci/check_api_compat.py 让任何破坏上一版写法的 PR 直接失败;后续 0.x 只增不改。另有几项能力一并进版:torch.compile(fullgraph=True) 对所有 manifest 条目含调用期张量输入且无组合的算子可用;H200 上验证了图测试组内算子的 CUDA Graph 捕获;op.request_tune() 后自动调优,编译产物与调优结果缓存在 TILELANG_CACHE_DIR 下跨进程复用;tileops.backend 允许后端分发版通过 register_kernel_builder 把自己设备上的整算子接进来;tileops.trace 记录内核内时间线。发布致谢指出 @FeathBow 一人贡献 98 笔合入。

来源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0、https://pypi.org/project/tileops/

3.2 TileOPs 0.1.0 基准:H200 上 2265 行里 2262 行不慢于最快实现,几何平均 1.37 倍(10-27)

夜间基准把每个 TileOPs 内核与同一算子的其他所有实现放在同一张 GPU 上比,单行比值取最快对照方耗时除以 TileOPs 耗时,大于 1 表示更快。全量 2265 行、16 个算子族,几何平均 1.37 倍,2262 行达到 0.99 以上。按族看差异明显:sequence_modeling 4.00 倍、sampling 2.95 倍、convolution 1.86 倍、mamba 1.61 倍、linear_attention 1.54 倍、pool 1.51 倍、quantization 1.45 倍、attention 1.42 倍;体量最大的 elementwise 527 行只有 1.09 倍,gemm 235 行 1.28 倍。对照方包括 torch、torch.compile、FlashAttention 2 与 3、FlashInfer、FlashMLA、vLLM、DeepGEMM、flash-linear-attention、FlagGems、QuACK、Mamba、Marlin 与 cuBLASLt。落后 0.99 的三行分别是 16M float32 元素上的 Where 与 ClampScalar、以及 DeepSeek-V3 预填充输出投影上的 FP8 GEMM,差距 1.1% 到 1.7%。这个分布说明收益主要来自长尾算子族而非主流 GEMM 与逐元素——后者本来就是各路实现卷得最狠的地方。

来源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.3 TileOPs 0.1.0 的已知缺口与下一步:融合算子、整模型、计算通信融合(10-27)

发布说明列了四条已知问题:FusedAddLayerNormFwdOp 在远离零的行上丢精度、在接近 1e36 的行上溢出(#2773,昨日日报已记录该缺陷);同一主机上若有不同架构的 GPU,为一个 GPU 编译的内核可能跑到另一个上(#2302);Bwd 类算子需直接调用,不接入 torch.autograd;单个算子实例同时只支持一个线程与一条 CUDA 流。验证范围只覆盖 H200、Python 3.12 与 3.13;部分内核另有 SM80、SM86、SM89 路径,这些路径在对应设备上跑但不属于发布验证范围。路线图三步:先做跨算子边界的复杂融合(注意力连同前后投影与归一化、整层 MoE、GEMM epilogue 内的量化),再做整模型的预填充与解码并与服务引擎端到端对比,最后做计算与通信融合(张量并行 GEMM 配 all-reduce、专家并行 MoE 配 all-to-all)。这三步若落地,TileOPs 的对手就从单算子库变成推理框架。

来源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.4 TileOPs 合入四条:where 与 masked_fill 可调优、D128 因果对角按 64 键打分、稀疏长 GDN 预填充拆分、变更日志按发布日定日期(#2782 / #2783 / #2779 / #2781)

四笔在 10-27 合入,恰是发布前的收尾。#2782 修 request_tune() 在 WhereFwdOp、MaskedFillTensorFwdOp、MaskedFillScalarFwdOp 上抛 ValueError 的问题:自动调优拒绝随机生成的整型输入,而这三个内核把 bool 条件或掩码当 uint8 传;随机选择器只改变每个元素取哪个操作数、不改变候选的工作量,故声明随机整型输入安全;新增三个测试节点,主干上三例全失败。#2783 优化 D128 稠密程序的因果对角:当因果调用只有一段 key 且两个长度都在 128 行边界上时,第一个消费 warpgroup 的对角 key 瓦片对全部行都屏蔽了高 64 键,于是这一步改按 m64n64 的 QK 配 64 键 softmax 跑;该选择是编译期标志(diag_aligned,进内核键),因为运行期在 64 与 128 宽 QK 之间分支会让 ptxas 在某些特化上串行化每条 WGMMA(C7511),改用标志后所有 manifest 里的 D128 预填充特化都不再出 C75xx。#2779 把 SM89 上稀疏长序列的 GDN 预填充拆成状态扫描加独立输出程序,适用于稠密、key 主序的 D128 输入且至少 16384 token、状态扫描 CTA 不超过 SM 数四分之一的情形;基准里 (1, 16384, 4, 128) bf16 的耗时从约 0.612 毫秒降到约 0.470 毫秒,与 flash-linear-attention 的 0.441 毫秒相比为 0.939 倍。#2781 把 CHANGELOG 里 0.1.0 的日期定为它进 PyPI 的那天(2026-10-27),并删掉贡献者文档里的发布一次配置与 Actions 点击路径。

来源:https://github.com/tile-ai/TileOPs/pull/2782、https://github.com/tile-ai/TileOPs/pull/2783、https://github.com/tile-ai/TileOPs/pull/2779、https://github.com/tile-ai/TileOPs/pull/2781

3.5 TileFoundry 合入两条:整窗口按操作数本身下发、动态共享内存按内核实际用量定尺寸(#273 / #271)

两笔均在 10-27 合入。#273 修 sm_89 GEMM 的下沉:schedule finalize 写出的内容编不过,因为下沉把整个累加器包进一个从零开始、跨度等于自身的窗口 T.ptr_of(acc[0:0+16, 0:0+8]),而 CUDA codegen 拒绝分片寄存器瓦片的指针切片,报「要求普通源布局」;MMA 又把寄存器片段读成 lhs_frag_dst[0:0+16, 0:0+32] 这样的裸 Slice,codegen 同样无处理器。修法在下沉侧:StoragePlan.operand 本来就把「从零开始且跨满操作数的窗口」当作操作数本身,但物化与指令下发仍构造 Slice;改为只判定一次——对 MLIR 的 TrivialSubViewOpFolder 做同样的替换。#271 修动态共享内存:合成的宿主入口给每个 TIR 内核都用 dynamic_smem=0 启动,经整型 TensorView 偏移访问共享内存的内核因此没有分到空间,超过 48 KiB 还需设置 cudaFuncAttributeMaxDynamicSharedMemorySize 而启动器从不设置;作者没有采纳 #267 里「按目标容量给满」的建议,改为按内核实际用量定尺寸,理由是静态共享内存与动态共享内存同计一个上限、且每块申请多少决定一个 SM 能放几块——sm_89 上要满 99 KiB 会让用 512 字节的内核也只放得下一块。CUTLASS 与 TileLang 都按同样方式定尺寸。

来源:https://github.com/tile-ai/TileFoundry/pull/273、https://github.com/tile-ai/TileFoundry/pull/271

3.6 外部生态:CANN 社区的 TileLang 算子仓、PPU 版零改动包装器、DeepSeek V4.1 Flash 片段示例(10-27)

三个外部仓在窗口内都有动作,合起来看是一条「昇腾侧算子资产」的成形过程。Atlas-Ascend-Fork/ops-tilelang 是 CANN 社区面向昇腾的 TileLang 算子仓,10-20 建仓、10-27 提交「自动发现增量测试并移除算子映射表」;README 明确只支持 Ascend 950 及以上架构、要求 CANN 9.3.0 以上并用社区 weekly 版,设计原则是接口与调度分离、按业务域聚合、正确性先于性能、benchmark 与测试分离且基线按测试文件与芯片集群分片保存。kevin3113/models-tilelang 用这个仓的七个算子串出 DeepSeek V4.1 Flash 的「Attention 输出加 FFN 前处理」片段(mhc_post_forward、带平方和 GEMM、归约加 RMSNorm、mix 拆分、sinkhorn、apply_mix、ffn_norm),HIDDEN=5120、MHC=4、默认 decode 形态 SEQ_LEN=1,支持 SuperKernel 与 npugraph_ex 的 ACL graph,B 从 1 到 256 实测 SuperKernel 收益 1.17% 到 2.75%。flytiger-eco/tilelang-wrapper-for-sail 解决另一个方向的迁移:把 CUDA 风格的 TileLang 代码零改动跑在 tilelang-sail(PPU 后端分发版)上,做法是纯运行时拦截 import 协议与 site 机制,不修改上游源码也不重新打包,无 PPU 对应物的 CUDA 特性显式报错而非静默降级。三者分别对应算子供给、模型落地、后端迁移三个环节。

来源:https://github.com/Atlas-Ascend-Fork/ops-tilelang、https://github.com/kevin3113/models-tilelang、https://github.com/flytiger-eco/tilelang-wrapper-for-sail


四、社区、教程与活动

4.1 文档站单日六笔自动更新(10-27)

tilelang.github.io 在窗口内有六次 Update docs 提交(05:20、06:22、07:19、09:28、22:42、22:43),均由机器人推送,内容来自主仓合入后的文档重生成。发布当天文档站更新频次与主干合入节奏同步,属正常联动,不单列内容。

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

4.2 昇腾仓的文档事实核查四条,单条评论最多 46 轮(#1628 / #1655 / #1656 / #1659)

昇腾仓在 10-27 有四条以「事实核查」为题的文档 PR 更新:T.tile.clear(#1628)、T.tile.bitwise_or(#1655)、T.tile.bitwise_not(#1656)、T.tile.bitwise_lshift(#1659),评论数分别为 10、46、18、20,其中三条带机器人标记。一轮核查能滚到 46 条评论,说明讨论的不只是措辞——这类逐 API 对行为的核对,实际是在为昇腾侧的语义约定留档。对一条新后端线来说,文档与实现对齐的成本往往被低估,这组 PR 是可见的投入。

来源:https://github.com/tile-ai/tilelang-ascend/pull/1655、https://github.com/tile-ai/tilelang-ascend/pull/1656

4.3 在审:把 lower trace 抽成独立的开发工具包(#3583)

10-27 更新的这条在审 PR 把 lower trace 从主仓抽到 tilelang-devkit。下沉过程的追踪信息对调试后端适配很有用,但放在主仓里会增加编译路径与依赖负担;抽成独立工具包后,主仓可以保持轻量,需要时再挂上。这类「把调试能力外置」的改动通常出现在项目开始有多个下游后端的时候——本窗口内 Hexagon、PPU、昇腾三条线同时在动,正是这种时点。

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

4.4 新闻侧:Google News 与 HN 24 小时内无本次窗口的新报道(10-27)

中英文四组查询(TileLang、TileOPs、TileLang 算子、tile-ai 智算 kernel)返回的条目全部落在 9 月底到 10 月中:一批是 DeepSeek 与华为在 9 月 29 至 30 日开源昇腾版 TileLang 的报道(SCMP、Reuters、Tom’s Hardware、雷峰网、InfoQ 等),最新一条是 10 月 22 日 SemiAnalysis 关于华为与 DeepSeek 共生关系的分析,均在窗口之外。TileOPs 精确查询零命中。Hacker News 窗口内只有一条与 GameMaker 重实现相关的条目,按噪音剔除。arXiv 上 TileLang 关键字最近一篇为 2026-07-24。今日的内容重心全在 GitHub 侧,这与发布日的时间点吻合——正式发布尚未进入媒体报道周期。

来源:https://news.google.com/


五、趋势观察

第一,TileOPs 0.1.0 把「算子库」这件事的成本结构换了一种算法。传统做法是算子数量随硬件与模型需求线性增长,人力是瓶颈;这里的做法是每个算子先写一份 YAML 规格,把签名、负载与 roofline 声明清楚,再由 agent 生成与调优内核,最后用规格派生的测试与基准做准入。206 个算子、195 笔 foundry 标签合入是这个流程的第一次全量交付。它能否持续,取决于两件事:规格能否覆盖复杂融合算子(路线图第一步正是这个),以及基准能否挡住「调优到基准上但实际不鲁棒」的内核——FusedAddLayerNormFwdOp 在远离零的行上丢精度这条已知问题,恰好是这类风险的一个样本。

第二,本窗口的合入几乎全是「静默错误」。ROCm 的数学函数默认走快速内建会让 sin/cos 差出三个数量级;address_of 被当成读会让不透明调用后的共享内存读到写前内容;T.ws(-1) 的区域在零个线程上运行,存全部丢弃且不报错;切片填充按长度而非末端校验,直接越界写;负索引回绕后仍为负,落到分配起点之前。这五条的共同点是编译通过、运行不报错、结果错,且都在最近两周内被集中发现。对一门要给大模型写内核的语言,这类缺陷的代价远高于编译期报错,它们指向同一个工程需求:把「不报错但行为不对」的写法在编译期拒掉。本窗口的边界校验两修与 ws(-1) 修正在做的就是这件事。

第三,后端覆盖正在从「厂商主导」转向「社区自发」。官方维护的昇腾、海光、沐曦、摩尔线程四条线之外,Hexagon DSP 后端在一天里落 15 笔并把线性代数原语铺到 Crouton 瓦片上,PPU 版出现零改动包装器,CANN 社区自建算子仓并用它串出 DeepSeek V4.1 Flash 的模型片段。这些动作都不在 tile-ai 组织的仓库里,却都依赖 TileLang 的语言层稳定性。这也是 0.1.0 冻结公开 API、并把 API 兼容性检查做成 CI 门禁的意义所在——下游越多,向后兼容的价值越高。

第四,缓存身份这件事被两个后端同时盯上。主仓把 compile_flags 纳入自动调优缓存键(#3338),昇腾把 Bisheng 构建身份纳入二进制缓存键(#3691)。同一个内核在不同编译选项或不同编译器下产生的二进制不同,缓存若只按源码与形状做键就会误命中,测出来的性能与真实编译结果对不上。两条独立路线在同一天做同一件事,说明这类问题在实际使用中已经足够痛。


附、素材与核查说明

信源核查表

信源 核查结果
GitHub tile-ai 组织(29 仓) 窗口内推送 7 仓:tilelang、TileOPs、TileOPs-nightly、TileOPs.github.io、tilelang-hygon、TileFoundry、tilelang.github.io
合入真值(is:merged) 20 笔,逐笔核对合入时间落在窗口内(tilelang 13、TileOPs 4、TileFoundry 2、tilelang-hygon 1)
在审真值(is:open 更新) 45 笔,按更新时间排序全量核对
新增 issue 7 条,逐条读正文,其中 2 条为昇腾仓每日测试机器人报告
发布页与 PyPI TileOPs v0.1.0 于 10-27 13:15 挂出,PyPI 制品上传 10-27 07:45;tilelang 最新仍为 0.1.15(09-30),tilefoundry 最新 0.0.4rc1(10-25)
tilelang-ascend 仓 窗口内无提交,但 PR 与每日测试 issue 活动密集
tilelang-metax 仓 组织视图有推送时间戳,六个分支均无窗口内提交,按无实质变更处理
tilelang-musa 仓 上次推送 09-30,窗口内无提交
采用方(TileKernels / FlashQLA) 均无窗口内推送(10-14)
外部后端仓 Hexagon 后端 15 笔提交;CANN 算子仓、PPU 包装器、模型片段仓均有窗口内动作,见 2.10 与 3.6
Google News 中英文多组查询 精确查询零命中,返回条目全部落在 9 月底至 10 月中,均在窗口外
Hacker News 窗口内一条 GameMaker 重实现条目,与 TileLang 无关,剔除
arXiv 关键字检索无窗口内新论文,最近一篇为 2026-07-24

完整信源清单