TileLang周报(2026-10-26 ~ 2026-10-30)
报告周期:2026 年 10 月 26 日(周一)至 10 月 30 日(周五),共 5 个工作日 素材来源:本刊每日 TileLang 社区动态日报 5 期(10-26、10-27、10-28、10-29、10-30),五期采集窗口合计 10-23 07:00 至 10-30 07:00(周一那期为 72 小时窗口,覆盖上周末;周二至周五为 24 小时);另经 tile-ai 组织各仓库 168 小时窗口补充核查、各仓发布页与标签、VERSION 文件复核(详见附录)
一、本周要闻
-
TileOPs 0.1.0 正式发布,206 个算子的首个正式版,公开 API 就此冻结(10-27):PyPI 制品于 10-27 07:45(北京时间)上传,GitHub 发布页 13:15 挂出,是仓库首个正式 tag。发布说明给出的硬指标是三组:206 个算子全部有 manifest 条目(签名、测试与基准负载、roofline 公式,CI 逐条校验);H200 上 2265 行夜间基准中 2262 行不慢于最快对照实现、几何平均 1.37 倍,对照方包括 torch.compile、FlashAttention 2 与 3、FlashInfer、FlashMLA、vLLM、DeepGEMM、flash-linear-attention、FlagGems 与 QuACK;206 个算子里 195 笔带 foundry 标签的合入由 TileFoundry 生成内核并通过同一套 CI 进库。发布同时冻结公开 API,新增 API 兼容性检查脚本,任何按上一版写法调用而会中断的 PR 直接失败,此后 0.x 只增不改。已知缺口四项与路线图三步同日公开,后者指向融合算子、整模型与计算通信融合。来源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0、https://pypi.org/project/tileops/
-
昇腾 910 A2/A3 的 NPUIR 后端整条并入 mlir-ascend 主干,验收口径首次写成数字(10-28 合入):一笔 55 文件 +21490/-1,Python 侧是 tilelang/ascend910 包(backend、adapter、codegen、compiler、execution_backend、jit、pipeline、runtime、target 加 language 下的 kernel、copy、gemm、npuir 方言面),C++ 侧是 src/ascend910/target 的开发者与专家两套发射器、算子层与 LowerNpuirBlock、NpuLoopVectorize、RewriteFDivToMul 三个 pass,运行时补 npu_launcher.h 与 npu_utils,target 别名接受 ascend910、ascend-910、a2、a3、npuir。作者给出的验收是 pytest testing/npuir 504 项全通过、examples 72 例中 71 例通过且 0 超时 0 报错,唯一失败的一例在旧后端同机同样失败,属环境与容差问题。来源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
-
主仓在审的 TIRx 重构把内置 TVM 换成外部依赖,规模从 1043 文件扩到 1194 文件(10-25 起,全周更新):方案是删掉 3rdparty/tvm 及其子模块条目,把前端、原生张量与作用域 IR、下降 pass、后端流水线与执行适配整体迁到独立构建发布的 TIRx,并对未打补丁的上游源码钉版本(apache/tvm-tirx 9af5544d,依赖 apache-tvm-tirx==0.1.0a2 与 apache-tvm-ffi==0.1.15.post0),原先带 +tilelang.1 后缀的集成补丁不再需要。周末到周中该 PR 三度更新,规模由 1043 文件 +71428/-104176 增至 1194 文件 +106648/-104497,是全周队列里体量最大的一笔,也决定后续 TVM 侧修复如何进入主线。至本周五仍在审,未合入。来源:https://github.com/tile-ai/tilelang/pull/3666
-
主仓一周的正确性硬化集中在「编译通过、运行不报错、结果错」这一族(10-26 至 10-30):典型条目是输出分配流(#3649,三个适配器接受调用方的 stream 却在当前流上分配 out_idx 输出,缓存分配器按分配流排序复用,H100 上五次跑三次输出损坏,修后归零)、全 warp 归约改走 redux.sync(#3692,四个 32 lane 的 FP32 取大从 20 条 shuffle 变 4 条硬件归约)、切片填充与负索引的静态边界(#3352 的 T.fill 在长度 8 的缓冲上写 As[6:12] 曾被放行,实际写到 6 到 11;#3347 的 A[-9] 曾静默回绕成 -1)、T.ws(-1) 区域在零个线程上运行(#3354,存全部丢弃且无诊断)、ROCm 七个默认数学函数不再降到快速内建(#3621,MI300X 上 sin/cos 误差从 7.8e-04 降到 6.6e-08)。来源:https://github.com/tile-ai/tilelang/pull/3649、https://github.com/tile-ai/tilelang/pull/3692、https://github.com/tile-ai/tilelang/pull/3621
-
昇腾侧的算子清单按上游模型结构逐个铺开(10-26 至 10-30):一周内落地的包括对齐 Kimi K3 Stable LatentMoE 门控激活的 situ_glu(#1891)、GLM-5.3 稀疏注意力 indexer 的 kpool 压缩算子(#1910,单 kernel 纯 Vector 路线,E4M3 编码用纯整数位运算实现、不依赖 FP8 类型也不做近似)、分块门控线性注意力的 Chunk Scaled Dot KKT 前向(#1803,静态四块展开、零 GM 工作区,L0 与 L1 各三例匹配率 1.0000)、W4A4 的 int4×int4 GEMM 结果全错根因与修法(#1731,半字节打包类型被当成 1 字节,所有拷贝 stride 与 repeat 次数 2 倍偏差,症状为 99.6% 不匹配),以及主仓 DeepGEMM 示例的 FP4×FP4 打包输入(#3589)。来源:https://github.com/tile-ai/tilelang-ascend/pull/1910、https://github.com/tile-ai/tilelang-ascend/pull/1803、https://github.com/tile-ai/tilelang-ascend/pull/1731
-
海光一周四笔主干动作,从「能跑」转向「按硬件结构摆」(10-26 至 10-30):寄存器流水线加 warp 发散合入(#18,28 文件 +5312/-64,跨两周多收敛,作者给出的结论是 example_gemm 达到与 rocblas 相当的性能)、原生 B16 存储与动态 WASP GEMM 合入(#28,6 文件 +340/-230)、scale 缓冲拆成独立双 bank 并与 LDS bank 分开(#29,12 文件 +1389/-314,本窗口软件侧最大的一笔)、普通拷贝路径的 LDS swizzle 特性分支进在审(#30,18 文件 +2000/-84,含五个新测试)。来源:https://github.com/tile-ai/tilelang-hygon/pull/18、https://github.com/tile-ai/tilelang-hygon/pull/29
-
发布镜像把算子库与 TileFoundry 首次打进同一镜像(10-29):运行器镜像在开发镜像之上安装 TileOPs 与调度器的 PyPI 发布版,标签为 tileops0.1.0 加 tilefoundry0.0.3 再加 cuda13.2、torch2.13 与 tilelang0.1.12 三段环境标识;附加脚本拉取镜像后在 GPU 上跑栈检查、基线检查、打包测试与可选基准,因为轮子里不带测试与基准,故都从镜像自身版本对应的标签取。同时移除 sgl-kernel 基线。来源:https://github.com/tile-ai/TileOPs/pull/2793
-
生态出现整建制的分支与社区仓,TileLang 被当作自有硬件栈的底座(10-26 至 10-30):T-HEAD 名下的 tilelang-for-sail 分支仓公开 13 个分支,除版本镜像外可见 PPU 矩阵共享布局修复、PPU language 包与运行时编译特性;CANN 社区的 ops-tilelang 算子仓 10-20 建仓后持续推送并已用七个算子串出 DeepSeek V4.1 Flash 的模型片段;高通 Hexagon 后端单日 15 笔,把 HVX 向量化、VTCM 竞技场、Crouton 布局与 DMA 内建一路铺到映射与归约;另有 RISC-V 社区枢纽仓、PTO-ISA 题库与基准双仓、tilelang-knowledge 知识仓同日推进。来源:https://github.com/t-head/tilelang-for-sail、https://github.com/Atlas-Ascend-Fork/ops-tilelang
二、核心项目进展(主仓 tilelang)
版本与发布
- 主仓最新标签与发布仍为 v0.1.15(09-30 发布),全周无新标签、无新发布页条目,至本周五已满 30 天。VERSION 文件核查仍为 0.1.15。
- TileOPs v0.1.0 是本周组织内唯一的新正式版本(10-27),也是 09-30 以来继 TileFoundry 0.0.3 之后的第二个新标签。TileFoundry 最新停在 0.0.4rc1(10-25),尚未转正。
- 适配线标签全周未动:tilelang-ascend v0.1.2.000(09-09)、tilelang-metax v0.1.15(10-16)、tilelang-musa v0.1.15+musa.1(09-30)、tilelang-mlir-ascend 停在 v0.1.2.020 系列。
- 合入量与节奏:周一那期为 72 小时窗口(10-23 07:00 至 10-26 07:00),合入 78 笔(TileOPs 51、tilelang 10、TileFoundry 10、文档站 4、昇腾仓 1、海光仓 1、Sunrise 1),与上一期周报(10-19 至 10-23)在 10-23 一天上有重叠,故不与后四期相加。周二至周五四个 24 小时窗口合计合入 86 笔:tilelang 47、TileOPs 26、tilelang-ascend 5、TileFoundry 4、tilelang-hygon 3、tilelang-mlir-ascend 1;单日依次为 30、20、24、12 笔,呈前高后低的收敛形态。在审更新全周维持在 45 至 76 笔之间,新开 issue 分别为 19、11、7、6、4 条。
编译器、语言层与执行路径
- 全 warp 归约改走硬件 redux(#3692,10-29 合入,4 文件 +337/-35):此前即使整组就是一个对齐的满 warp,普通跨线程归约也照样展开成 shuffle 树。主分支 837486a 上四个独立的 32 lane FP32 取大在 sm_103a 配 CUDA 13.1 编译出 20 条 shuffle、零条 redux,改后是四条 redux.sync.max.f32、零条 shuffle。做法是复用标量与批量全归约已有的 warp_reduce 助手,调用方只判线程几何。整数归约在 SM80 及以上走硬件 redux;FP32 的取大取小只在 sm_100a、sm_100f 与 sm_103a、sm_103f 上可用;窄浮点提升到 FP32 后再进助手,FP64 不降精度仍在 shuffle 路径。作者明确不声称端到端加速,只给指令级对照。
- autotuner 的独立 JIT 配置改为并发编译(#3476,10-29 合入,2 文件 +54/-46):原实现整段编译期持一把锁,连只用默认编译选项的配置也串行,前端、下沉与设备编译全部排队。改法是给每份覆盖配置用 dataclasses 复制出独立 JITImpl,默认选项仍走原有惰性路径以保留内核缓存与对象身份,只对共享第一阶段键的候选串行,调试输出的写盘仍串行。跨后端验收数据:昇腾 1797 项通过(Python 3.12.14、torch_npu 2.10.0、CANN 9.3.0,AscendC 示例 50 例全过、PTO 示例 44 过 6 跳)、ROCm 7.2 在 gfx942 上 3220 项通过、Metal 95 项通过、CUDA 在 H100 PCIe cu130 上 5500 项通过、CuTeDSL 75 过 17 跳。
- 昇腾缓冲区视图的布局与指针单位不再跨别名传播(#3623,10-29 合入,23 文件 +1139/-438):由 T.view 或 T.reshape 造出的多个缓冲可共用同一段存储而要求不同布局,此前统一按共享的数据变量解析布局,可能选中错误的 Cube 映射或 DMA 地址。本笔把布局附着到单个 Buffer 视图上,贯穿布局推断、分形存储归一化、多缓冲物化与拷贝和矩阵乘的下沉;裸指针偏移按各自元素单位重映射。同一带版本分配上,Cube 视图现在可用不同 MajorK 与 MajorMN 布局,条件是逻辑步长稠密、基址偏移为零、操作数类型兼容且填充后存储大小相等。
- GM 到 L1 的拷贝与越界填充合并成一个调度任务(#3543,10-29 合入):两者共同准备同一块 L1 数据,此前填充可能被拆成独立调度任务,同步插入就会多一道 MTE2 管道栅栏。改后消费者等待的是包含填充在内的完整数据,FP8 MX 的 K 尾填充同样并入并保留任务内条件填充,原来的 MX 延迟填充队列被删除。47 项聚焦测试通过。
- ROCm 与执行上下文的一组修正:exp、exp10、log、log2、log10、sin、cos 七个默认函数在 HIP 上不再无论 fast_math 设置都降到快速内建,八个带下划线的快速算子新增最低下沉规则(#3621);自动目标改读当前 HIP 设备而非固定 0 号卡(#3594);NVRTC 与 CuTeDSL 按首个输入张量所在设备取当前流(#3359);profiler 未指定设备时改用当前设备,缓存冲刷不再打到 0 号卡(#3357)。
- 语言层与下降的小修:TMA gather/scatter 的 Python 整型行索引转为 int32,字面量写法在 Blackwell 上不再编译崩(#3365);运行期行程数的流水线 epilogue 直接展开,分页解码恢复可编译(该缺陷由下游 TileOPs#1852 报出)(#3356);VectorizeLoop 不再静默忽略循环步长(#3345);T.clamp 的 NaN 传播下降被回退、回到 min/max 语义(#3628);T.Parallel 的 coalesced_width 统一转 IntImm(#3374);TMA 描述符只记录已校验的名字(#3484);非 restrict 参数的 handle 后缀不再被裁掉(#3366)。
正确性与缺陷修复
- 流与缓存身份:输出在调用方流上分配(#3649)与 NVRTC、CuTeDSL 取流(#3359)之后,本窗口把自动调优缓存键纳入 compile_flags(#3338,同一内核换编译选项不再复用旧二进制),昇腾侧把 Bisheng 构建身份纳入二进制缓存键(#3691)。两条独立路线在同一周做同一件事:缓存若只按源码与形状做键,换编译器就会误命中。
- 同步与内存规划:ThreadSync 把 address_of 记为写,不透明调用后的栅栏补上(#3346,此前 T.address_of(B[0]) 形式在 extern 之后没有栅栏,而 T.access_ptr(B, “w”) 形式有);T.ws(-1) 改为在所有线程上运行(#3354);跨存储域命名屏障重号(#3643)与 sm_100 的 TMEM 释放缺 CTA 屏障(#3645 / #3663 / #3664)在周初处理。周三浮出三笔 ThreadSync 错位栅栏缺陷(#3680):栅栏被放进发散分支、分支里的栅栏错误清掉外层挂起依赖、按线程计数循环的栅栏位置,配 #3682 与 #3681 两笔在审修复。
- 静态边界与类型:切片填充按区域末端而非长度校验、负索引在回绕后仍为负时拒绝(#3352 / #3347);SIMT 原子区域的形状校验补上(#3465);稀疏 MMA 的 K 尾块不再被静默放过(#3367);atomic_addx4 的目标 dtype 收窄到 float*(#3369);profiler 输出索引统一校验(#3486);标量读的索引系数不再混进 lane 数(#3453,行宽 32768 以上的编译崩溃关闭);批量 all-reduce 按标量分支解析线程区间(#3341,batch>1 时 2048 个输出全错);int64 原子 max/min 与 AtomicStore 的 int64、bf16、fp16 崩溃(#3307);归约的 FP8/FP4 非有限常量写出非法字面量(#3189)。周四把六个假定冷编译的测试改走 without_kernel_cache(#3695 / #3696)。
- 新开缺陷的指向:周四开出的三条都落在静态检查边界——带填充的共享缓冲走批量拷贝报形状不符(#3704,末维非 16 对齐的全局张量拷进带填充共享缓冲,BK 取 cdiv 乘 16,下沉报 112 对 100)、负缩放迭代图误编译反向索引(#3705,复现环境是 v0.1.15 与 T-HEAD 分支仓源码构建,属上游共有路径)、视图位宽按 int32 相乘溢出(#3707,Qwen3.8 的 lm_head 权重 248320 乘 2560 字节 NVFP4 共 5,085,593,600 位,回绕成 790,626,304 后抛字面量超限)。同批还有 sm_120 上 13 个示例失败、IfStmtBinding 指数级编译、without_kernel_cache 不强制重编(#3701 / #3700 / #3699),以及 CuTeDSL 的 CUDA graph 捕获为空图(#3693)。周五三条修复进在审:位宽改 int64、CPU 标量 GEMM 只清目标区域、关缓存时绕过两层进程内缓存(#3708 / #3709 / #3710)。
运行时、CI 与工具链
- CI 把 ROCm 与 Metal 任务锁回兼容的 PyTorch wheel(#3706,10-29 合入):ROCm 的 CI 通过索引变量暴露匹配的 PyTorch ROCm 源,但工作流级的「不安全最优匹配」策略让更晚的 PyPI 候选胜出,解析到 PyTorch 2.15.0(CUDA 13.2)而不是 2.14.1 加 ROCm 7.2 构建;同一批 2.15.0 的铺开也让 macOS 26.6.2 的 ARM64 运行器拒绝 Apple 半虚拟 GPU。改法三条:按配置的稳定 ROCm 版本设置后端变量、夜间安装改用首个索引、新增测试前快速失败的校验步(确认是 HIP 构建、非 CUDA 构建、且能访问运行器 GPU)。
- 昇腾 CI 把 AscendC 与 PTO 拆成两个矩阵作业,可分别重跑(#3687);CI 镜像升到 cann9.3.0 周更版(#3630)。
- 贡献流程治理四连(10-26,全围绕评审成本):CONTRIBUTING.md 写明 PR 描述应有问题、方案与取舍,证据贴在其支持的论断旁(#3671);纯文档 PR 跳过 CUDA、ROCm、Metal、Ascend 与 CuTeDSL 的构建与测试,白名单刻意收窄(#3672);提交技能加确认步,先把拟用标题给贡献者确认(#3673);关掉评审机器人的自动摘要与文档字符串覆盖率检查,要求把发现锚定在当前 diff 与实际调用路径上(#3674)。
- 语言服务上线实验性 CUDA 调试(10-26):tilelang-lsp 推入首个调试提交,新增 tlsp.debug 启动器做 Python 行号映射、加 -G 与 -lineinfo 编译、隔离缓存并抓取生成源,把 Nsight 的 cuda-gdb 接进 VS Code,支持 Python 断点、启动前检查与生成 CUDA 源的浏览。这是该仓自 8 月以来第一次实质提交。
- 文档站全周持续自动同步:10-26 八笔、10-27 六笔、10-28 两笔,另有 TileOPs 文档站与本地化仓的常规更新。
评审队列观察
- 体量最大的一笔仍是 TIRx 重构(#3666,1194 文件 +106648/-104497),对固定的上游基线打零补丁,是队列里唯一会决定「后续 TVM 侧修复以何种节奏进入主线」的改动。
- 昇腾仓的 a2a3 调度器迁移占 100 文件 +18453/-666(#1951),是专用仓里最大的一笔在审,说明下一阶段变动量在调度层而不在算子层。
- PTO 线的在审队列一周内持续推进:T.Parallel 向量化、AutoSimVF 通道、批量三维 ND2NZ、写计数多缓冲(#3512 / #3320 / #3689 / #3629),PTO 目标默认走 tvm_ffi 执行后端(#3516),另有三笔视图与广播修正(#1962 / #1960 / #1961)与按位异或的无符号字节工作区视图修正(#1959)。
- 在审更新数全周在 45 至 76 笔之间,周五回升到 54 笔,与合入量的回落形成对照:本周是「清账周」而非「开单周」。
三、多后端适配(昇腾 / 海光 / ROCm / Hexagon)
昇腾(Ascend):后端并入主干,算子按模型铺开
- 后端层面:910 A2/A3 的 NPUIR 路径并入 mlir-ascend 主干(#207,55 文件 +21490/-1),模式开关 TILELANG_ASCEND_MODE 取 expert 或 developer、默认 expert;第二笔提交逐条对齐旧后端行为差异,重点是 bf16 合法化(HIVM 不接受 bf16 操作数,改为 fp32 计算后显式转换)与动态 extent 的拷贝下沉。
- 调度与定址:主仓新增 T.annotate_buffer_addresses,允许在调度后的昇腾流水线里对 UB、L1、L0A、L0B、L0C 与 SS 各池内的分配做精确字节定址,放入定址的缓冲须先用 T.annotate_buffer_versions 声明固定份数,校验覆盖缺少或不支持的份数、重叠、未对齐、动态尺寸与容量溢出(#3607,14 文件 +982/-26)。PerCoreTask 候选体允许内部 if/else,条件 L1 填充不再被整体拒绝(#3686);手写流水线保留配置寄存器相位(#3624)。
- 算子供给:一周内合入 situ_glu(#1891)、GLM-5.3 的 kpool 压缩(#1910)、Chunk Scaled Dot KKT 前向(#1803)、W4A4 int4×int4 的根因修复(#1731)、depthwise_conv_2d 示例重提(#1920,改用清单化测试通道以避开并发批次);在审还有 GroupNorm 与 fla_chunk 两个示例(#1607 / #1953 / #1952)、NPUIR 视图与广播三笔(#1962 / #1960 / #1961)。这些算子的共同点是把量化与融合写进 kernel 本体,而不是留给宿主侧补救。
- 工程口径:新增算子性能剖析 skill,用 msprof op 输出定优化方向(#1931);安装器的 SHMEM 段双修,空壳 wheel 不再假成功(#1943 / #1946);超过两维的 GM 到 UB 拷贝按 DMA 行数折叠前导维(#1805);Cython 包装器每次调用重建分析器的问题(约 0.28 毫秒)先记录(#1945)后由在审的惰性构造修掉(#1948)。
- 回归状态:每日全量测试 10-25 报 2903 项全通,10-26 凌晨测试作业报失败(#1944),10-27 与 10-28 连续两日 2903 项全通过(#1950)。用例规模已足以当门禁,结论有效期以周计。
- 在审的兼容性改动:Carver 按所选 SDK 读 A3 资源预算,避免把 256×256 的 FP32 累加器放到只有 128 KiB 的 L0C 上(#1955);AI Core Exception Dump 的依赖降到 CANN 9.2.0 并补齐四个必填字段(#1956);文档改用 T.Tensor、弃用 T.Buffer DSL 接口(#1957)。昇腾仓的文档事实核查四条单条评论最多 46 轮(#1628 / #1655 / #1656 / #1659)。
海光(Hygon):一周四笔,按硬件结构摆放
- 寄存器流水线可与既有的共享内存流水线协同,对应的提交与屏障指令按流水级插入,warp 数超过 4 时启用 warp 发散(#18)。这一笔 10-08 开出、跨两周多收敛,合入时是先并主干再走分支的大合并,故行数远大于同批其他改动。
- 原生 B16 存储与动态 WASP GEMM(#28):保持 MLS 的 Let 参数发射顺序、把 GEMM 的 warp 划分传播到矩阵存储、拒绝打包的 B16 到 B32 输出别名,同时让 WASP 维度可动态并按轴推断对齐。
- scale 缓冲的行规划改为独立双 bank 分区、与 LDS bank 分开,同一套行映射同时交给分配、scale 拷贝与 MMAC 下沉三处,避免三处各自推一遍布局(#29,12 文件 +1389/-314)。这一笔接上 #28,海光在两天里完成三笔主干合入。
- 普通拷贝路径的 LDS 写相位与 swizzle 策略进在审,配 b8 与 fp8 两类数据宽度(#30,18 文件 +2000/-84,新增五个测试)。
来源:https://github.com/tile-ai/tilelang-hygon/pull/18、https://github.com/tile-ai/tilelang-hygon/pull/28、https://github.com/tile-ai/tilelang-hygon/pull/29、https://github.com/tile-ai/tilelang-hygon/pull/30
ROCm:pip 版 SDK 与浮点语义
- ROCm 10 的 pip 版 SDK 支持合入(#3549,2 文件 +106/-19):ROCm 10.x 的 rocm/pytorch 镜像把 ROCm 以 Python wheel 形式发布,既没有 /opt/rocm 目录树也没有 rocm-* 的 apt 包;get_rocm_arch() 此前硬编码 /opt/rocm 前缀,路径不存在时回落到 gfx900,于是 MI355X(gfx950)上 rocwmma 的静态断言直接失败,设 ROCM_PATH 也无效。修法是按调用方前缀、find_rocm_path()、torch 的 gcnArchName、gfx900 兜底的顺序解析,每次 rocminfo 探测加 10 秒超时,并新增基于 rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12 的 Dockerfile。作者在 MI355X 上实测架构探测从 gfx900 变为 gfx950。
- exp/sin/cos 的精度修正与自动目标、取流、profiler 三处上下文修正见第二章。ROCm 侧的 FP4 转换与 E8M0 类型、pip 版镜像的架构探测仍在审(#3665 / #3549 相关配套)。
沐曦(MetaX)/摩尔线程(MUSA):窗口内无实质提交
- tilelang-metax 组织视图有推送时间戳(10-26 15:08),但逐分支核对后无窗口内提交:默认分支 main 最新仍是 09-29 的版本号提升,dev 分支最新 10-16,0.1.14 与更早的版本镜像未动,事件流里也没有窗口内推送事件。最近一次正式发布仍是 10-16 的 0.1.15。
- tilelang-musa 全周无窗口内主分支提交,最近标签仍是 09-30 的 v0.1.15+musa.1。摩尔线程与沐曦在本周只出现在核查表里,不构成条目。
来源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa
第三方:Hexagon DSP 后端
- 10-27 单日 15 笔,把 HVX 向量化、VTCM 竞技场、Crouton 布局与 DMA 内建一路铺到 Crouton 上的映射与归约。至 10-29 主分支最新提交仍停在 08-24 的 CI 依赖升级,仓库推送时间更新但主分支无新代码,故只作为社区投入强度的观察,不计入上游进展。
- 另有 PPU 版零改动包装器(运行时拦截 import 协议与 site 机制,不修改上游源码也不重新打包,无 PPU 对应物的 CUDA 特性显式报错而非静默降级),是后端迁移的第三种形态。
来源:https://github.com/xwhzz/tilelang-hexagon、https://github.com/flytiger-eco/tilelang-wrapper-for-sail
四、生态与采用方
TileOPs:0.1.0 之后的第二周
- 发布之后的合入节奏明显放缓:周二 14 笔、周三 4 笔、周四 7 笔、周五 1 笔,重心从「把接口冻住」转到性能与正确性的收尾。
- 性能侧的四处硬数字:选择器不再贴着 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%(#2784);非 NT 布局改走转置持久内核 coop2_swap,2112×7168×4096 TN 从 195.3 微秒降到 175.3 微秒,同机另一块 H200 从 174.3 降到 158.5 微秒,均压过 QuACK 的 187.9 与 167.0 微秒(#2789);cumsum 长行改两段式种子扫描,(8,65536) 从 12.75 微秒降到 5.18 微秒并修掉 65535 tile 以上的发射失败(#2790);extrema 归约改用 PTX 的 max.NaN / min.NaN,amax (2048,4096) bf16 从 8.352 微秒回到 6.976 微秒,快于 FlagGems(#2792)。
- 背靠背口径被单独立项:周四新开议题指出单次调用下本实现领先 QuACK,但背靠背、功率封顶的训练场景里十条大行有六条落后(#2791);随后在审的 GEMM 改动按功耗封顶重算栅格分组与常驻块数,DeepSeek V3 预填充的 q-b 形状上内存读从 168 MB 降到 94 MB(#2794)。
- 正确性与能力补齐:D128 稠密前向拆 key,prefill-rope 从 31.04 微秒降到 17.09(#2777);D128 持久化前向按线程重写并写出 LSE(#2774);融合加法的 LayerNorm 行跨 CTA 拆分,精度缺陷另开新单(#2768 / #2773);偏移存储三类修复与 SM90 分组 GEMM、输出 ABI(#2761 / #2765 / #2762)。
- 发布形态:发布镜像合入(#2793),算子库与 TileFoundry 首次打进同一镜像;已知缺口中的 FusedAddLayerNormFwdOp 精度问题在周内仍未被修掉,路线图第一步的融合算子尚无合入。
来源:https://github.com/tile-ai/TileOPs/pull/2784、https://github.com/tile-ai/TileOPs/pull/2789、https://github.com/tile-ai/TileOPs/pull/2790、https://github.com/tile-ai/TileOPs/pull/2792
TileFoundry:从三包拆分到静态契约
- 周初拆成 core / nvidia / apple 三个发行包,测试输入只装载一次(#268 / #269);sm_89 补上 E4M3 的 warp MMA,并把能力校验从「只查 MMA」扩到每条指令,bulk copy、TMA copy 与 mbarrier 形式此前能通过验证、到 codegen 或 ptxas 才失败(#261 / #262 / #264)。
- 周中修 cp.async 目标布局寻址与 MLA 反向布局(#263 / #260 / #257),以及整窗口按操作数本身下发(#273,sm_89 的 GEMM 恢复可编译)与动态共享内存按内核实际用量定尺寸(#271,作者未采纳「按目标容量给满」的建议,理由是静态与动态共享内存同计一个上限、每块申请多少决定一个 SM 能放几块)。
- 周四合入一笔破坏性变更(#279):把 24 个 AMX 运行时算子族描述成 39 个静态注册的 TIR 契约,带参数域、操作数与布局模式、副作用集合,以及 M1 到 M4 的目标能力;新增 Python 评估器与原始 C 模拟器作为金标准,补 59 个代表性原子用例覆盖全部 39 个静态算子,高阶与中阶中间表示用例跑真实的 2、4、8、12 线程域并验证候选的接受与拒绝。验证为 46 项聚焦测试通过。
- 新开缺陷侧记一笔:分析侧把共享内存峰值算重、且不传到 TIR(#272),与主仓 ThreadSync 那一族同源。
采用方与社区分支
- 两个已知采用方全周静默:DeepSeek 的 TileKernels(2018 星)与通义侧的 FlashQLA(742 星)主分支都停在 10-14,窗口内无新推送。本周生态信号主要来自社区自建仓,而非这两个内部使用方。
- T-HEAD 的 tilelang-for-sail:09-22 建仓,是上游主仓的分支,公开 13 个分支,除主分支与 v0.1.11 到 v0.1.14 的版本镜像外,可见 PPU 矩阵共享布局修复、PPU language 包、运行时编译特性与若干缺陷修复分支。主分支最近提交停在 10-22,仓库推送时间更新到 10-29 19:27;仓内无说明文档也无发布记录,故只能作为观察线索。值得注意的是周五第 #3705 号缺陷正是用该仓源码构建复现的——分支用户比上游更早撞上真实构建问题。
- CANN 社区的 ops-tilelang:10-20 建仓,定位面向昇腾 NPU 的社区高性能算子仓,统一管理 kernel 实现、自动化用例、标准工作负载与性能基线,README 明确只支持 Ascend 950 及以上、要求 CANN 9.3.0 以上并用社区 weekly 版。已用七个算子串出 DeepSeek V4.1 Flash 的「Attention 输出加 FFN 前处理」片段,HIDDEN=5120、MHC=4,B 从 1 到 256 实测 SuperKernel 收益 1.17% 到 2.75%。它与官方 tilelang-ascend 的关系尚不明确,只作生态观察。
- PTO-ISA 的双配套仓:题库仓把 PTO 的广播加载迁移到 PTOAS 0.70,此前一周在做量化缩放符号澄清与题面改写;基准仓加 SV2 分块缩放与逐列分析两笔提交,依据是库内绝对最大值指令的变体计数 31.39,附运行时指令吞吐与屋顶线分析,明确不是统一缓冲口径。
- 分发版:Sunrise 的 0.1.15+sunrise.1.2.0 发布 PR 于周初合入(#9),是组织内第二例下游分发形态;主仓的 TIRx 迁移若落地,这类分发版需要跟着改依赖声明,是下周值得跟踪的连带影响。
- RISC-V 方向的社区枢纽仓:TongmingLAIC/tilelang-riscv 于 10-23 建仓、10-26 推送,定位是 TileLang 在 RISC-V 上的社区枢纽,聚合编译器后端、AI 算子、开发工具与文档。目前仍是空骨架,未见可运行算子或后端落地。
来源:https://github.com/t-head/tilelang-for-sail、https://github.com/Atlas-Ascend-Fork/ops-tilelang、https://github.com/PTO-ISA/tilelang-puzzles-ascend、https://github.com/TongmingLAIC/tilelang-riscv
五、社区、教程与活动
- 知识资产成为本周社区侧的新品类:tilelang-knowledge 于 10-29 建仓,当天四笔提交构成一套面向智能体的知识底座——仓库结构、33 个算子设计案例(每条带固定出处)、20 张带出处的编程知识卡、7 条有证据支撑的 GPU 优化模式。仓库自述用途是让智能体面对一个算子时能判断哪些经验适用、应当如何改代码、如何验证、失败后往哪里找。这与 TileFoundry 的算子契约、昇腾仓的剖析 skill 是同一类尝试。
- 教程与练习方向持续投入:tilelang-puzzles-ascend 做 PTO 对 Ascend 的量化 kernel 教程;中文 TileLang 算子面试练习站 9 模块 18 题;图解课配套仓自述用原创场景加动画讲透每个知识点并随视频发布;分布式布局实验仓同日新建。
- 文档与本地化:tilelang.github.io 全周随主仓合入自动同步(10-26 八笔、10-27 六笔、10-28 两笔),第三方本地化仓 tilelang-docs-l10n 常规更新,中文 RISC-V 社区枢纽仓完成首次内容提交,说明文档同步链路正常运转。
- 讨论侧:TileScale 用户询问所用 TileLang 版本能否升级(#72)。该仓最后推送停在 08-25,文档与示例仍指向较早版本;这类提问通常出现在上游接口进入快速改名期之后——本周主仓在审的 TIRx 替换与 TileOPs 的 26 个算子类改名,都会放大下游跟版成本。
- 新闻侧全周零实质命中:Google News 的中英文六组查询(覆盖精确词加两日、中文加昇腾、中文加算子、组织名加两日、英文加 GPU、中文加编译器加七日)在五个采集窗口里只返回一条临床人工智能的转载稿,与算子语言无关;Hacker News 的相关命中全部落在窗口之外,最近一条是 10-02 的昇腾编程栈转载讨论;arXiv 以 TileLang 检索无新论文。日报主体因此继续以仓库与生态仓为准。产业侧仅 10-26 有一篇昇腾与 DeepSeek 生态差距收窄的报道,属二次评述,不含新的技术数据。
六、趋势观察
发布工程从「卡在最后一格」走到「镜像化交付」。 周一那期日报还在写 TileOPs 首发「文档就绪、制品未上传」,PyPI 上最新仍是 2025-10-09 的 0.0.1.dev2;周二 07:45 制品上传、13:15 发布页挂出,周四发布镜像就把算子库、TileFoundry 与三段环境标识(cuda13.2、torch2.13、tilelang0.1.12)锁在一起。一周之内,这条线从「能不能发」变成「怎么组合交付」。对一个 200 余算子的库来说,可安装性比再优化一个百分点更能决定外部采用速度,本周的答案是把安装路径也做成制品。
缺陷类型高度集中在「编译通过、运行不报错、结果错」,且与硬件或上下文强相关。 输出在错误的流上分配(只在非默认流且有排队工作时损坏)、T.ws(-1) 在零个线程上运行、切片填充按长度而非末端校验、负索引回绕后仍为负、address_of 被当成读——五条都在编译期无警告、运行期无异常。这类缺陷的代价远高于编译错误,因为它们污染的是「测试全过」这个前提本身;周三浮出的六个假定冷编译的测试与周四的 without_kernel_cache 语义漏洞,恰好说明测试基础设施也在同一片风险区。本周的应对是把边界校验前移到编译期,方向正确,但#3704、#3705、#3707 三条新开缺陷说明这条边界还远未画完。
昇腾从「搬后端」转到「补算子」,并第一次给出可复算的验收口径。 周三是 NPUIR 并入主干(504 项自测、71/72 示例),周四到周五的五笔合入里四笔是算子与示例,且都按上游模型的需求排列:GLM-5.3 的 indexer 压缩、分块门控注意力的内积、四比特量化矩阵乘、量化视觉模型的深度卷积。共同点是把量化与融合写进 kernel 本体(E4M3 纯整数编码、四比特打包约定、双 kernel 加 1 乘 1 快路径),而不是留给宿主侧补救;同时把剖析 skill 与清单化测试通道立成流程。主要矛盾正从「能不能编出来」变成「算子的数值与性能能不能被稳定验收」,而 a2a3 调度器迁移的 100 文件预告了下一阶段的变动在调度层。
执行路径做减法,后端清单做加法。 主仓这一周把全 warp 归约收进 redux.sync、把 autotuner 的编译锁拆细、把 GM 到 L1 的拷贝与填充并成一个任务、把 buffer 视图的布局从共享变量收回到单视图;与此同时后端清单继续扩张——ROCm 补 pip 版 SDK 与 FP4、Hexagon 单日 15 笔、PPU 出现零改动包装器、Intel AMX 在 TileFoundry 里变成 39 个静态契约。两条方向并不矛盾:上游收敛执行路径,下游各自打包分发,这与 TIRx 把 TVM 变成外部依赖是同一条思路。
「fork 化」是本周最硬的采用信号,也带来最实的同步成本。 T-HEAD 把 TileLang 当自有硬件栈底座做 PPU 分支、CANN 社区自建算子仓、PTO-ISA 把题库与基准都做成配套仓、RISC-V 出现社区枢纽仓——这些都不在 tile-ai 组织里,却都依赖语言层的稳定性,这正是 0.1.0 冻结公开 API 并做成 CI 门禁的意义。代价同样清楚:周五第 #3705 号缺陷由分支用户的真实构建先撞上,上游才补修;分支越多,这类「下游先发现、上游后修复」的路径就越长,主仓需要开始回答分支与上游如何同步的长期问题。
口径问题开始被当作产品问题处理。 TileOPs 周四把「单次调用领先、背靠背落后」单独立项(#2791),周五的 GEMM 改动按功耗封顶重算栅格分组并把 DRAM 读从 168 MB 降到 94 MB(#2794);选择器也从贴 manifest 拟合改为按忙时校准。相比「某行快了多少微秒」,这类改动承认了训练与推理的真实调用形态与基准不同,并把差异写进调度策略。对一门要给大模型写内核的语言,这比再堆算子数量更能决定它是否被长期采用。
附录:素材与核查说明
信源核查表
| 信源 | 核查结果 |
|---|---|
| GitHub tile-ai 组织逐仓(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文档站 / 夜间镜像仓) | 窗口内有推送:tilelang、TileOPs、TileFoundry、tilelang-ascend、tilelang-mlir-ascend、tilelang-hygon、tilelang.github.io、TileOPs.github.io、TileOPs-nightly;tilelang-metax、tilelang-musa、TileRT、tilescale、DeepStack、tilelang-puzzles 无窗口内主分支提交 |
| 合入真值检索(is:pr is:merged) | 周一 72 小时窗口 78 笔;周二至周五四个 24 小时窗口合计 86 笔(tilelang 47、TileOPs 26、tilelang-ascend 5、TileFoundry 4、tilelang-hygon 3、tilelang-mlir-ascend 1),单日依次 30、20、24、12 笔 |
| 在审真值检索(is:pr is:open updated) | 五个窗口分别为 76、49、45、48、54 笔,含 TIRx 重构、a2a3 调度器迁移、昇腾文档核查批次与 TileOPs 的 GEMM 调度改动 |
| 新增 issue 核对 | 五个窗口分别为 19、11、7、6、4 条,其中每日包含昇腾仓测试机器人的失败通知 |
| 主分支提交级核查 | 各窗口合入数与主分支提交数一致;tilelang 单日 10、14、13、15、5 笔 |
| 发布页与标签 | tilelang 最新仍为 v0.1.15(09-30 发布),全周无新版本;TileOPs 最新为 v0.1.0(10-27);TileFoundry 最新为 0.0.4rc1(10-25);tilelang-ascend 最新为 v0.1.2.000(09-09);tilelang-metax 为 v0.1.15(10-16) |
| 第三方与生态仓 | 昇腾社区算子仓、PTO-ISA 双仓、开发工具仓、知识仓、图解课仓、分布式布局仓、DCU 融合算子仓、模型移植仓、Hexagon 后端仓、T-HEAD 分支仓、RISC-V 社区枢纽仓均有窗口内推送;Hexagon 主分支无新提交,仅仓库推送时间更新 |
| 采用方 | TileKernels 与 FlashQLA 主分支停在 10-14,窗口内无推送 |
| Google News(中英六组) | 仅一条中文加昇腾的临床人工智能转载稿与一条产业侧评述,均不含新的技术数据,不收录为条目 |
| Hacker News | 近四日相关命中全部落在窗口之外 |
| arXiv | 以 TileLang 检索无新论文 |
主要引用
- [1] TileOPs v0.1.0 发布页 — https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0
- [2] TileOPs PyPI 制品 — https://pypi.org/project/tileops/
- [3] 昇腾 910 NPUIR 后端并入主干(#207)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
- [4] TIRx 重构(#3666)— https://github.com/tile-ai/tilelang/pull/3666
- [5] 输出分配流修复(#3649)— https://github.com/tile-ai/tilelang/pull/3649
- [6] 全 warp 归约改走 redux.sync(#3692)— https://github.com/tile-ai/tilelang/pull/3692
- [7] autotuner 并发编译(#3476)— https://github.com/tile-ai/tilelang/pull/3476
- [8] ROCm 数学函数精度修正(#3621)— https://github.com/tile-ai/tilelang/pull/3621
- [9] 昇腾缓冲精确定址接口(#3607)— https://github.com/tile-ai/tilelang/pull/3607
- [10] 昇腾 kpool 压缩算子(#1910)— https://github.com/tile-ai/tilelang-ascend/pull/1910
- [11] W4A4 int4×int4 根因修复(#1731)— https://github.com/tile-ai/tilelang-ascend/pull/1731
- [12] 海光 scale 缓冲双 bank(#29)— https://github.com/tile-ai/tilelang-hygon/pull/29
- [13] 海光寄存器流水线(#18)— https://github.com/tile-ai/tilelang-hygon/pull/18
- [14] ROCm 10 pip 版 SDK 支持(#3549)— https://github.com/tile-ai/tilelang/pull/3549
- [15] TileFoundry AMX 静态契约(#279)— https://github.com/tile-ai/TileFoundry/pull/279
- [16] TileOPs 发布镜像(#2793)— https://github.com/tile-ai/TileOPs/pull/2793
- [17] TileOPs 选择器按忙时校准(#2784)— https://github.com/tile-ai/TileOPs/pull/2784
- [18] T-HEAD 分支仓 — https://github.com/t-head/tilelang-for-sail
- [19] CANN 社区算子仓 — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [20] Hexagon 后端仓 — https://github.com/xwhzz/tilelang-hexagon
- [21] RISC-V 社区枢纽仓 — https://github.com/TongmingLAIC/tilelang-riscv
- [22] tilelang-knowledge 知识仓 — https://github.com/qiuuuuu622/tilelang-knowledge
说明:本报告事实、版本号与日期均取自周期内五期日报原文及其对应的一手来源,未作推测性补充;周一那期日报的采集窗口为 72 小时(覆盖上周末),其合入数含 10-23 一天,与上一期周报存在重叠,故本报告未将其与后四期相加。周期内五个工作日日报齐全,无缺失。