TileLang 每日动态报告(2026-10-29)
调研窗口:2026-10-28 07:00 ~ 2026-10-29 07:00(北京时间,24 小时)。
信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文档站)+ 合入与在审真值检索(合入 24 笔、在审更新 48 笔)+ 新增 issue 全量核对(6 条)+ 提交级 diff 与文件清单核查 + 发布页与 tag 核查 + 第三方后端仓核查 + Google News 中英文六组查询(走代理)+ Hacker News + arXiv。
本期索引
- 今日重点:昇腾 910 A2/A3 的 NPUIR 后端并入 mlir-ascend 主干,504 项自测全过;TileOPs 选择器改按忙时校准,manifest 遗憾从 5.79% 降到 0.48%;24 小时合入 24 笔、在审更新 48 笔、新开 issue 6 条
- 一、核心项目进展
- 1.1 主仓:标量读的索引系数不再混进 lane 数,行宽 32768 以上的编译崩溃关闭(#3453)(10-28)
- 1.2 主仓:批量 all-reduce 按标量分支解析线程区间,batch>1 时 2048 个输出全错的问题修复(#3341)(10-28)
- 1.3 主仓:SIMT 原子区域的形状校验补上,越界区域在编译期拒绝(#3465)(10-28)
- 1.4 主仓:稀疏 MMA 的 K 尾块不再被静默放过(#3367)(10-28)
- 1.5 主仓:atomic_addx4 的目标 dtype 收窄到 float*(#3369)(10-28)
- 1.6 主仓:非 restrict 参数的 handle 后缀不再被裁掉(#3366)(10-28)
- 1.7 主仓:TMA 描述符只记录已校验的名字(#3484)(10-28)
- 1.8 主仓:profiler 输出索引统一校验,标量与列表形式同一套边界规则(#3486)(10-28)
- 1.9 主仓:T.Parallel 的 coalesced_width 统一转 IntImm(#3374)(10-28)
- 1.10 主仓:测试基建把 cache 敏感用例改走 without_kernel_cache(#3695 / #3696)(10-28)
- 1.11 新增缺陷三条:SM120 上 13 个示例失败、IfStmtBinding 指数级编译、without_kernel_cache 不强制重编(#3701 / #3700 / #3699)(10-28)
- 二、多后端适配(昇腾 / 海光 / ROCm)
- 2.1 昇腾:910 A2/A3 的 NPUIR 后端并入 mlir-ascend 主干,504 项自测全通过、72 个示例仅 1 例环境性失败(#207)(10-28)
- 2.2 昇腾主仓:Bisheng 构建身份进二进制缓存键(#3691)(10-28)
- 2.3 昇腾主仓:手写流水线保留配置寄存器相位(#3624)(10-28)
- 2.4 昇腾仓在审:W4A4 int4×int4 GEMM 结果错误的根因与修法(#1731)(10-28)
- 2.5 昇腾仓在审:a2a3 调度器迁移,100 文件 18453 行(#1951)(10-28)
- 2.6 昇腾仓在审:Carver 按所选 SDK 读 A3 资源预算(#1955)(10-28)
- 2.7 昇腾仓在审:AI Core Exception Dump 的依赖降到 CANN 9.2.0(#1956)(10-28)
- 2.8 昇腾仓在审:文档改用 T.Tensor,弃用 T.Buffer DSL 接口(#1957)(10-28)
- 2.9 海光:HCU 普通拷贝的 LDS swizzle 特性分支,18 文件含 5 个新测试(#30)(10-28)
- 2.10 ROCm:pip 版 SDK 的 ROCm 10 支持合入,MI355X 不再被误判成 gfx900(#3549)(10-28)
- 三、生态与采用方(TileOPs / TileFoundry)
- 3.1 TileOPs:选择器改按忙时校准,manifest 遗憾 5.79% 降到 0.48%(#2784)(10-28)
- 3.2 TileOPs:非 NT 布局改走转置持久内核,2112×7168×4096 TN 从 195.3 降到 175.3 微秒(#2789)(10-28)
- 3.3 TileOPs:cumsum 长行改两段式种子扫描,(8,65536) 从 12.75 降到 5.18 微秒(#2790)(10-28)
- 3.4 TileOPs:extrema 归约改用 PTX 的 max.NaN,NaN 传播不变而回归消失(#2792)(10-28)
- 3.5 TileOPs:并行 cumsum 的有效行 tile、GDN 窄预填充分块两笔(#2788 / #2786)(10-28)
- 3.6 TileOPs 新开议题:背靠背调用下大 GEMM 行落后最快基线(#2791)(10-28)
- 3.7 TileFoundry:ring 阶段绑定与编译流程统一,70 文件(#277)(10-28)
- 3.8 TileOPs 在审:发布镜像把 TileOPs 0.1.0 与 TileFoundry 0.0.3 打在一起(#2793)(10-28)
- 四、社区、教程与活动
- 4.1 文档站两笔自动更新(10-28)
- 4.2 昇腾仓文档事实核查与安装修复批量在审,七条同批推进(10-28)
- 4.3 昇腾仓新增算子示例在审:GroupNorm 与 fla_chunk 两个算子(#1607 / #1953 / #1952)(10-28)
- 4.4 第三方:社区自动调优代理刷新中文说明(#74),Hexagon 后端仓库主分支窗口内无新提交(10-28)
- 4.5 新闻侧:Google News 六组中英文查询、Hacker News 与 arXiv 均无窗口内实质命中(10-28)
- 五、趋势观察
- 附、素材与核查说明
今日重点
今天窗口的重心在昇腾侧。tilelang-mlir-ascend 把 Ascend 910(A2/A3)的 NPUIR 后端整条并入主干,一笔 55 文件、21490 行新增:Python 侧是 tilelang/ascend910 包(backend / adapter / codegen / compiler / execution_backend / jit / pipeline / runtime / target 加 language 下的 kernel、copy、gemm、npuir 方言面),C++ 侧是 src/ascend910/target 下的开发者与专家两套发射器、op 层与 LowerNpuirBlock、NpuLoopVectorize、RewriteFDivToMul 三个 pass,运行时补了 npu_launcher.h 与 npu_utils,安装走 install_npuir.sh,target 别名接受 ascend910 / ascend-910 / a2 / a3 / npuir。作者给出的验收是仓库自带的两套用例:pytest testing/npuir 504 项全通过、examples 72 例中 71 例通过且 0 超时 0 报错,唯一失败的 deepseek_v32/fp8_lighting_indexer.py 在旧后端上同机同样失败,属环境与容差问题。第二笔提交是逐条对齐旧后端的行为差异,重点在 bf16 合法化(HIVM 不接受 bf16 操作数,改为 fp32 计算后显式转换)与动态 extent 的拷贝下沉。
TileOPs 今天合入七笔,两条主线都在性能侧。选择器不再贴着 manifest 拟合,改为在忙时(device_busy_ms)上按 160 个稠密 shape、60 个 m 9–32 的 shape、264 个 m≤8 的 shape 重新标定,用 softmin 期望对数遗憾做坐标下降;端到端忙时遗憾在留出的 manifest 上从 5.79% 降到 0.48%,m≤8 的留出集从 14.18% 降到 1.79%。GEMM 侧新增 coop2_swap:把非 NT 布局换成转置的持久内核,两个消费者 warpgroup 各持 128 行,尾写经 stmatrix.trans 与 TMA 转置落回,并支持宽窄 tile 混排(例如 2112 = 13×160 + 32 拆成 160 宽 tile 加 128 窄 tile),2112×7168×4096 TN 从 195.3 微秒降到 175.3 微秒,同机上另一块 H200 读到 174.3 降到 158.5 微秒,均压过 QuACK 的 187.9 与 167.0 微秒。
归约侧则同日先坏后修。上午 #2787 用 int32 键实现 extrema 归约的 NaN 传播、#2788 改并行 cumsum 的行 tile,两笔都在 07 时前后落地,随后 #2790 与 #2792 分别把两个回归收掉:cumsum 改成按种子分段、两次发射((8,65536) float32 从 12.75 微秒降到 5.18 微秒,同时修掉主干的 65535 tile 以上发射失败),extrema 改用 PTX 的 max.NaN / min.NaN 并在非整 16 字节向量的行上走单独内核(amax (2048,4096) bf16 从 8.352 微秒回到 6.976 微秒,快于 FlagGems)。这类”同一天内引入再修掉”的节奏说明 TileOPs 的夜间基准已经把回归拦在了当天。
一、核心项目进展
1.1 主仓:标量读的索引系数不再混进 lane 数,行宽 32768 以上的编译崩溃关闭(#3453)
合于 10-28 11:06,单文件 +8/-5。一个用向量化方式读整行、同时又按下标读该行单个元素的 kernel,在行宽对 65536 取模后不小于 32768 时编译报 InternalError: Can’t fetch the lanes of a scalable vector。根因在 VectorTypeAccessChecker::OnArrayAccess:它把标量读记为 access_dtype.with_lanes(me->coeff),而 me->coeff 是读下标的模系数,对 x[row, 0] 就是行宽;DataType 的 lane 数只有 16 位并按 int16 读回,32768 到 65535 因此被当成可伸缩向量,lanes() 在 get_preferred_dtype 里直接抛错,更宽的倍数则静默回绕。修法是把系数单独存进 int64_t,兼容性判断改用它取模,标量读的 dtype 只保留元素类型。这个缺陷的原始触发场景是 TileOPs 的 GroupNorm:Stable Diffusion 的 40960 与 163840 两个组宽都会踩到。验证用两次 wheel 构建对比:主干在 32768 与 40960 两个宽度上编译失败,本笔在所有测试宽度上都能构建。
来源:https://github.com/tile-ai/tilelang/pull/3453
1.2 主仓:批量 all-reduce 按标量分支解析线程区间,batch>1 时 2048 个输出全错的问题修复(#3341)
合于 10-28 05:46,2 文件 +170/-8。标量 all-reduce 分支在发射调用前会先把参与线程的区间解析出来(reducing_threads 大于 32 且目标支持命名栅栏时走 ResolveAllReduceThreadRange),批量分支却直接取块边界当偏移与栅栏宽度。当归约线程是 [32, 96) 这种不从 0 开始的局部区间时,批量调用从 0 号线程索引、按整块等栅栏,非参与线程未初始化的寄存器被折进归约:batch=1 正确,batch=2 在同一 kernel 上 2048 个输出元素全部错,最大绝对误差 1083,且没有任何诊断。修法是让批量分支与标量分支同一套解析,把解析出的偏移与范围传给 MakeBatchAllReduce,工作区 stride 仍保持块大小。新增测试在同一 kernel 上跑 batch=1 与 batch=2 并对 torch 参考比对。
来源:https://github.com/tile-ai/tilelang/pull/3341
1.3 主仓:SIMT 原子区域的形状校验补上,越界区域在编译期拒绝(#3465)
合于 10-28 08:14。SIMT 原子操作此前不校验传入区域的形状,形状不匹配的写法一路走到 codegen 才表现出错误结果或崩。本笔在 CUDA 与 ROCm 两条路径上补齐区域形状校验,把不合法区域提前到编译期拒绝。
来源:https://github.com/tile-ai/tilelang/pull/3465
1.4 主仓:稀疏 MMA 的 K 尾块不再被静默放过(#3367)
合于 10-28 16:01。稀疏 MMA 在 K 方向上有未处理的尾块时,此前不报错而是沿用整块的搬运与计算参数,数值静默偏差。本笔显式拒绝这类 K 尾块,把问题挪到编译期。
来源:https://github.com/tile-ai/tilelang/pull/3367
1.5 主仓:atomic_addx4 的目标 dtype 收窄到 float*(#3369)
合于 10-28 16:08,3 文件 +27/-9。CUDA 的 atomic_addx4 与 atomic_addx4_ret 通用实现此前接受任意目标指针,int32 或 int64 目标会落到 float 语义上算错。本笔把两个通用重载限制为 float* 目标,fp16 与 bf16 的既有特化不动;Python 侧 API 保持类型中立,CPU 后端的整数目标不受影响。CUDA 侧新增整数目标编译失败的回归用例,CPU C 后端新增两种整数宽度的累加正确性用例。
来源:https://github.com/tile-ai/tilelang/pull/3369
1.6 主仓:非 restrict 参数的 handle 后缀不再被裁掉(#3366)
合于 10-28 13:19。参数去重时按去掉 _handle 后缀的名字比较,导致同一个内核里 A 与 A_handle 两个不同参数被当成同一个,代码生成只发一份。修法改为按精确名加变量身份去重,并新增 CUDA 侧 A 与 A_handle 并存的回归用例。
来源:https://github.com/tile-ai/tilelang/pull/3366
1.7 主仓:TMA 描述符只记录已校验的名字(#3484)
合于 10-28 17:12。TMA 描述符的登记此前在校验之前完成,校验失败的描述符名留在表里,后续同名的合法描述符会被误判为重复。本笔把登记挪到校验之后。
来源:https://github.com/tile-ai/tilelang/pull/3484
1.8 主仓:profiler 输出索引统一校验,标量与列表形式同一套边界规则(#3486)
合于 10-28 19:24,2 文件 +15/-9,修 #2884。profiler 与 adapter 两侧此前对 result_idx 的处理不一致:标量形式做过边界检查、列表形式没有,负索引也未归一。本笔让 BaseKernelAdapter 与 Profiler 共用一套规则,接受 None、整数或整数列表,按 -len(params) 到 len(params)-1 判定,合法负索引归一、越界与非整数元素抛 ValueError。新增 28 项 CPU 侧回归覆盖标量与列表两种路径。
来源:https://github.com/tile-ai/tilelang/pull/3486
1.9 主仓:T.Parallel 的 coalesced_width 统一转 IntImm(#3374)
合于 10-28 16:27。coalesced_width 传 Python 整数时,部分路径直接以原生整数进入 pass,类型断言在后续变换里失败。本笔在语言层统一转成 IntImm,消除写法差异带来的失败。
来源:https://github.com/tile-ai/tilelang/pull/3374
1.10 主仓:测试基建把 cache 敏感用例改走 without_kernel_cache(#3695 / #3696)
合于 10-28 01:59 与 09:19。承接昨日记录的 issue #3694(内核缓存命中破坏假定冷编译的测试),两笔把检查编译产物、源码与诊断信息的用例改走 without_kernel_cache() 或显式 .compile(…),让断言看到的是本次构建的结果而不是进程内缓存里的旧 kernel。
来源:https://github.com/tile-ai/tilelang/pull/3696
来源:https://github.com/tile-ai/tilelang/pull/3695
1.11 新增缺陷三条:SM120 上 13 个示例失败、IfStmtBinding 指数级编译、without_kernel_cache 不强制重编(#3701 / #3700 / #3699)
三条新开 issue 指向测试与编译时间。其一,#3701 在 RTX PRO 6000 Blackwell(sm_120,动态共享内存上限 99 KiB)上跑完整 examples 套件,CUDA 与 CuTeDSL 两个后端都是 13 失败 199 通过 27 跳过,失败集逐字相同:7 例是示例在非 –tune 路径上硬编码 Hopper 尺寸的 tile(如 attention_sink 要 163840 字节、deepseek_mla 解码要 231424 字节),4 例是参照实现 flash-attn 2.5.8 没有 sm_12x 内核,另两例分别是 ReduceOp 下沉与 ThreadSync 无法证明线程范围。其二,#3700 报告在循环里重复写 x = x + x 并加上 i < N 守卫时,tl.LayoutInference 从 0.087 秒涨到 118.5 秒,关掉 replayable bind 内联即退回 0.100 秒。其三,#3699 指出 without_kernel_cache() 只翻转 CacheState(只被磁盘缓存读取),而 @tilelang.jit 包装器在 JITImpl.compile() 之前还会查进程内的 _call_form_cache 与 _kernel_cache,因此该块内的一次普通调用并不保证重新编译。
来源:https://github.com/tile-ai/tilelang/issues/3701
来源:https://github.com/tile-ai/tilelang/issues/3700
来源:https://github.com/tile-ai/tilelang/issues/3699
二、多后端适配(昇腾 / 海光 / ROCm)
2.1 昇腾:910 A2/A3 的 NPUIR 后端并入 mlir-ascend 主干,504 项自测全通过、72 个示例仅 1 例环境性失败(#207)
合于 10-28 08:04,55 文件 +21490/-1,作者 CeleNewYear,基于 v0.1.15 的 a35f8ddf,两个提交分别负责后端落地与行为对齐。后端面包含 tilelang/ascend910 下的 backend、adapter、codegen、compiler、execution_backend、jit、pipeline、runtime、target,以及 language 下的 kernel、copy、gemm、npuir 方言面;C++ 侧是 src/ascend910/target 的开发者与专家两套发射器、src/ascend910/op 算子层与 src/ascend910/transform 的 LowerNpuirBlock、NpuLoopVectorize、RewriteFDivToMul;运行时要 npu_launcher.h、npu_utils.cpp、npu_utils.py 生成 torch 兼容的启动包装并按 *_infer_workspace_shape_function 估每块 scratch;构建侧接 CMakeLists 并给 install_npuir.sh 的 –bishengir-path 选项;target 别名接受 ascend910、ascend-910、a2、a3、npuir,模式开关 TILELANG_ASCEND_MODE 取 expert 或 developer,默认 expert。验收数据:pytest testing/npuir 504 通过 0 失败;examples/run_all.py –sequential 72 例中 71 通过 1 失败,0 超时 0 报错,唯一失败的 deepseek_v32/fp8_lighting_indexer.py 在旧后端同机同样失败。第二笔提交对齐的差异包括 bf16 合法化(HIVM 拒绝 bf16 操作数,改为 fp32 计算后显式转换)与动态 extent 区域的拷贝下沉。
来源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
2.2 昇腾主仓:Bisheng 构建身份进二进制缓存键(#3691)
合于 10-28 10:26。昇腾侧的二进制缓存此前只按内核与 target 做键,换 Bisheng 工具链版本后仍会命中旧二进制。本笔把构建身份并入缓存键,工具链变化即失效重编。
来源:https://github.com/tile-ai/tilelang/pull/3691
2.3 昇腾主仓:手写流水线保留配置寄存器相位(#3624)
合于 10-28 09:57。手动流水线在改写阶段丢失配置寄存器的相位信息,导致同一份配置在正确的拍上不生效。本笔在改写中保留相位,恢复与自动流水线一致的时序语义。
来源:https://github.com/tile-ai/tilelang/pull/3624
2.4 昇腾仓在审:W4A4 int4×int4 GEMM 结果错误的根因与修法(#1731)
在审,7 文件 +634/-24,修 #438。根因是整条 C++ 模板栈都用 sizeof(AscendC::int4b_t) == 1 推分形几何,而 int4b_t 是半字节打包类型:32 字节的 C0 块正确元素数是 64 而不是 32,于是所有拷贝 stride、repeat 次数与字节计数全部 2 倍偏差。三层叠加:L1 到 L0A/L0B 的 LoadData 参数由该常数推出,数据堆到错误位置;B 操作数的 zN 到 nZ 转置被 s4 专用指令 load_cbuf_to_cb_s4 静默关掉,必须改走 LoadDataWithTranspose 到 load_cbuf_to_cb_transpose_s4,而其参数语义在 CANN 头文件里没有文档,作者只能用探针 kernel 在设备上逆向(正确参数 repeatTimes = N/64、srcStride = roundUp16(K)/64、dstGap = 3,且 s4 版本按 2048 字节块而不是文档写的 512 字节 fractal 工作);UB 到 GM 的字节长度按每元素 1 字节算,每行多搬一倍并混入相邻行。mad_s4 指令本身正确,症状是 99.6% 不匹配、输出 -501/627 而正确答案 641/1785。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1731
2.5 昇腾仓在审:a2a3 调度器迁移,100 文件 18453 行(#1951)
在审,100 文件 +18453/-666,作者 erhsh。把 A2/A3 的调度器按新架构迁移,是窗口内最大的一笔在审改动。PR 描述表格未填,具体边界需按文件清单核对。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1951
2.6 昇腾仓在审:Carver 按所选 SDK 读 A3 资源预算(#1955)
在审,3 文件 +203/-27。在 Ascend910_9392 上 Carver 回落到 Ascend910A,把 256×256 的 FP32 累加器放到 256 KiB 的 L0C 预算上,而所选 SDK 只给 128 KiB。修法是识别 9392 与 9382 两个具体 SKU、用既有的 ASCEND_HOME 选择读各自平台的 INI(另留跨编译用的显式文件选项),把 SDK 容量裁剪到当前后端 arena 同时保留原始容量、配置身份、BF16 支持、核数与 L2 属性。六个定向用例覆盖设备、target 与显式芯片探测、超出与合法的 L0 tile、以及 9382 关闭 BF16 的独立配置;作者给出的回归是原生昇腾 CI 主套 1352 通过 3 跳过、示例套 153 通过。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1955
2.7 昇腾仓在审:AI Core Exception Dump 的依赖降到 CANN 9.2.0(#1956)
在审,5 文件 +49/-32。该特性此前声明依赖 CANN 9.3.0,作者在 9.2.0 上验证可行并修两处兼容:acldumpTensorInfo 补齐 addrType=ACL_DUMP_ADDR_RAW、placement、shapeNum、originShapeNum 四个必填字段(9.2.0 的 acldumpSaveExceptionInfo 强校验 addrType,原实现被 calloc 置零成 ACL_DUMP_ADDR_PTR 而被拒),并跳过 NULL 地址或零大小的 tensor;注册接口从已弃用的 aclrtSetExceptionInfoCallback 换成 aclrtExceptionInfoCallbackRegister。解析侧兼容 msaicerr.py 的两种产物(9.2.0 出 .npy,旧版出 .bin),文档的版本要求同步改为 CANN 不低于 9.2.0。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1956
2.8 昇腾仓在审:文档改用 T.Tensor,弃用 T.Buffer DSL 接口(#1957)
在审,单文件 +124/-122。编程指南 3.1 节的 kernel 参数定义与上游 API 对齐:文本与示例统一用 T.Tensor,并补一句说明 T.Buffer(…) 与 T.Buffer[…] 的弃用只针对 TileLang DSL 接口,底层 TVM 的 Buffer 类型与概念并未弃用。依据是 ascendc_pto 分支 tilelang/language/proxy.py 里的 @deprecated 装饰器。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1957
2.9 海光:HCU 普通拷贝的 LDS swizzle 特性分支,18 文件含 5 个新测试(#30)
在审,18 文件 +2000/-84,分支 Feat/hcu plain copy lds swizzle 于 10-28 13:59 推送。改动集中在新增的 gemm_lds_strategy_utils(+148/-0 与 +29/-0)、annotate_mls_gemm_dep.cc(+332/-45)与 ds_read_format.cc(+61/-22),copy 与 merge_shared_memory_allocations 有配套小改;测试侧一次新增 test_hcu_plain_copy_lds_swizzle(396 行)、_formula(237)、_formula_b8(220)、fp8(200)与 _b8_write_phases(80)。方向是普通拷贝路径的 LDS 写相位与 swizzle 策略,配 b8/fp8 两类数据宽度。
来源:https://github.com/tile-ai/tilelang-hygon/pull/30
2.10 ROCm:pip 版 SDK 的 ROCm 10 支持合入,MI355X 不再被误判成 gfx900(#3549)
合于 10-28 07:38,2 文件 +106/-19。ROCm 10.x 的 rocm/pytorch 镜像把 ROCm 以 Python wheel 形式发布(rocm-sdk 与 _rocm_sdk_devel),既没有 /opt/rocm 目录树也没有 rocm-* 的 apt 包。get_rocm_arch() 此前硬编码 /opt/rocm 前缀,路径不存在时回落到 gfx900,于是 MI355X(gfx950)上 rocwmma 的静态断言直接失败;设 ROCM_PATH 也无效,因为 TVM 的 tvm_callback_rocm_get_arch 用默认前缀调用该函数。修法是按调用方前缀、find_rocm_path()、torch 的 gcnArchName、gfx900 兜底的顺序解析,每次 rocminfo 探测加 10 秒超时;另新增 docker/Dockerfile.rocm10,基于 rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12,去掉 rocm-* 的 apt 安装、把 ROCM_PATH 指向 rocm-sdk path –root、保留基础镜像的 torch 构建并加 import 自检。作者在 MI355X 上实测架构探测从 gfx900 变为 gfx950,testing/python/amd 套件按正确架构构建并通过。
来源:https://github.com/tile-ai/tilelang/pull/3549
三、生态与采用方(TileOPs / TileFoundry)
3.1 TileOPs:选择器改按忙时校准,manifest 遗憾 5.79% 降到 0.48%(#2784)
合于 10-28 10:51,4 文件 +259/-203。稠密选择器的打分改为一个 tensor_core_tflops 乘结构因子,此前结构差异只体现在一个 max() 里,内存受限的 shape 会把 basic、simple 与 coop2 打成分;launch_us 项删除(在忙时上拟合为 0)。m≤8 的 NT 带用流式模型替换手测阈值:每 CTA 速率取 min(ns·rows·bk·2/latency, cta_gbps),波次按 SMEM 受限的驻留算,再加 DRAM 下限 n·k·2/dram_tbps 与 split-K 的归约项。TN/TT 的持久内核路径打通(A/B 的 SMEM tile 跟随存储布局,transpose_A/B),176/192 宽度、窄 tile、ping-pong 与 split-K 仍限 NT。标定在忙时上做,三组网格 160/60/264 个 shape、7 轮交错 CUPTI 取中位,softmin 期望对数遗憾按坐标下降、tau 0.1→0.05→0.02 退火。端到端结果:留出的 manifest 5.79% 到 0.48%,标定网格 9.28% 到 2.65%,m 9–32 3.13% 到 2.09%,m≤8 11.20% 到 1.17%,m≤8 留出集 14.18% 到 1.79%。
来源:https://github.com/tile-ai/TileOPs/pull/2784
3.2 TileOPs:非 NT 布局改走转置持久内核,2112×7168×4096 TN 从 195.3 降到 175.3 微秒(#2789)
合于 10-28 21:06,4 文件 +620/-15,解决 #2785。新增 coop2_swap:对 TN/TT/NN 计算 Cᵀ = op(B)ᵀ·op(A)ᵀ,按 256×block_n 的 tile 走持久内核,两个消费者 warpgroup 各持 128 行,尾写经 stmatrix.trans 与 TMA 转置写回。混排 tile 宽度(narrow_n 与 row_mix)让 M 不被 padding 浪费,例如 2112 = 13×160 + 32 拆成宽 160 加窄 128,宽的先跑、窄的收尾,拆分点沿用 coop2 窄 tile 的 _busiest_cta 指标。选择器在 NT 之外枚举 coop2_swap 的 block_n 128 与 160、结构因子 1.2 并按忙时重拟合。H200 忙时:2112×7168×4096 TN 从 195.3 降到 175.3 微秒(QuACK 187.9),24576×1536×4096 TN 从 421.9 降到 416.4 微秒;另一块 H200 上分别为 174.3 降到 158.5(QuACK 167.0)与 384–402 微秒区间。行组规则把固定 4 改成 16,在 24 个 TN/TT/NN 标定 shape 上与 group 4 相差 0.989–1.022 倍;另有 9 个 TN shape 上混排比均匀 coop2_swap 快 3.6% 到 11%。
来源:https://github.com/tile-ai/TileOPs/pull/2789
3.3 TileOPs:cumsum 长行改两段式种子扫描,(8,65536) 从 12.75 降到 5.18 微秒(#2790)
合于 10-28 20:45,4 文件 +255/-246。少量长行的 cumsum 改用种子分段扫描,发射次数从三次降到两次,既消掉 #2788 引入的回归,也修掉其布局失败。规则是行不短于 64 KiB 且行数不多时走分段扫描,cumprod 不变;另修掉主干上行数超过 65535 个 tile(如 (1, 2**24) float32)时的发射失败。H200 @1500 MHz 忙时: (8,65536) float32 12.75 到 5.18 微秒,(16,65536) 19.14 到 7.01 微秒,(1,65536) bf16 9.03 到 3.45 微秒,(3,8193) fp16 7.52 到 3.29 微秒。384 个 shape 的扫描里没有一项比 #2788 之前更慢,CumsumFwdOp 与 CumprodFwdOp 的 manifest 行均快于最快对照实现;新增 few-long-seq 负载(8×65536 float32)5.20 微秒对 FlagGems 6.30 微秒。
来源:https://github.com/tile-ai/TileOPs/pull/2790
3.4 TileOPs:extrema 归约改用 PTX 的 max.NaN,NaN 传播不变而回归消失(#2792)
合于 10-28 23:31,4 文件 +346/-119。amax/amin 改为在 fp32 上用 PTX 的 max.NaN 与 min.NaN,不再走 #2787 的 int32 键方案,NaN 传播语义不变而当天早些时候的回归被收掉;不是整 16 字节向量的行另走一个 extrema 行内核,因为 T.reduce_max 在 fp32 上会丢掉 NaN。H200 @1500 MHz:amax (2048,4096) bf16 在 #2787 之前是 7.008 微秒、主干 8.352、本笔 6.976;fp16 7.008 到 7.775 到 6.944;amax (64,32768) bf16 3.552 到 4.544 到 3.520;amin 同形 3.552 到 4.224 到 3.520。
来源:https://github.com/tile-ai/TileOPs/pull/2792
3.5 TileOPs:并行 cumsum 的有效行 tile、GDN 窄预填充分块两笔(#2788 / #2786)
#2788 于 10-28 07:13 合入,让并行 cumsum 只处理有效的行 tile,是全窗口内最早的一笔,随后被 #2790 以两段式扫描替换掉问题路径。#2786 于 07:09 合入,2 文件 +30/-2,在 SM89 上把 32 列的 GDN 预填充计划改用既有的带融合输出的状态扫描,把整块分段与尾块分开处理,保留准备阶段、输出与最终状态三个接口;选用条件为稠密输入、batch 1、16 个 value head、状态宽 128、key major、序列长大于 16384 且不超过 65536。H=4 时中位设备执行时间相对 TileOPs 基线降 13.65% 到 18.31%,设备活动跨度降 13.48% 到 17.07%;与 FLA 交替对比的四轮里 H=16、长 65536、BF16 的跨度差异在 -2.57% 到 +1.77% 之间,作者自述该行落在测量波动内。
来源:https://github.com/tile-ai/TileOPs/pull/2788
来源:https://github.com/tile-ai/TileOPs/pull/2786
3.6 TileOPs 新开议题:背靠背调用下大 GEMM 行落后最快基线(#2791)
新开,承接 #2133 与 #2789。bench_gemm 单次读一个调用、中间有 L2 冲刷与同步,调用从功耗上限以下起步,能耗代价被掩盖;改成训练里的背靠背连续调用后,十条 manifest 大行里有六条落后最快基线。典型一行:prefill-o-proj 4096×7168×16384 NT 本实现 1412.7 微秒,QuACK 1371.1、cuBLAS 1360.1、DeepGEMM 1359.2,比值 0.962。方法为单卡 H200、时钟解锁、预热 1 秒后计时 1.5 秒、四轮交替取中位。
来源:https://github.com/tile-ai/TileOPs/issues/2791
3.7 TileFoundry:ring 阶段绑定与编译流程统一,70 文件(#277)
合于 10-28 16:55,70 文件 +1567/-1298,取代 #276。动机有两处:未绑定的 ring 视图与阶段选择会带着没有 SSA 名字的视图走到 CUDA 操作数上;三个编译入口各自拼装不同的 pass 列表。改法是把 ring 阶段保存为一组已绑定的 TensorView、每个消费作用域只绑定一次,让 default_pipeline 在 finalize、compile 与 scaffold 三处共用,核心宿主入口 pass 经 Target.resolve 解析并由 CUDA 提供 @bind 实现;Module.entry 成为宿主起点,分析的共享内存峰值带进宿主启动;调用链上的所有 CPU 宿主都参与链接,根宿主排在最前。契约变化:link_program 去掉 entries,加载器绑定第一个被链接的宿主,build 要求存在 CPU 宿主根。遗留风险写在 PR 里:#278 跟踪 Broadcast 布局归一,#255 跟踪 SymbolRef 与函数引用。
来源:https://github.com/tile-ai/TileFoundry/pull/277
3.8 TileOPs 在审:发布镜像把 TileOPs 0.1.0 与 TileFoundry 0.0.3 打在一起(#2793)
在审,9 文件 +120/-44。.github/runner/Dockerfile 新增 release 目标,在开发镜像之上装 tileops 与 tilefoundry 的 PyPI 发布版,镜像名以 TILEOPS_VERSION 与 TILEFOUNDRY_VERSION 命名,推为 ghcr.io/tile-ai/tileops-release:tileops0.1.0-tilefoundry0.0.3-cuda13.2-torch2.13-tilelang0.1.12;配套 scripts/verify_release_image.sh 拉镜像后在 GPU 上跑栈与基线检查、打包测试与指定的 –bench 文件,用例按镜像内 tileops 版本对应的 git tag 取(wheel 里不带测试与基准)。同时移除 sgl-kernel 基线:没有任何镜像再装它,bench_moe_permute_align 一直跳过它。
来源:https://github.com/tile-ai/TileOPs/pull/2793
四、社区、教程与活动
4.1 文档站两笔自动更新(10-28)
tilelang.github.io 在 10-28 08:20 与 13:31 各有一笔由 github-actions 提交的 Update docs,内容为文档构建产物随主干同步,无独立内容变更。
来源:https://github.com/tile-ai/tilelang.github.io
4.2 昇腾仓文档事实核查与安装修复批量在审,七条同批推进(10-28)
tilelang-ascend 的文档事实核查流水线在窗口内更新七条:T.tile.bitwise_lshift、bitwise_or、bitwise_and、fill、clear 五个 API 的核查,clamp 在 PTO 上生效并核对文档(#1620),以及 axpy 的 docstring 与测试加强(#1600)。另有 #1943 把 SHMEM 安装改成构建失败即停止,与 #1942 的快照分析器校验同属机器人提交。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1659
来源:https://github.com/tile-ai/tilelang-ascend/pull/1620
4.3 昇腾仓新增算子示例在审:GroupNorm 与 fla_chunk 两个算子(#1607 / #1953 / #1952)
三条示例类 PR 在窗口内更新:#1607 给 GroupNorm 加 20 个 cann-bench 用例,#1953 加 fla_chunk_o 算子与测试,#1952 加 fla_chunk_intra_token_parallel 算子与测试。这也是 1.1 那个标量读崩溃缺陷的原始场景来源(GroupNorm 读每行首元素)。
来源:https://github.com/tile-ai/tilelang-ascend/pull/1607
来源:https://github.com/tile-ai/tilelang-ascend/pull/1953
4.4 第三方:社区自动调优代理刷新中文说明(#74),Hexagon 后端仓库主分支窗口内无新提交(10-28)
社区仓库 cc-c122/TileLang-Remote-Autotuning-Agent 于 10-28 09:21 合入 #74,按当前 V2 工作流刷新中文 README,仓库此前在 10-24 还有一次基线取消状态与 mcprofiler 归因的合并。另一个第三方后端仓库 xwhzz/tilelang-hexagon 的组织视图显示 10-29 06:56 有推送时间戳,但主分支最新提交停在 08-24(CI 依赖升级),窗口内主分支无代码提交,昨天的 15 笔属特性分支活动,本条不作为新增内容计入。
来源:https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent/pull/74
来源:https://github.com/xwhzz/tilelang-hexagon
4.5 新闻侧:Google News 六组中英文查询、Hacker News 与 arXiv 均无窗口内实质命中(10-28)
六组查询包含 TileLang 精确词加两日、中文算子词加七日、英文 GPU kernel DSL 加三日、组织名加七日、中文厂商词(昇腾 / 沐曦 / 海光)加七日以及算子库词加两日。返回结果里与 TileLang 相关但落在窗口内的只有一篇转载稿(昇腾 950PR 矩阵算力对比,属二手市场评论,无原始数据,不收录),其余为 DeepSeek 融资与技术类泛报道,标题语义不含 TileLang 相关内容。英文查询全部空返回。Hacker News 当日条目为一门名为 Vu 的新语言发布,与 TileLang 无关。arXiv 无窗口内新增的 tilelang 或 tiled DSL 论文。新闻侧本日与 GitHub 侧的信号强度差距继续扩大,日报主体仍以仓库为准。
五、趋势观察
昇腾侧的推进方式这两天已经很清楚:主仓按小时合入小修(缓存键、寄存器相位),专用仓 mlir-ascend 用一次两万行的合并把 910 的 NPUIR 路径从特性分支搬进主干,tilelang-ascend 则在 PTO 与 AscendC 双后端的既定节奏上做文档与算子补齐。三处节奏不同但指向同一件事:910 的编译器路径正在从”能跑”转向”可验收”,504 项自测与 71/72 示例就是这次给出的验收口径。值得注意的是在审队列里 a2a3 的调度器迁移占到了 100 文件,说明下一阶段的变动量在调度层而不在算子层。
性能侧的看点在测量口径本身。TileOPs 这两笔把选择器和 GEMM 的结论都挪到了忙时(device_busy_ms),并在同一份 PR 里给出单次调用与背靠背调用的差别:单次读法下本实现领先 QuACK,背靠背、功率封顶的训练场景里十条大行有六条落后。这是一个比”某行快了多少微秒”更重要的信号,因为训练实际就是背靠背。#2791 把这件事单独立项,说明团队已经把口径问题当成产品问题处理,而不是当作基准偏差粉饰过去。
工程卫生上今天有三条值得记:一是同一批内的自伤与自修(#2787/#2788 上午合入、#2790/#2792 下午收回归),说明夜间基准的兜底在起作用,但也说明归约这套代码的语义边界仍然脆;二是 without_kernel_cache() 的语义漏洞(#3699)被从迁移过程中揪出来,进程内两层缓存不吃 CacheState,测试自以为在冷编译其实拿的是旧 kernel,这类缺陷会让”我们的测试全过”失去意义;三是 SM120 上 13 个示例失败(#3701)暴露出示例硬编码 Hopper 尺寸的问题,Blackwell 装机量上来后这类”示例不是产品”的债务会被反复计入。
附、素材与核查说明
信源核查表
| 信源 | 核查结果 |
|---|---|
| GitHub tile-ai 组织逐仓(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文档站) | 窗口内有推送:tilelang、TileOPs、TileFoundry、tilelang-mlir-ascend、tilelang-hygon、tilelang.github.io;tilelang-metax、tilelang-musa、TileRT、tilescale、DeepStack、tilelang-puzzles 无窗口内主分支提交 |
| 合入真值检索(is:pr is:merged) | 24 笔:tilelang 15、TileOPs 7、TileFoundry 1、tilelang-mlir-ascend 1 |
| 在审真值检索(is:pr is:open updated) | 48 笔,含昇腾仓文档核查批次与 a2a3 迁移 |
| 新增 issue 核对 | 6 条,其中 1 条为昇腾仓每日测试机器人 |
| 主分支提交级核查 | tilelang 15 笔、TileOPs 7 笔、TileFoundry 1 笔、文档站 2 笔;合入数与主分支提交数一致 |
| 发布页与 tag | tilelang 最新仍为 v0.1.15(09-30 发布),窗口内无新版本;TileOPs 最新为 v0.1.0 |
| 第三方后端仓 | Hexagon 后端主分支最新提交停在 08-24,窗口内无新提交;自动调优代理合入中文说明刷新;沐曦仓六个分支无窗口内提交 |
| Google News(中英六组) | 仅一篇二手转载稿涉及,无原始数据,不收录 |
| Hacker News | 当日无 TileLang 相关条目 |
| arXiv | 无窗口内新增 tilelang 或 tiled DSL 论文 |
完整信源清单
- [1] tilelang 标量读索引系数修正(#3453)— https://github.com/tile-ai/tilelang/pull/3453
- [2] tilelang 批量 all-reduce 线程区间(#3341)— https://github.com/tile-ai/tilelang/pull/3341
- [3] tilelang SIMT 原子区域形状校验(#3465)— https://github.com/tile-ai/tilelang/pull/3465
- [4] tilelang 稀疏 MMA K 尾块拒绝(#3367)— https://github.com/tile-ai/tilelang/pull/3367
- [5] tilelang atomic_addx4 目标 dtype(#3369)— https://github.com/tile-ai/tilelang/pull/3369
- [6] tilelang 非 restrict 参数 handle 后缀(#3366)— https://github.com/tile-ai/tilelang/pull/3366
- [7] tilelang TMA 描述符名校验(#3484)— https://github.com/tile-ai/tilelang/pull/3484
- [8] tilelang profiler 输出索引校验(#3486)— https://github.com/tile-ai/tilelang/pull/3486
- [9] tilelang coalesced_width 转 IntImm(#3374)— https://github.com/tile-ai/tilelang/pull/3374
- [10] tilelang 测试改走 without_kernel_cache(#3696)— https://github.com/tile-ai/tilelang/pull/3696
- [11] tilelang 缓存敏感用例迁移(#3695)— https://github.com/tile-ai/tilelang/pull/3695
- [12] tilelang SM120 示例失败议题(#3701)— https://github.com/tile-ai/tilelang/issues/3701
- [13] tilelang IfStmtBinding 编译时间议题(#3700)— https://github.com/tile-ai/tilelang/issues/3700
- [14] tilelang without_kernel_cache 语义议题(#3699)— https://github.com/tile-ai/tilelang/issues/3699
- [15] tilelang-mlir-ascend 910 A2/A3 NPUIR 后端(#207)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
- [16] tilelang Bisheng 构建身份进缓存键(#3691)— https://github.com/tile-ai/tilelang/pull/3691
- [17] tilelang 手写流水线配置寄存器相位(#3624)— https://github.com/tile-ai/tilelang/pull/3624
- [18] tilelang-ascend W4A4 int4 GEMM 修复(#1731)— https://github.com/tile-ai/tilelang-ascend/pull/1731
- [19] tilelang-ascend a2a3 调度器迁移(#1951)— https://github.com/tile-ai/tilelang-ascend/pull/1951
- [20] tilelang-ascend Carver A3 资源预算(#1955)— https://github.com/tile-ai/tilelang-ascend/pull/1955
- [21] tilelang-ascend CANN 9.2.0 异常转储兼容(#1956)— https://github.com/tile-ai/tilelang-ascend/pull/1956
- [22] tilelang-ascend 文档改用 T.Tensor(#1957)— https://github.com/tile-ai/tilelang-ascend/pull/1957
- [23] tilelang-hygon HCU 普通拷贝 LDS swizzle(#30)— https://github.com/tile-ai/tilelang-hygon/pull/30
- [24] tilelang pip 版 ROCm 10 支持(#3549)— https://github.com/tile-ai/tilelang/pull/3549
- [25] TileOPs 选择器忙时校准(#2784)— https://github.com/tile-ai/TileOPs/pull/2784
- [26] TileOPs coop2_swap 转置持久内核(#2789)— https://github.com/tile-ai/TileOPs/pull/2789
- [27] TileOPs cumsum 两段式种子扫描(#2790)— https://github.com/tile-ai/TileOPs/pull/2790
- [28] TileOPs extrema 用 max.NaN(#2792)— https://github.com/tile-ai/TileOPs/pull/2792
- [29] TileOPs 并行 cumsum 有效行 tile(#2788)— https://github.com/tile-ai/TileOPs/pull/2788
- [30] TileOPs GDN 窄预填充分块(#2786)— https://github.com/tile-ai/TileOPs/pull/2786
- [31] TileOPs 背靠背 GEMM 议题(#2791)— https://github.com/tile-ai/TileOPs/issues/2791
- [32] TileFoundry ring 阶段绑定与编译流程统一(#277)— https://github.com/tile-ai/TileFoundry/pull/277
- [33] TileOPs 发布镜像(#2793)— https://github.com/tile-ai/TileOPs/pull/2793
- [34] tilelang 文档站 — https://github.com/tile-ai/tilelang.github.io
- [35] tilelang-ascend 文档核查 bitwise_lshift(#1659)— https://github.com/tile-ai/tilelang-ascend/pull/1659
- [36] tilelang-ascend clamp 在 PTO 生效(#1620)— https://github.com/tile-ai/tilelang-ascend/pull/1620
- [37] tilelang-ascend GroupNorm 示例(#1607)— https://github.com/tile-ai/tilelang-ascend/pull/1607
- [38] tilelang-ascend fla_chunk_o 示例(#1953)— https://github.com/tile-ai/tilelang-ascend/pull/1953
- [39] 社区自动调优代理中文说明(#74)— https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent/pull/74
- [40] tilelang-hexagon 后端仓 — https://github.com/xwhzz/tilelang-hexagon
- [41] tilelang 发布页 — https://github.com/tile-ai/tilelang/releases
- [42] TileOPs 发布页 — https://github.com/tile-ai/TileOPs/releases
- [43] Google News 检索 — https://news.google.com/