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

信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文档站 / 夜间镜像仓)+ 合入与在审真值检索(合入 12 笔、在审更新 54 笔)+ 新增 issue 全量核对(4 条)+ 主分支逐笔提交核查(tilelang 5、tilelang-ascend 5、TileOPs 1、TileFoundry 1、文档站 5)+ 发布页与 tag 核查 + 第三方后端与生态仓核查(10 个仓库)+ Google News 中英文六组查询(走代理)+ Hacker News + arXiv。


本期索引

  • 今日重点:昇腾仓五笔合入补齐算法侧算子(GLM-5.3 kpool 压缩、Chunk Scaled Dot KKT、W4A4 GEMM 修复、depthwise 卷积示例、性能剖析 skill);主仓把全 warp 归约改走 redux.sync、把 autotuner 的独立配置改成并发编译;24 小时合入 12 笔、在审更新 54 笔、新开 issue 4 条
  • 一、核心项目进展
    • 1.1 主仓:全 warp 归约改走 redux.sync,四个 32 lane 的 FP32 取大从 20 条 shuffle 变 4 条硬件归约(#3692)(10-29)
    • 1.2 主仓:autotuner 的独立 JIT 配置改为并发编译,锁只保留在同键候选(#3476)(10-29)
    • 1.3 主仓:昇腾缓冲区视图的布局与指针单位不再跨别名传播(#3623)(10-29)
    • 1.4 主仓:GM 到 L1 的拷贝与越界填充合并成一个调度任务,去掉多余 MTE2 栅栏(#3543)(10-29)
    • 1.5 主仓:CI 把 ROCm 与 Metal 任务锁回兼容的 PyTorch wheel(#3706)(10-29)
    • 1.6 新开缺陷三条:带填充的共享缓冲走批量拷贝报形状不符、负缩放迭代图误编译反向索引、视图位宽在 int32 溢出(#3704 / #3705 / #3707)(10-29)
    • 1.7 三条修复在审:位宽改 int64、CPU 标量 GEMM 只清目标区域、关缓存时绕过两层进程内缓存(#3708 / #3709 / #3710)(10-29)
  • 二、多后端适配(昇腾 / 海光 / 摩尔线程 / 沐曦)
    • 2.1 昇腾:GLM-5.3 稀疏注意力 indexer 的 kpool 压缩算子合入,单 kernel 纯 Vector 路线(#1910)(10-29)
    • 2.2 昇腾:分块门控线性注意力的 Chunk Scaled Dot KKT 前向算子合入,静态四块展开、零 GM 工作区(#1803)(10-29)
    • 2.3 昇腾:W4A4 的 int4 乘 int4 GEMM 结果全错问题修复合入,根因是类型尺寸谎报与专用转置指令被静默关闭(#1731)(10-29)
    • 2.4 昇腾:depthwise_conv_2d 示例重提,改用清单化测试通道以避开并发批次(#1920)(10-29)
    • 2.5 昇腾:新增算子性能剖析 skill,用 msprof op 输出定优化方向(#1931)(10-29)
    • 2.6 昇腾在审:NPUIR 语言层三笔视图与广播修正(#1962 / #1960 / #1961)(10-29)
    • 2.7 昇腾在审:PTO 后端按位异或的无符号字节工作区视图修正(#1959)(10-29)
    • 2.8 海光:HCU 两笔在审有更新、社区出现 DCU 融合算子仓;摩尔线程与沐曦窗口内无推送(10-29)
  • 三、生态与采用方(TileOPs / TileFoundry / TileKernels / FlashQLA / 第三方)
    • 3.1 TileOPs:发布镜像合入,算子库与 TileFoundry 首次打进同一镜像(#2793)(10-29)
    • 3.2 TileOPs 在审:GEMM 按功耗封顶重算栅格分组与常驻块数,q-b 的 DRAM 读从 168 MB 降到 94 MB(#2794)(10-29)
    • 3.3 TileFoundry:AMX 运行时算子改成 39 个静态 TIR 契约,属破坏性变更(#279)(10-29)
    • 3.4 采用方:TileKernels 与 FlashQLA 主分支窗口内无推送(10-29)
    • 3.5 昇腾侧出现独立的社区算子仓 ops-tilelang,窗口内有推送(10-29)
    • 3.6 PTO-ISA 配套双仓活跃:题库仓迁移到 PTOAS 0.70,基准仓加 SV2 分块缩放与逐列分析(10-29)
    • 3.7 观察:T-HEAD 名下的 tilelang 分支仓窗口内有推送,分支名指向 PPU 的矩阵指令布局(10-29)
  • 四、社区、教程与活动
    • 4.1 文档站把内建归约的适用边界写进指令页,与 redux 合入同步(10-29)
    • 4.2 开发工具仓完成下沉跟踪抽取后的三笔收尾,与上游在审工具化改动同向(10-29)
    • 4.3 新仓 tilelang-knowledge 一天四笔:33 个算子设计案例、20 张带出处的知识卡、7 条优化模式(10-29)
    • 4.4 图解课配套仓与分布式布局实验仓同日新建(10-29)
    • 4.5 新闻侧:Google News 六组中英文查询、Hacker News 与 arXiv 均无窗口内实质命中(10-29)
  • 五、趋势观察
  • 附、素材与核查说明

今日重点

今天的合入量从昨天的 24 笔回落到 12 笔,但结构更集中:昇腾专用仓占五笔、主仓五笔、TileFoundry 与 TileOPs 各一笔。昇腾仓这一批全是算法侧算子,而且都是能对上游模型直接用的那一类。GLM-5.3-Flash 稀疏注意力 indexer 的 kpool 压缩算子把逐维 softmax 池化、BF16 边界、归一化 Hadamard-128 旋转与 E4M3 量化压进一个纯 Vector 的 kernel,每块覆盖 64 个 pool、两个向量子核各管 32 个,宿主侧零重排,直接按调用方的三维布局做跨步搬运;E4M3 编码用纯整数位运算实现(round-half-even、含次正规、饱和到 448),不依赖 FP8 类型也不做近似。Chunk Scaled Dot KKT 是分块门控线性注意力前向里的 A = gate(K K^T),静态四块展开、零 GM 工作区、零宿主算子,L0 与 L1 三例的匹配率都是 1.0。

主仓今天的主题是编译与发射层。全 warp 归约改用 redux.sync:主分支上四个独立的三十二线 FP32 取大在 CUDA 13.1、sm_103a 上展开成 20 条 shuffle,改后变成四条 redux.sync.max.f32 且不含 shuffle;整数归约在 SM80 以上走硬件归约,FP32 的取大取小只在 sm_100a、sm_100f、sm_103a、sm_103f 上可用,窄类型(含 CUTLASS 的 FP16、BF16)提升到 FP32,FP64 仍留 shuffle 路径。作者明确不声称端到端加速,只给了指令级对照。autotuner 的并发编译则是把整段持有的编译锁拆开:每份覆盖配置用 dataclasses 复制出独立的 JITImpl,只有共享第一阶段键的候选才串行,调试输出的写盘仍串行。作者给出的跨后端验收是昇腾 1797 项、ROCm 3220 项、Metal 95 项、CUDA 5500 项。

另有一处值得记的细节:今天新开的反向索引误编译缺陷(#3705)是拿 T-HEAD 名下的 tilelang 分支仓做运行时复现的,这说明已有整建制团队把这条上游代码当作自有硬件栈的底座在跟,而不只是当一个外部依赖。


一、核心项目进展

1.1 主仓:全 warp 归约改走 redux.sync,四个 32 lane 的 FP32 取大从 20 条 shuffle 变 4 条硬件归约(#3692)

合于 10-29 15:08,4 文件 +337/-35。此前普通的跨线程归约(如 T.reduce_max)即使整组就是一个对齐的满 warp,也照样展开成 shuffle 树:主分支 837486a 上四个独立的 32 lane FP32 取大在 sm_103a 配 CUDA 13.1 编译出 20 条 shuffle、零条 redux,改后是四条 redux.sync.max.f32、零条 shuffle。做法是复用标量与批量全归约里已有的 warp_reduce 助手:调用方只判线程几何(是否与 warp 对齐、单位步长、满 warp),架构、数据类型与算子分发交给助手。整数归约在 SM80 及以上走硬件 redux,更小的整数提升到 32 位;FP32 的取大取小只在 sm_100a、sm_100f 与 sm_103a、sm_103f 上可用;窄浮点提升到 FP32 后再进助手,FP64 不降精度、仍在 shuffle 路径。子 warp、跨步、打包归约以及不支持的算子与目标一律回退 shuffle,跨 warp 的共享内存交换与栅栏不变。顺带修掉一处编译失败:自定义无符号整数 reducer 此前因内联分发返回空类型而编译不过,现在落到 shuffle 回退。PTX 回归断言真实指令与 shuffle 缺席,运行时覆盖有符号无符号边界值、批量不清零、NVRTC 执行、非有限值、有符号零与 FP64 精确结果。作者明确不声称端到端加速。既有的一处遗留问题也在描述里点明:批量归约在 batch=4 每线程两输出、batch=8 每线程四输出时存在 FP16 与 BF16 失败,主分支原模板同样复现,且失败发生在下沉阶段,与本笔无关。

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

1.2 主仓:autotuner 的独立 JIT 配置改为并发编译,锁只保留在同键候选(#3476)

合于 10-29 16:24,2 文件 +54/-46。主分支的 _make_jit_compile_func 在整段编译期间持一把锁,连只用默认编译选项的配置也串行,前端、下沉与设备编译因此全部被排队。改法是给每份覆盖配置用 dataclasses.replace 构造候选 JITImpl,后端、目标、输出选择、日志级别、编译标志与调试选项保持原值;默认选项仍走原有惰性调用路径,保留其内核缓存与对象身份;只对共享第一阶段键(由 get_tir 用的参数绑定器算出,保留位置参数、默认值与带张量注解的工厂)的候选串行;调试文件的共享输出仍在编译内部串行,覆盖 debug_root_path、IR 转储设置与布局可视化。验收数据在跨后端流水线上:昇腾任务合并 3be11be6 与主分支 c58a86f0,Python 3.12.14、torch_npu 2.10.0、CANN 9.3.0,1797 项检查通过,AscendC 示例 50 例全过、PTO 示例 44 例过 6 例跳过;ROCm 7.2 任务在 gfx942 上 3220 项通过、2765 项跳过;Metal 任务在 macOS arm64 上 95 项通过、5 项跳过;CUDA 任务在 H100 PCIe、cu130 上 5500 项检查通过、491 项跳过,示例 76 过 14 跳;CuTeDSL 任务 75 过 17 跳 2 例意外通过。

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

1.3 主仓:昇腾缓冲区视图的布局与指针单位不再跨别名传播(#3623)

合于 10-29 13:17,23 文件 +1139/-438。由 T.view 或 T.reshape 造出的多个缓冲可以共用同一段存储而要求不同布局,此前统一按共享的数据变量解析布局,于是可能选中错误的 Cube 映射或 DMA 地址,陈旧的分块几何还会丢掉版本偏移、或把带类型的指针偏移按错的单位解释。本笔把布局附着到单个 Buffer 视图上,贯穿布局推断、分形存储归一化、多缓冲物化与拷贝和矩阵乘的下沉;视图身份保留到布局映射与紧凑 L0 间距应用之后,裸指针偏移按各自的元素单位重映射。跨存储别名的布局不再传播,每个视图用自己的注解、算子推断出的布局或普通寻址;同一个带版本的分配上,Cube 视图现在可以用不同的 MajorK 与 MajorMN 布局,条件是逻辑步长稠密、基址偏移为零、操作数类型兼容且填充后的存储大小相等。多缓冲别名共用一个物理槽位大小,按各视图的元素单位换算而不传播索引映射;视图只是重新解释已有字节,写入方必须按消费者期望的表示写入,需要选视图给内建指令时可走访问指针接口。

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

1.4 主仓:GM 到 L1 的拷贝与越界填充合并成一个调度任务,去掉多余 MTE2 栅栏(#3543)

合于 10-29 11:39。全局内存到 L1 的拷贝与其越界填充共同准备同一块 L1 数据,此前填充可能被拆成独立的调度任务,同步插入就会在拷贝与填充之间加一道多余的 MTE2 管道栅栏。本笔把每个拷贝与其填充编进同一个昇腾任务,消费者等待的是包含填充在内的完整数据;FP8 MX 的 K 尾填充同样并入并保留任务内的条件填充,已有的显式任务作用域保持不变、不引入嵌套;原来的 MX 延迟填充队列被删除,每个操作数的填充紧跟自己的拷贝。47 项聚焦测试通过,覆盖普通越界填充与 MX 尾块两类。

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

1.5 主仓:CI 把 ROCm 与 Metal 任务锁回兼容的 PyTorch wheel(#3706)

合于 10-29 21:38。ROCm 的 CI 通过索引变量暴露匹配的 PyTorch ROCm 源,但工作流级的「不安全最优匹配」策略仍让更晚的 PyPI 候选胜出:#3534 的 ROCm 任务里解析器选了 PyPI 的 PyTorch 2.15.0(CUDA 13.2),而不是可用的 2.14.1 加 ROCm 7.2 构建,官方在运行器上直接复现了这两种解析结果。同一批 2.15.0 的铺开也打断了 Metal 任务:macOS 26.6.2 的 ARM64 运行器上,2.15.0 因为不暴露 Metal 2.0 而拒绝 Apple 半虚拟 GPU,同一镜像上上次通过用的是 2.14.1。改法有三条:按配置的稳定 ROCm 版本设置后端变量,让 uv 把 PyTorch 留在 ROCm 源上而其余依赖仍用最优匹配;夜间 PyTorch 安装改用首个索引,使夜间任务无法逃出配置的源;新增一个快速失败的校验步,在测试套件启动前确认 PyTorch 是 HIP 构建、不是 CUDA 构建、并能访问运行器 GPU。

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

1.6 新开缺陷三条:带填充的共享缓冲走批量拷贝报形状不符、负缩放迭代图误编译反向索引、视图位宽在 int32 溢出(#3704 / #3705 / #3707)

三条 issue 都指向静态检查的边界。其一,#3704 报告把末维不是 16 对齐的全局张量(例如 K=100)拷进带填充的共享缓冲(BK 取 cdiv(K,16) 乘 16 得 112)时,一旦这支拷贝被路由到批量拷贝路径,下沉直接报共享块与全局块形状不符(112 对 100),位置在 CUDA 拷贝的批量下沉处;逐元素路径下的带填充缓冲是正常的,而带填充恰好是满足矩阵乘对齐的常见写法,任何与批量拷贝同用的内核都可能踩到。其二,#3705 报告打平缓冲的简化在负缩放迭代图上误编译反向索引,复现环境是 v0.1.15 与 T-HEAD 名下 tilelang 分支仓的源码构建,报告同时核对了上游主分支与对应 TVM 修订里受影响的拆分函数未变,属上游共有路径。其三,#3707 报告视图与重塑在核对缓冲位宽时按 int32 相乘,超过 2 的 31 次方即回绕:Qwen3.8 的 lm_head 权重为 248320 乘 2560 字节的 NVFP4,共 5,085,593,600 位,位宽乘积回绕成 790,626,304,随后把静态视图位数转 int32 常量时抛出字面量超限。

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

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

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

1.7 三条修复在审:位宽改 int64、CPU 标量 GEMM 只清目标区域、关缓存时绕过两层进程内缓存(#3708 / #3709 / #3710)

对应上面三条缺陷与昨天的缓存语义议题,修复同日进审。位宽那笔把静态维度乘积改为按 Python 整数相乘、符号化维度转 int64,并把超出 int32 的字面量转成 int64 常量,回归用例直接取该 issue 的缓冲(248320 乘 5120 的视图形状),主分支报错、改后通过,符号化维度的同形视图也覆盖。CPU 那笔修标量矩阵乘在清零累加时清整块底层缓冲的问题:对 4 乘 4 输出缓冲把 7 初始化后调用 1 到 3 行列的子区域,原实现会把区域外 12 个元素一并归零,改法只用输出区域自身的偏移与范围初始化,回归用例照此写了全 1 的 2 乘 2。缓存那笔修昨天记录的 #3699:关闭内核缓存只翻转了缓存状态,而 jit 包装器在真正编译前还查两个进程内缓存,块内普通调用会拿到同进程早先编好的内核;改法是调用时读缓存开关(同时覆盖环境变量禁用),关闭期间跳过两层进程内缓存,出块后恢复。测试从改前的 1 失败 4 通过变成 5 全过。

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

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

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

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

2.1 昇腾:GLM-5.3 稀疏注意力 indexer 的 kpool 压缩算子合入,单 kernel 纯 Vector 路线(#1910)

合于 10-29 15:41。算子对每个 pool(pool_size 为 4)把打分张量与位置编码相加后做逐维 softmax 池化,池化向量过 BF16 边界,再乘归一化 Sylvester Hadamard-128 旋转(除以 128 的平方根),最后按每个向量的 FP32 二次幂绝对最大值缩放做 E4M3 量化:E4M3 行与缩放因子写进连续暂存区,由宿主适配器散射进分页缓存(键缓存用无符号字节承载 E4M3,缩放缓存原地写 FP32)。实现是单 kernel 纯 Vector 路线,开发者模式打开自动同步、自动向量与矩阵核合并、内存规划三个开关,每块覆盖 64 个 pool、两个向量子核各处理 32 个独立并行;宿主零重排,kernel 直接按调用方的三维布局跨步读取,位置编码折入池化;E4M3 编码是纯整数位运算,round-half-even、含次正规、饱和到 448,不使用 FP8 类型也不做近似。

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

2.2 昇腾:分块门控线性注意力的 Chunk Scaled Dot KKT 前向算子合入,静态四块展开、零 GM 工作区(#1803)

合于 10-29 15:01。新增分块内积的缩放点积前向:对分块门控线性注意力的前向计算 A 等于门口控的 K 乘 K 转置。实现为单个编译器装饰的 kernel,静态展开四个分块做片上融合,零全局内存工作区、零宿主算子调用,开发者模式四个开关全开。测试分层为 L0、L1、L2 与边界四档并可命令行选档,L0 覆盖三种情形(带门控、不带门控、小规模带门控),精度对照 PyTorch CPU 的 FP32 金标准,采用混合容差双门。报告给出的测形为批 1、序列长度 32768、头 32、每头维度 128、分块 64、开启门控,输入输出 BF16、累加 FP32,L0 与 L1 各三例全部通过,匹配率 1.0000,最大绝对误差在 0 到 0.25 之间。

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

2.3 昇腾:W4A4 的 int4 乘 int4 GEMM 结果全错问题修复合入,根因是类型尺寸谎报与专用转置指令被静默关闭(#1731)

昨天仍在审的这笔今天 15:08 合入,对应 #438。根因是三层错误叠加。其一是分形几何按类型尺寸等于 1 推导,而半字节打包类型每个元素只占 4 位、每字节装两个元素,于是 32 字节的 C0 块正确元素数是 64 而不是 32,所有拷贝步长、重复次数与字节计算全部偏差一倍,数据在 L0A 与 L0B 里堆到错的位置。其二是 B 操作数的转置加载被静默丢弃:通用加载路径对 zN 到 nZ 的转换设了转置标志,但四比特专用指令硬编码关闭转置并忽略该标志,B 永远得不到矩阵乘指令期望的布局;正确路径是带转置的专用加载,而它在 CANN 头文件里没有任何文档,只能用探针 kernel 在设备上实测反推参数(一次重复搬完完整的块,重复次数取 N 除以 64、源步长取 K 向上取整 16 后除以 64、目的间隙取 3),且该版本实际按 2048 字节块工作,与通用文档的 512 字节分形单位不同。其三是向量与全局内存之间的字节计数同样按每元素 1 字节而非 0.5 计算,每行拷贝长度偏长一倍,混入相邻行数据。计算指令本身始终正确,只是忠实地把摆错的数据相乘,症状是 99.6% 不匹配、输出为负 501 与 627 而正确值是 641 与 1785。报告同时写清了全局内存的 int4 打包约定:每字节两个半字节,低半字节为偶数元素、高半字节为奇数元素,第 r 行起始于 r 乘 N 除 2 的字节位置,与 CANN 自家格式一致。

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

2.4 昇腾:depthwise_conv_2d 示例重提,改用清单化测试通道以避开并发批次(#1920)

合于 10-29 15:43。这是 #1915 的重提:原笔合并后夜间全量 CI 失败,经 #1919 回退,根因是 L0 最大形状用例在八路并发的遗留脚本批次下触发未知错误;本笔把该用例瘦身四分之一,并把算子注册进持续集成的算子测试清单走 Pytest 通道,与另外二十五个已迁移算子同路线,彻底撤出并发批次,kernel 本体与原笔逐字节一致。算子本身支持 fp16、fp32、bf16 与任意核大小、步长、填充、膨胀,步长 1、2、4 走向量化区域路径,其余走标量内积兜底;实现是填充加卷积双 kernel 加 1 乘 1 快路径,纯 Vector 路线,网格按几何自适应不写死核数。报告里几处实测约束值得记:填充阶段耗时约等于语句条数乘 35 纳秒、与字节数无关,故末带用截断收尾;共享内存侧的拷贝目标偏移若不是 32 字节对齐会触发向量核的 0x26 异常;同步用手写冠军模型(卷积段关自动同步),通道间用深度 1 的相邻搬运与计算事件协议,深度 2 的协议或中途插入全栅栏都会死锁。

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

2.5 昇腾:新增算子性能剖析 skill,用 msprof op 输出定优化方向(#1931)

合于 10-29 17:17。新增一个算子性能剖析的开发流程资产:在性能优化前用 msprof 的算子级跑一遍,再按输出判断下一步优化方向。作者已在加法逐元素、归约取小示例与动态量化三个算子上验证过流程可用。这类资产的价值在于把「先量后改」写进流程,而不是每次重新摸索剖析口径。

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

2.6 昇腾在审:NPUIR 语言层三笔视图与广播修正(#1962 / #1960 / #1961)

三笔同日进审且互相依赖。其一,#1962 修本地统一缓冲的视图描述符:桥接层按共享数据变量查缓冲,可能返回未应用描述符的分配内存引用,同类型偏移别名于是读到分配区前缀,不同类型别名则带着分配的元素类型走进拷贝下沉,同类型重塑还可能带着错误维度进入归约;改法是引入统一的视图获取助手,同类型非恒等视图用带检查的重解释转换,多元素类型访问的分配用单块字节分配加带类型的视图操作保持共同存储身份。其二,#1960 修视图丢掉源偏移:视图从数据变量重建缓冲时丢掉了元素偏移,偏移别名因此从分配偏移零开始而不是源首元素地址;改法对同类型保留偏移,字节级类型变化时按元素宽度换算并要求可整除,不可对齐或不支持的转换直接拒绝,不引入存储或数据拷贝。其三,#1961 改一元向量广播按访问区域推断:此前按原缓冲形状推断广播维度,把 8 乘 8 源写到 24 乘 8 目标的行 8 到 16 会失败;改法改按生成的内存引用类型推断,保留访问区段的形状与既有兼容性检查。三笔都附了独立的干净基线构建与逐位验证说明。

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

2.7 昇腾在审:PTO 后端按位异或的无符号字节工作区视图修正(#1959)

进审于 10-29 16:36,关闭 #1721。PTO 后端里按位异或的临时工作区被当作原始一维字节区传给指令,而不是所需二维视图:生成时把目的端的二维形状复制到工作区的形状信息,却没标记为切片,于是对无符号字节(工作区不做类型重解释)解析出一维块描述,指令运行时越界读写。工作区本身是一维区按源元素数分配,容量对对齐的二维块是够的,失败在描述而非容量。修法是继承目的形状后强制构造二维视图并标记为切片,测试补了无符号与有符号字节的多组覆盖(二维自动临时、显式一维临时、一维、对齐非零偏移切片加哨兵、有符号字节),并保留 16 位回归。

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

2.8 海光:HCU 两笔在审有更新、社区出现 DCU 融合算子仓;摩尔线程与沐曦窗口内无推送(10-29)

海光侧仍在审:普通拷贝的共享内存纵向交错特性分支(#30)与分布式那笔(#19)窗口内均有更新,主分支最新提交仍停在 10-27 的缓冲区分组特性。社区侧同日出现一个 DCU 方向的仓库,把海光平台的融合算子与对 TileLang 编译器的改动一并保留,C++ 为主、约 15 MB,历史提交里可见注意力算子支持与可变长支持,属个人实践仓而非海光官方。摩尔线程适配仓(tilelang-musa)与沐曦适配仓(tilelang-metax)窗口内没有推送,最新分别是 09-30 与 10-26。

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

来源:https://github.com/SealSoul117/TileLang_DCU_FuseCmpComm

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

3.1 TileOPs:发布镜像合入,算子库与 TileFoundry 首次打进同一镜像(#2793)

合于 10-29 11:18,9 文件 +123/-49。运行器镜像新增发布目标:在开发镜像之上安装算子库与调度器的 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

3.2 TileOPs 在审:GEMM 按功耗封顶重算栅格分组与常驻块数,q-b 的 DRAM 读从 168 MB 降到 94 MB(#2794)

进审于 10-29 21:44,2 文件 +35/-12,承接背靠背调用的性能议题。两处改动:其一,当 A 能落进一个二级缓存分区而 B 不能时,栅格分组让一个分组跨全部行,于是每个 B 块只从内存读一次而不是每个十六行分组读一次,在 DeepSeek V3 预填充的 q-b 形状(4096 乘 24576 乘 1536,NT 布局)上内存读从 168 MB 降到 94 MB;窄块的分裂边界仍按十六行分组约束,选择不变。其二,常驻块数在宽度均匀的双协作组网格上改为按块数、总块数与处理器数算出向上取整的常驻数,而不是直接用处理器数,因为最忙的那块无论如何都要跑同样多的块,而每一波能读到更少的不同面板;乒乓、窄块与混合宽度网格仍用处理器数。测试在 H200、CUDA 13.2、PyTorch 2.13、TileLang 0.1.12 上跑算子测试,212 项通过、24 项跳过。

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

3.3 TileFoundry:AMX 运行时算子改成 39 个静态 TIR 契约,属破坏性变更(#279)

合于 10-29 15:59,破坏性变更(提交标题带感叹号)。动机是让 AMX 有一份显式、带类型的算子声明,智能体可以从同一份计算描述里匹配与评估。做法是把 24 个 AMX 运行时算子族描述成 39 个静态注册的契约版本,带参数域、操作数与布局模式、副作用集合,以及 M1 到 M4 的目标能力;新增 Python 评估器、原始 C 模拟器作为金标准,并补 59 个代表性原子用例覆盖全部 39 个静态算子,高阶与中阶中间表示用例跑真实的 2、4、8、12 线程域并验证候选的接受与拒绝;高阶候选从算子内的匹配声明发现,保留针对目标的能力推断。命名统一到 AMX 命名空间,别名由注册入口持有,注册名由统一查询给出;CPU 执行用线程拓扑,私有逻辑 AMX 寄存器状态,线程数不是 AMX 引擎数。验证为 46 项聚焦测试通过,含一处 57 个原始 C 结果的比对与两项生命周期检查。

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

3.4 采用方:TileKernels 与 FlashQLA 主分支窗口内无推送(10-29)

两个已知采用方的主分支都停在 10-14,窗口内无新推送:深度求索的算子仓(2018 星)与通义侧的 FlashQLA(742 星)。本日生态侧的新增信号主要来自社区自建仓而非这两个内部使用方。

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

来源:https://github.com/QwenLM/FlashQLA

3.5 昇腾侧出现独立的社区算子仓 ops-tilelang,窗口内有推送(10-29)

该仓 10-20 建立、10-29 13:06 有推送,定位是面向昇腾 NPU 的社区高性能算子仓,用统一方式管理 kernel 实现、自动化测试用例、标准工作负载与性能基线,目标是让算子资产可复用、可验证、可持续演进。它挂在昇腾生态的分支组织名下而不是 tile-ai 组织里,说明昇腾侧的算子供给正在从官方仓外溢出一层社区协作面;与官方 tilelang-ascend 的关系尚不明确,收录时只作生态观察,不视作官方动作。

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

3.6 PTO-ISA 配套双仓活跃:题库仓迁移到 PTOAS 0.70,基准仓加 SV2 分块缩放与逐列分析(10-29)

PTO-ISA 组织下的两个配套仓窗口内都有提交。题库仓把 PTO 的广播加载迁移到 PTOAS 0.70,此前一周在做量化缩放符号的澄清与题面「指向自己的解而不是答案」的改写;基准仓加了两笔分析提交:SV2 的分块缩放与逐列分析,以及一条依据(库内绝对最大值指令的变体计数为 31.39,附运行时指令吞吐与屋顶线分析,明确不是统一缓冲口径)。这两个仓与官方仓的关系是教学与测量配套,值得作为昇腾侧方法论的旁证。

来源:https://github.com/PTO-ISA/tilelang-puzzles-ascend

来源:https://github.com/PTO-ISA/tilelangGym

3.7 观察:T-HEAD 名下的 tilelang 分支仓窗口内有推送,分支名指向 PPU 的矩阵指令布局(10-29)

T-HEAD 组织(官网为开发者门户域名)名下有一个 tilelang 的分支仓,2026-09-22 建立,是上游主仓的分支,公开 13 个分支,除主分支与 v0.1.11 到 v0.1.14 的版本镜像外,还可见 PPU 矩阵共享布局修复、PPU language 包、运行时编译特性与若干缺陷修复分支;该仓主分支最近提交停在 10-22,但仓库推送时间更新到 10-29 19:27(北京时间)。仓内没有说明文档,也没有发布记录,因此只能作为观察线索:有整建制团队在把 TileLang 作为自有硬件栈底座做分支开发,方向与 PPU 的矩阵指令布局相关。本日 #3705 缺陷报告正是用该仓的源码构建复现的。

来源:https://github.com/t-head/tilelang-for-sail

四、社区、教程与活动

4.1 文档站把内建归约的适用边界写进指令页,与 redux 合入同步(10-29)

文档站窗口内有五笔自动更新,其中指令页新增了一段归约说明,与今天合入的硬件归约改造对齐:CUDA 的跨线程归约把与 warp 对齐、单位线程步长的满 warp 组委托给与 warp 归约同源的助手,助手在 SM80 以上的整数求和、取大、取小与按位运算上选硬件归约指令,更小的整数提升到 32 位;标量 FP32 的取大取小只在 sm_100a、sm_100f 与 sm_103a、sm_103f 上使用,标量 FP16 与 BF16 提升到 FP32;同一规则也适用于更大归约的最后一段与未打包的批量归约。子 warp、跨步、打包、FP64 以及不支持的类型算子组合保留 shuffle 实现,跨 warp 的共享内存交换与栅栏不变。把回退边界写进文档,比只写「支持硬件归约」对使用者更有用。

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

4.2 开发工具仓完成下沉跟踪抽取后的三笔收尾,与上游在审工具化改动同向(10-29)

独立开发工具仓在 10-20 完成从主仓抽出的下沉跟踪后,窗口内补了三笔收尾:避免跟踪激活时的循环导入、把持续集成固定到支持插桩插件的版本、在加载核心兼容别名前先初始化 TileLang。这与上游在审的「把下沉跟踪抽到开发工具仓」那笔同向,说明工具化已经不只是上游的分支实验,而是形成了可独立发布的包。

来源:https://github.com/easy-tilelang/tilelang-devkit

4.3 新仓 tilelang-knowledge 一天四笔:33 个算子设计案例、20 张带出处的知识卡、7 条优化模式(10-29)

该仓 10-29 建立,当天四笔提交构成一套面向智能体的知识底座:先是仓库结构,随后整理 33 个算子设计案例(每条带固定出处),再补 20 张带出处的编程知识卡,最后文档化 7 条有证据支撑的 GPU 优化模式。仓库自述的用途是让智能体面对一个算子时能判断哪些经验适用、应当如何改代码、如何验证、失败后往哪里找。这与 TileFoundry 的算子契约、昇腾仓的剖析 skill 是同一类尝试:把「怎么改、怎么验」写成可检索的资产。

来源:https://github.com/qiuuuuu622/tilelang-knowledge

4.4 图解课配套仓与分布式布局实验仓同日新建(10-29)

两个个人仓同日建立:一个图解课配套仓,自述用原创场景加动画把每个知识点讲透、随视频发布;一个分布式布局实验仓,提交说明是加入分布式的 TileLang 布局实验。加上今天有推送的模型移植仓与 Hexagon 后端仓(仓库推送时间更新到 10-30 06:10,其主分支最近提交仍是 08-24 的依赖升级),社区在教育与方言后端两个方向的自发投入都在继续。

来源:https://github.com/zachshi-ai/tilelang-course

来源:https://github.com/Leewin239/tilelang-distributed-layout

4.5 新闻侧:Google News 六组中英文查询、Hacker News 与 arXiv 均无窗口内实质命中(10-29)

六组查询覆盖精确词加两日、中文加昇腾、中文加算子、组织名加两日、英文加 GPU 与中文加编译器加七日。仅中文加昇腾返回一条,是临床人工智能进展的转载稿,与算子语言无关,其余全部空返回。Hacker News 近四日的相关命中全部落在窗口之外(最近的是一条 10-02 的昇腾编程栈转载讨论)。arXiv 以 TileLang 检索无新论文。新闻侧连续多日零命中,日报主体继续以仓库与生态仓为准。

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

五、趋势观察

昇腾侧的节奏已经从「搬后端」转到「补算子」。昨天记的是 NPUIR 后端并入主干、验收口径是 504 项自测与 71 比 72 的示例;今天合入的五笔里四笔是算子与示例,而且是按上游模型的需求排的:GLM-5.3 的 indexer 压缩、分块门控注意力的内积、四比特量化的矩阵乘、量化视觉模型的深度卷积。这些算子的共同点是都把量化与融合写进了 kernel 本体(E4M3 纯整数编码、四比特打包约定、双 kernel 加 1 乘 1 快路径),而不是留给宿主侧补救。同时两笔顺带把工程口径立起来:算子性能剖析进流程、并发批次造成的不稳定用例改走清单化测试通道。这说明昇腾侧的主要矛盾正从「能不能编出来」变成「算子的数值与性能能不能被稳定验收」。

主仓今天的三笔(硬件归约、并发编译、越界拷贝分组)分别对应指令选择、编译时延与发射质量,都是编译器内核问题,而不是语言层特性。这与昨天的新增缺陷方向一致:越界、别名、位宽这类静态检查的边界正在被逐个打开。今天三条新缺陷里有两条(带填充缓冲的批量拷贝、视图位宽溢出)都在真实模型形状上被触发,说明 TileLang 的使用面已经深入到会碰到边界条件的规模:248320 乘 5120 的 NVFP4 权重、K 不是 16 整数倍的注意力头维度。

生态侧今天最值得注意的是「fork 化」。昇腾社区出现独立算子仓,T-HEAD 名下有带 PPU 分支的 TileLang 分支仓,PTO-ISA 把题库与基准都做成了配套仓,另有开发工具仓把下沉跟踪独立发包。这些动作的共同前提是把 TileLang 当作自有硬件栈的底座来对待:需要按自己的指令集改布局、按自己的测试体系做验收、按自己的知识形态沉淀经验。对一个领域专用语言来说,这是比星标数更硬的采用信号,但也意味着上游要开始面对「分支与上游如何同步」的长期成本:今天 #3705 的缺陷就是分支用户在真实构建里先撞上的。

附、素材与核查说明

信源核查表

信源 核查结果
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-hygon、tilelang.github.io、TileOPs.github.io、TileOPs-nightly;tilelang-metax(10-26)、tilelang-musa(09-30)、TileRT、tilescale、DeepStack、tilelang-puzzles 无窗口内主分支提交
合入真值检索(is:pr is:merged) 12 笔:tilelang 5、tilelang-ascend 5、TileFoundry 1、TileOPs 1
在审真值检索(is:pr is:open updated) 54 笔,含昇腾文档核查批次、NPUIR 三笔视图修正与 TileOPs 的 GEMM 调度改动
新增 issue 核对 4 条:#3704、#3705、#3707 三条缺陷,另 1 条为昇腾仓每日测试机器人失败通知
主分支提交级核查 tilelang 5 笔、tilelang-ascend 5 笔、TileOPs 1 笔、TileFoundry 1 笔、文档站 5 笔;tilelang 与两仓合入数一致
发布页与 tag 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)
第三方与生态仓 10 个仓库逐一核对:昇腾社区算子仓、PTO-ISA 双仓、开发工具仓、知识仓、图解课仓、分布式布局仓、DCU 融合算子仓、模型移植仓、Hexagon 后端仓、T-HEAD 分支仓均有窗口内推送(Hexagon 主分支无新提交,仅仓库推送时间更新)
采用方 TileKernels 与 FlashQLA 主分支停在 10-14,窗口内无推送
Google News(中英六组) 仅一条中文加昇腾的临床人工智能转载稿,与主题无关,不收录
Hacker News 近四日相关命中全部落在窗口之外
arXiv 以 TileLang 检索无新论文

完整信源清单