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

信源:GitHub tile-ai 组织逐仓核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 文档站 / 语言服务仓 / 夜间基准仓)+ 合入与在审真值检索(合入 30 笔、在审更新 49 笔)+ 新增 issue 全量核对(11 条)+ 提交级 diff 与发布页核查 + PyPI 制品状态核查 + 采用方仓核查(TileKernels / FlashQLA)+ 社区新仓检索 + Google News 中英文多组查询(走代理)+ Hacker News + arXiv + 产业媒体报道。


本期索引

  • 今日重点:24 小时合入 30 笔、在审更新 49 笔、新开 issue 11 条;主线围绕「流与屏障语义」三处静默错;TileOPs 首发仍卡在 PyPI 最后一格;海光两笔大合入
  • 一、核心项目进展
    • 1.1 主仓:out_idx 输出改在调用方流上分配,H100 复现的五次三次损坏归零(#3649)(10-26)
    • 1.2 主仓:NVRTC 设备端编译接进 TVM-FFI 执行路径,新增编译键项(#3285)(10-26)
    • 1.3 主仓:批量 TMA 加载的 cache hint 少一个模板参数,DeepSelect 的 TMA 路径编译不过(#3683)(10-26)
    • 1.4 主仓:归约的 FP8/FP4 非有限常量写出非法字面量,亚字节累加器改为早拒(#3189)(10-26)
    • 1.5 主仓:SM120 分块缩放 GEMM 显式展开 K 原子,FullCol 配置提升 51%(#3286)(10-26)
    • 1.6 主仓:JIT 输出形状里的标量参数在分配前解析(#3602)(10-26)
    • 1.7 主仓:sm8x 稀疏元数据的操作数 dtype 不再无兜底落入(#3335)(10-26)
    • 1.8 主仓:DeepSelect Top-K 示例合入,129 组配置里相对 torch.topk 最高 8.00 倍(#3304)(10-26)
    • 1.9 在审:TIRx 重构规模再扩到 1194 文件,锚定未打补丁的上游版本(#3666)(10-26)
    • 1.10 新增缺陷:ThreadSync 三类错位栅栏,配三笔在审修复(#3680 / #3682 / #3681)(10-26)
    • 1.11 新增缺陷:sm_100 以上归约 batch 与布局冲突,已有一笔在审修复(#3667 / #3684)(10-26)
    • 1.12 在审:转置目标形状校验、区域化 kill 重排(#3677 / #3670)(10-26)
    • 1.13 已关闭:模糊测试系列的六条编译期崩溃(10-26)
  • 二、多后端适配(昇腾 / 海光 / 沐曦)
    • 2.1 昇腾主仓:带版本片上缓冲的精确定址接口(#3607)(10-26)
    • 2.2 昇腾在审:手工调度的 FlashMLA,BF16 预填充与 FP8 解码(#3678)(10-26)
    • 2.3 昇腾在审:片上内存溢出提前拒绝、E4M3 的 UB 到 L1 拷贝、视图布局与指针单位(#3655 / #3639 / #3623)(10-24~26)
    • 2.4 昇腾在审:批量调度可选、跨核同步要求批模式、L0C 双拷贝直转(#3675 / #3614 / #3468)(10-26)
    • 2.5 昇腾仓:安装器的 SHMEM 段双修,空壳 wheel 不再假成功(#1943 / #1946)(10-26)
    • 2.6 昇腾仓:W4A4 的 int4 矩阵乘结果错误,根因是半字节打包被当成 1 字节(#1731)(10-26)
    • 2.7 昇腾仓:超过两维的 GM 到 UB 拷贝按 DMA 行数折叠前导维(#1805)(10-26)
    • 2.8 昇腾新缺陷:Cython 包装器每次调用都造一个分析器,约 0.28 毫秒(#1945)(10-26)
    • 2.9 海光:寄存器流水线与 warp 发散合入,28 文件(#18)(10-26)
    • 2.10 海光:原生 B16 存储与动态 WASP GEMM 合入,scale 缓冲双 bank 在审(#28 / #29)(10-26)
    • 2.11 沐曦:组织视图有推送记录,六个分支均无窗口内新提交(10-26)
  • 三、生态与采用方(TileOPs / TileFoundry)
    • 3.1 TileOPs:首发仍在最后一格,制品仍是 2025-10-09 的开发版(#2738)(10-26)
    • 3.2 TileOPs:D128 稠密前向拆 key,prefill-rope 从 31.04 微秒降到 17.09(#2777)(10-26)
    • 3.3 TileOPs:D128 持久化前向按线程重写并写出 LSE(#2774)(10-26)
    • 3.4 TileOPs:融合加法的 LayerNorm 行跨 CTA 拆分,精度缺陷另开新单(#2768 / #2773)(10-26)
    • 3.5 TileOPs:逐元素与归约四条,clamp 的 bf16 行提速 5.9%(#2766 / #2767 / #2769 / #2770)(10-26)
    • 3.6 TileOPs:调用张量的对齐判断上移到算子层,三笔收尾(#2772 / #2778 / #2780)(10-26)
    • 3.7 TileOPs:偏移存储三类修复,SM90 分组 GEMM 与输出 ABI(#2761 / #2765 / #2762)(10-26)
    • 3.8 TileOPs 在审:FP8 查询的分页预填充改按 e4m3 乘 value(#2775)(10-26)
    • 3.9 TileOPs 在审:稀疏长序列的 GDN 预填充拆分扫描(#2779)(10-26)
    • 3.10 TileFoundry:动态共享内存按内核实际视图定尺寸(#271)(10-26)
    • 3.11 TileFoundry:整窗口按操作数本身下发,sm_89 的 GEMM 恢复可编译(#273)(10-26)
    • 3.12 TileFoundry 新缺陷:分析侧把共享内存峰值算重、且不传到 TIR(#272)(10-26)
  • 四、社区、教程与活动
    • 4.1 语言服务:实验性 CUDA 调试上线,接 cuda-gdb 与 Nsight(10-26)
    • 4.2 文档站:单日八笔更新(10-26)
    • 4.3 贡献流程治理四连:PR 描述规范、纯文档跳过后端测试、Agent 标题确认、评审机器人降噪(#3671 / #3672 / #3673 / #3674)(10-26)
    • 4.4 社区新仓:RISC-V 社区枢纽仓、CANN 算子仓、Hexagon 后端仓、调试工具仓(10-26)
    • 4.5 产业侧:昇腾与 DeepSeek 生态差距收窄的报道(10-26)
  • 五、趋势观察
  • 附、素材与核查说明

今日重点

24 小时内 tile-ai 组织合入 30 笔(tilelang 14、TileOPs 14、tilelang-hygon 2),在审更新 49 笔,新开 issue 11 条、关闭 12 条。昨日 72 小时窗口里 78 笔的量在单日回落到 30 笔,节奏正常,但内容重心明显从「性能清扫」转向「流与屏障语义」:主仓合入的输出分配修复(#3649)把 H100 上五次跑三次输出损坏的问题降到零,原因是三个适配器接受了调用方的 stream 却在当前流上分配 out_idx 输出,缓存分配器按分配流排序复用,于是别的流上排队的活覆盖了内核结果;同一天浮出的三笔 ThreadSync 缺陷(#3680)指出栅栏被放进发散分支、分支里的栅栏错误地清掉外层挂起依赖、以及按线程计数循环的栅栏位置,配 #3682 与 #3681 两笔在审修复,另有 tvm 分叉仓一笔把布尔位运算译成逻辑运算,让证明器不再在 10000 步预算内放弃(#88)。发布侧没有进展:TileOPs 的首发追踪单(#2738)只剩 TestPyPI 试跑与首个 tag,PyPI 上最新仍是 2025-10-09 的 0.0.1.dev2,仓库无 tag 无 release。后端侧海光一天合入两笔大会:寄存器流水线加 warp 发散(#18,28 文件 +5312/-64)和原生 B16 存储加动态 WASP GEMM(#28);昇腾主线合入了带版本片上缓冲的精确定址接口(#3607,14 文件 +982/-26)。在审最重的仍是 #3666,规模从昨日的一千出头扩到 1194 文件 +106648/-104497。


一、核心项目进展

1.1 主仓:out_idx 输出改在调用方流上分配,H100 复现的五次三次损坏归零(#3649)

合于 10-26 22:36,5 文件 +183/-31,修 #3648。cython、nvrtc、cutedsl 三个适配器都接受 stream= 并在该流上启动,但 out_idx 输出是用 torch.empty 在当前 torch 流上分配的;缓存分配器只按「分配流」排序复用,于是一块仍有排队工作的显存被当成空闲块返回,那批工作随后覆盖内核结果。修法是新增 output_allocation_stream(stream, device):调用方流与当前流不同时返回调用方流,否则返回空;句柄为零时映射到默认流,因为旧版 PyTorch 的 ExternalStream(0) 会给出池化流而不是默认流。H100 80GB 上的复现从 4/5、5/5、5/5 损坏变为三个适配器各 0/5,新增测试文件 10 项通过。作者明确否决了加 wait_stream 的写法:那会让每次调用都把启动流排在当前流的全部工作之后。

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

1.2 主仓:NVRTC 设备端编译接进 TVM-FFI 执行路径,新增编译键项(#3285)

合于 10-27 00:02,12 文件 +133/-133。新增 tl.cuda_compiler="nvrtc",与既有 TVM-FFI 宿主运行时配合:NVRTC 把设备端 CUDA 源编成 cubin,共享 CUDA 运行时模块导入这些字节,宿主侧绑定、标量转换、条件、循环与启动准备仍由 TVM-FFI 生成,默认值仍是 nvcc,旧的 execution_backend="nvrtc" 适配器保持原实现。二进制缓存键里加进了编译器名与 NVCC/NVRTC 版本,与目标细节和有效编译选项一起参与哈希;架构 token 如 90a 与 100f 原样保留,不兼容目标在编译前报诊断;每个成功创建的 NVRTC program 都在 finally 里销毁。该 PR 的 CI 在 H20 上跑出 5506 通过、491 跳过,示例 76 通过,ROCm 7.2 通道 3220 通过,昇腾通道 1797 项通过。

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

1.3 主仓:批量 TMA 加载的 cache hint 少一个模板参数,DeepSelect 的 TMA 路径编译不过(#3683)

合于 10-27 00:11,2 文件 +20/-11。eviction_policy="evict_first" 或 "evict_last" 的批量 TMA 加载在 main 上编译失败:代码生成发出一个 CacheHintSm90 模板实参,而裸指针版的 tl::tma_load 重载只接受屏障类型。这条缺陷直接压住了 DeepSelect 的 TMA 路径,在 GB200 加 CUDA 13.1 上用该示例的 ring 与 replay 用例复现。修法是从 #3296 里抽出缓存提示这一块,补上缺失的模板参数并把选定的策略传给 PTX 指令;共享内存对齐那一部分仍留在 #3296。修复后同一台机器上 DeepSelect 全文件 144 项通过、TMA 拷贝测试 22 项通过,CUDA 12.9 的 SM90 编译检查也确认三种策略的 PTX 都带上了应有的操作数(#3296 的其余改动尚未合入)。

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

1.4 主仓:归约的 FP8/FP4 非有限常量写出非法字面量,亚字节累加器改为早拒(#3189)

合于 10-26 15:44,7 文件 +204/-4,修 #2998。T.reduce_max 与 T.reduce_min 在输出为 FP8 E4M3 时发出非法的 inff 字面量;现在非有限的 E4M3、E8M0、FP4 常量改走既有的转换构造器并传入合法的 CUDART_INF_F 与 CUDART_NAN_F 表达式,饱和与 NaN 行为保持不变。同一笔改动让普通与延迟归约提前拒绝亚字节数值累加器(FP4、int4 之类),给出一句明确诊断;FP4 输入配 FP16 或 FP32 累加仍受支持,布尔归约豁免,E8M0 归约支持仍不在本次范围内。B300 加 CUDA 13.1 上 28 项聚焦测试全过,完整归约测试文件 150 项通过,另有三项批处理校验失败是 main 上已有的。

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

1.5 主仓:SM120 分块缩放 GEMM 显式展开 K 原子,FullCol 配置提升 51%(#3286)

合于 10-26 16:36,4 文件 +308/-2。这是 #3257 引入 T.gemm_blockscaled 与 T.mma_gemm_blockscaled 后的收尾:A 操作数取自寄存器分片时,K 原子循环原先用 T.serial,每次片访问都是运行期下标,nvcc 在不自行展开的流水化与 warp 特化内核里会把分片留在局部内存,也就是把 MMA 路径溢出掉;改为 T.unroll(..., explicit=True) 后每个访问都是编译期常量。第二处是 SFA 与 SFB 都在共享内存时,原来每对 MMA 要读三次共享,现在每 lane 每 K 原子只装一次紧凑选择包,行主序与分块 K 主序两种布局都覆盖。RTX PRO 6000 Blackwell 上 143 项通过、4 项跳过,输出与 float32 参考逐位一致;寄存器 A 基准里 FullRow 提升 2.7%、Square 11.1%、FullCol 51%(348.5 到 526.7 TFLOPS),ptxas 报告的溢写从 664 字节与 484 字节降到 320 与 160。作者同时标注:block_K = 128 配行主序共享 scale 的路径在 main 上本就不过校验,与本次改动无关,留待另查。

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

1.6 主仓:JIT 输出形状里的标量参数在分配前解析(#3602)

合于 10-26 10:06,7 文件 +242/-3,修 #3406。Cython 与 PTO 包装器以前在分配输出时,遇到用显式标量 PrimFunc 参数表达的输出维度就解析不出来;现在适配器传出「输出维度到标量参数位置」的映射,包装器直接用实参值,且不追加重复的启动参数。回归覆盖两个后端、零维与非零维、标量与多维输出、输出顺序、设备放置与启动参数个数。同一族缺陷的后续两笔仍在审:#3636 处理 (n, 2*n) 这类复合维度表达式,#3638 处理同名但彼此独立的动态符号被按名字合并的问题。

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

1.7 主仓:sm8x 稀疏元数据的操作数 dtype 不再无兜底落入(#3335)

合于 10-26 16:54,1 文件 +12/-5。稀疏元数据的布局构造函数按架构分派,把 MMA 操作数 dtype 与元数据缓冲宽度配对;sm90 分支会拒绝不支持的 dtype,sm8x 分支却写成两条白名单检查且没有兜底,于是其他拼写都会落到下面的放置逻辑,而那段逻辑只从缓冲宽度推 group 与 interweave。后果是 float8_e4m3fn(torch 的正规拼写)被按 16 位元数据接受,并套用了 fp16 的 group 32 与 interweave 4 约定;任何未识别的拼写包括打字错误也在任意宽度被接受。现在不支持的 dtype 直接报错。

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

1.8 主仓:DeepSelect Top-K 示例合入,129 组配置里相对 torch.topk 最高 8.00 倍(#3304)

合于 10-26 16:53,4 文件 +1265/-0。示例把 DeepSeek 的 DeepSelect 用 TileLang 重写,不含内联 PTX;作者实测 129 组配置、10 组不可用,在 RTX 5090(170 个 SM)加 CUDA 12.9 与 PyTorch 2.7 上跑重复输入与清缓存两种方法。重复输入口径下 BF16、批 6、N 为 16384、K 为 512 的一组时延 5.059 微秒对 torch.topk 的 40.448,加速 8.00 倍;N 为 1048576 的一组 15.503 对 100.254,6.47 倍。FP32 与 K 为 4096 的组合在该封装里不可用,K 为 1 的原生路径有对齐地址错误,作者在表里用横线标注「不可用」而非零。该示例依赖的缓存提示修复正是同日合入的 #3683。

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

1.9 在审:TIRx 重构规模再扩到 1194 文件,锚定未打补丁的上游版本(#3666)

10-25 开出、10-26 三度更新,当前 1194 文件 +106648/-104497。方案是把 TileLang 从内置 TVM 迁到独立发布的 TIRx,且对固定的上游基线打零补丁:删除 3rdparty/tvm 及其子模块项,3rdparty/ 只剩 CUTLASS、HIP 头文件与忽略文件;未修改的 apache/tvm-tirx 源码钉在 9af5544d,依赖 apache-tvm-tirx==0.1.0a2 与 apache-tvm-ffi==0.1.15.post0,原先带 +tilelang.1 后缀的集成补丁不再需要。这是当前队列里体量最大的一笔,也决定后续 TVM 侧修复(例如同日 tvm 分叉仓的 Z3 布尔位运算)如何进入主线。

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

1.10 新增缺陷:ThreadSync 三类错位栅栏,配三笔在审修复(#3680 / #3682 / #3681)

#3680 于 10-26 19:56 开出,指出 tl.ThreadSync 在三处放错栅栏:把针对逐线程取值的守卫当成块内一致,于是栅栏留在分支里,等的是跳过分支的线程;if 某一分支里的栅栏清掉了外层序列的挂起依赖,块跳过该分支时 if 前的写与 if 后的读就冲突;循环的行程数是逐线程的(如按线程号取模)时,栅栏被留在循环体内,部分同步的下降直接以「无法证明线程矩形」失败。#3682(3 文件 +39/-7 级别改动,12 项新测试)用一遍前置分析收集可能逐线程取值的缓冲,并把同步摘要按分支与循环的可达性收缩;作者声明该改动只报告竞争、不消除竞争,也不做分支拆分。#3681 修的是另一条同源崩溃:sm_100 起 float32 循环被向量化到 256 位,广播到 8 个 lane 时生成的表达式取了临时对象地址,NVRTC 以「必须是左值」拒绝编译,改为用 __float_as_uint 加移位直接拼位。配套的 tvm 分叉 PR #88 把布尔位运算译成逻辑运算,使证明器在 10000 步预算内能得出同一结论。

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

1.11 新增缺陷:sm_100 以上归约 batch 与布局冲突,已有一笔在审修复(#3667 / #3684)

#3667 于 10-26 08:31 开出:T.reduce_* 在 sm_100 与 sm_120 上指定 batch 等于 N 时无法下降,因为片段拷贝在 sm_100 以上走 256 位加载,每线程持有的输出数变少,一个在 sm_90 上合法的 batch 触发了硬校验。#3684 于 10-27 05:44 开出并附修法:ReduceOp 改为取 N 的不超过 batch 的最大真因数并给警告,finalize_reducer 保持严格因为其测试钉住了行为;新增的 sm_100a 与 sm_120a 下降测试在 main 上失败、在本分支通过,实测机为 RTX 5070 Ti。同一批归约问题的另一笔 #3341(批量全归约的线程范围解析)仍在审。

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

1.12 在审:转置目标形状校验、区域化 kill 重排(#3677 / #3670)

#3677 把 tl.transpose 的目标形状校验放进 TransposeNode 构造,在 NormalizeToAccessRegion 之后执行,因此绕过 Python 包装器直接构造也能覆盖;相等性用 arith::Analyzer::CanProveEqual 判定而不是比较 Python 维度列表,(3, n+1) 与独立构造的 (n+1, 3) 仍被接受,常量不匹配则以轴号与两侧长度报错。此前不匹配的目标会下降成越界下标,产生错误结果或越界写。CPU 目标 13 项测试通过。#3670 换掉共享内存活跃性分析里残留的 scope_level 机制:该值只在含 SeqStmt 的 For 体内递增、While 从不递增,于是两个方向都会错——While 情形下 kill 点没越过回边,打包器把互斥区域重叠,sm90 上实测结果每轮翻倍;兄弟循环情形下又会让 kill 越过整段区域边界。新规则按真实的外层 For 与 While 链与每段循环的连续区间计算,越界与提前停止两个方向都被夹住。

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

1.13 已关闭:模糊测试系列的六条编译期崩溃(10-26)

10-26 一次性关闭六条来自模糊测试的编译期崩溃或错误代码类缺陷:int16 与 uint16 目标上的 T.atomic_add 因 nvcc 无 16 位重载直接中止(#2572);非布尔的缓冲元素直接当 if 或 assert 条件时以内部 is_bool() 断言中止(#3009);浮点操作数喂给只接受整数的 div 与 mod 族时同样内部断言(#3008);int2、uint2 这类不在 1、4、8、16、32、64 之列的整型宽度让 CUDA 代码生成中止(#2996);reduce_max 与 reduce_min 写进 FP8 或 FP4 分片时崩溃(#2998);以及整缓冲目的地的原子归约被静默丢掉除 0 号线程之外的全部线程(#2959)。最后一条另有文档侧修正合入(#3340),把示例里的目的地从整缓冲改成具名元素,避免示例本身继续教出错误写法。

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


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

2.1 昇腾主仓:带版本片上缓冲的精确定址接口(#3607)

合于 10-26 10:04,14 文件 +982/-26。新增 T.annotate_buffer_addresses,允许在调度后的昇腾流水线里,对 UB、L1、L0A、L0B、L0C 与 SS 各池内的分配做精确字节定址;放入定址的缓冲必须先用 T.annotate_buffer_versions 声明固定份数。校验发生在版本展开与布局补齐之后,覆盖缺少或不支持的份数、重叠、未对齐、动态尺寸与容量溢出;分配器还要为 SIMT 的 UB 上下文预留与 AIV 私有地址空间让位,L0 的 scale factor 句柄跟随其数据分配。部分映射只钉住给定区间,其余缓冲在对齐后的空闲空间里安置且不允许复用已占用地址;没有定址标注的池保持原行为。子核投影时只保留活跃存储副本上的标注。

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

2.2 昇腾在审:手工调度的 FlashMLA,BF16 预填充与 FP8 解码(#3678)

10-26 19:31 开出,覆盖 Ascend 950 上的 BF16 预填充与紧凑 FP8 解码,含可选的 FP8 与 FP4 额外 KV cache、attention sink、非法下标与逐查询长度。注意力循环、在线 softmax、反量化与调度都用 TileLang 表达,只有小段 intrinsic 头文件包裹单条硬件指令。手工实现显式控制 AIC 与 AIV 流水与同步;持久化版本用 T.Stage 加编译器管理的缓冲版本与同步,QK 与 PV 使用同一 L0B 存储环的不同布局视图。

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

2.3 昇腾在审:片上内存溢出提前拒绝、E4M3 的 UB 到 L1 拷贝、视图布局与指针单位(#3655 / #3639 / #3623)

#3655 让 MergeUBAllocations 在下降期就按目标物理容量拒绝溢出:两个独立循环各自合法地选两份 32 KB 的 L0A 分片,合并后 64 KB 正好占满,此前会发出越过物理缓冲末尾的指针偏移且不带任何尺寸元数据,现在改为给出一句短的诊断。#3639 打开 E4M3 的 UB 到 L1 拷贝:ND 输入经既有的 scatter 与后拷贝流水自动重排,或直接拷入预排的 NZ 输入,复用既有的几何检查并给足迹与延迟分析补上 FP8 元素尺寸处理。#3623 让 T.view 与 T.reshape 造出的缓冲各自保持布局,避免经共享数据变量解析布局时选错 Cube 映射或 DMA 地址,并让裸指针偏移以自己的元素单位重映射。

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

2.4 昇腾在审:批量调度可选、跨核同步要求批模式、L0C 双拷贝直转(#3675 / #3614 / #3468)

#3675 给昇腾加可选的批调度(schedmode 1):正常模式下调度器只要所需核资源部分可用就启动算子,多个流的全核同步算子可能交错,批模式要等全部资源就绪再整体启动,全核同步场景必须用它;#3614 描述的是同一机制在 mode 0 下的症状与修法。#3468 处理 A5 上 FixPipe 不能同时开双目标模式与 quant_pre 的限制,让 T.dual_copy 直接做 L0C 的 fp32 到 f16/bf16 转换,省掉一次 Vector 转换与一块 fp32 的 UB 缓冲,也省掉 L0C 到 UB 的 fp32 带宽。

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

2.5 昇腾仓:安装器的 SHMEM 段双修,空壳 wheel 不再假成功(#1943 / #1946)

两笔在审 PR 指向同一个缺陷(#1728):install_ascend.sh --enable-shmem 在 build.sh 失败后不检查返回码,继续 bdist_wheel,package_data 的通配匹配不到文件也不报错,于是打出一个只有 __init__.py 的空壳包并 Successfully installed shmem-1.0.0,二十多行后才在 source ../../../install/set_env.sh 处报「设置 shmem 环境失败」,把排障方向引向环境配置。#1946 只改两处:加返回码检查立即退出,把 set_env.sh 改成基于脚本目录的绝对路径,并用逐字提取 shmem 段的受控 harness 验证(修复前 10 项中 6 失败 4 通过,修复后 10 项全过)。#1943 走另一条路:在卸载现有包之前依次检查原生构建、环境脚本、wheel 构建与 wheel 存在性,新增 10 项 CPU 控制流测试,并说明它修的是安装器续行与诊断,不含上游 SHMEM 编译失败本身。

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

2.6 昇腾仓:W4A4 的 int4 矩阵乘结果错误,根因是半字节打包被当成 1 字节(#1731)

#1731 在 10-26 更新,定位与复现都相当完整。整个 C++ 模板栈用 sizeof(AscendC::int4b_t) == 1 推导分形几何,但该类型是半字节打包,32 字节的 C0 块正确元素数是 64 而非 32,于是三层搬运全部出现二倍偏差:L1 到 L0A/L0B 的 LoadData 参数按错误常数计算,数据在片上堆叠到错误位置;B 操作数的转置加载被 load_cbuf_to_cb_s4 硬编码关闭并静默丢弃标志,B 永远拿不到 mad_s4 期望的 nZ 布局;UB 与 GM 之间的字节长度按每元素 1 字节算,每行多拷一倍并混入相邻行。计算指令本身是对的,于是症状是「量级正常、数值全错」(原报告 99.6% 不匹配)。修复全部落在 src/tl_templates/ascend/common.h 与两处变换,新增 W4A4 示例与 10 项 int4 回归,1024 立方端到端从 99.6% 不匹配变为零不匹配,既有 gemm_v0 套件 48 项、l1_to_l0 与 gm_to_ub 263 项、逐元素 492 项均通过。

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

2.7 昇腾仓:超过两维的 GM 到 UB 拷贝按 DMA 行数折叠前导维(#1805)

#1805 修 #1231。copy_gm_to_ub 与 copy_ub_to_gm 的 DMA 模型是二维矩形,而 AscendCopy::Lower 从 region 的活跃维里只取最后两个当行与列,超过两个活跃维时前导维被静默丢弃:3D 块拷贝 A[t:t+16, :, :](region 为 16、4、8,共 512 元素)被降级成 4 与 8,DMA 只搬 32 个元素,UB 分片其余 15 个 token 未初始化,后续计算读到垃圾值。修法是把前导维折叠进 DMA 的行数,两个方向同改。同日昇腾仓另有一批文档事实核查 PR 更新(sin 与 cos、axpy、clamp、bitwise 系列、clear、fill),其中 clamp 那笔顺带修了 PTO 上 T.tile.clamp 下界无效的问题。

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

2.8 昇腾新缺陷:Cython 包装器每次调用都造一个分析器,约 0.28 毫秒(#1945)

10-26 15:52 开出,报告者给出逐层定位与因果验证。CythonKernelWrapper.forward() 每次调用都无条件构造一个 TVM Analyzer(),本机中位数约 0.28 毫秒,而它只在动态 shape 分支里被用到一次;静态 shape 内核的 dynamic_symbolic_map 为空字典,simplify() 从不执行。因果验证的做法是把模块里的 Analyzer 全局换成「一调 simplify 就抛错」的桩,所有调用照常成功,全调用耗时从 442 至 467 微秒降到 104 至 107 微秒,且与内核规模无关(1 个 block 与 32 个 block 同量级)。实测环境为 Ascend 910B3 加 CANN 9.2.0。这对小算子密集调用的推理链路是固定开销,值得关注修法走向。

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

2.9 海光:寄存器流水线与 warp 发散合入,28 文件(#18)

合于 10-26 13:33,28 文件 +5312/-64,10-08 开出、跨两周多收敛。改动给 tilelang-hygon 加寄存器流水线,可与既有的共享内存流水线协同,对应的提交与屏障指令按流水级插入;warp 数超过 4 时启用 warp 发散。作者给出的结论是 example_gemm 达到与 rocblas 相当的性能。同一仓库当天还需注意 #18 的合并顺序:它先于 #28 合入,属于把主干改动并入分支后的大合并,因此行数远大于 #28。

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

2.10 海光:原生 B16 存储与动态 WASP GEMM 合入,scale 缓冲双 bank 在审(#28 / #29)

#28 合于 10-26 16:23,6 文件 +340/-230,三件事:保持 MLS 的 Let 参数发射顺序、把 GEMM 的 warp 划分传播到矩阵存储、拒绝打包的 B16 到 B32 输出别名;同时让 WASP 维度可动态,按轴推断对齐,修 pong 的 K 尾过滤,并把 WASP 端到端回归与基础 MLS 测试分开。#29 于 10-26 17:13 开出并在审:给 scale 缓冲的行规划一个与 LDS bank 分开的双 bank 分区,让分配、scale 拷贝与 MMAC 下降共用同一套行映射,保留物理区间生命周期与 m0 对齐分页,把空闲 wave 的工作摊到两个 bank 上,并用于手工 warp 特化下的分块缩放 GEMM。

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

2.11 沐曦:组织视图有推送记录,六个分支均无窗口内新提交(10-26)

组织页显示 tilelang-metax 在窗口内有推送时间戳(北京时间 10-26 15:08),但逐分支核对后没有找到对应提交:默认分支 main 最新仍是 2026-09-29 的版本号提升,dev 分支最新 2026-10-16 的一次合并,0.1.14 分支 2026-09-17,0.1.12 与 race 更早;事件流里也没有窗口内的推送事件。据此按「无实质变更」处理,不计入本期条目,也不猜测其性质。该仓最近一次正式发布仍是 10-16 的 0.1.15。

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


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

3.1 TileOPs:首发仍在最后一格,制品仍是 2025-10-09 的开发版(#2738)

首发追踪单 10-26 更新,八项清理已完成六项:命名与清单参数、打包与仓库文件、Python 3.13、公共 API 与扩展点、算子行为约定、输入校验生成、文档站。剩下两格:夜间基准(修两条坏掉的基线 #2763 已合入;把新行拉到所有对照器之上,累计扫描改成单次启动 #2766 已合入,GQADenseWithLSEFwdOp 的 D128 程序 #2774 已合入,但 MHA 行仍只有 FA3 的 0.996,调优继续)与发布本身(第三方声明、PyPI 安装文档、开发者依赖组、CHANGELOG 与发布步骤 #2753,以及需要维护者执行的 TestPyPI 试跑与首个 tag)。核对 PyPI 与仓库:tileops 最新仍是 2025-10-09 的 0.0.1.dev2,共三个版本,仓库没有任何 tag 与 release。首发处在「文档就绪、制品未上传」的最后一格,已持续多日。

来源:https://github.com/tile-ai/TileOPs/issues/2738

3.2 TileOPs:D128 稠密前向拆 key,prefill-rope 从 31.04 微秒降到 17.09(#2777)

合于 10-26 22:45,2 文件 +527/-124。D128 稠密调用若 128 行的瓦片让 SM 闲着,就把每个瓦片的 key 拆成 min(SM 数除以瓦片数, 最短瓦片的 key 步数) 份,且只在拆分能从最长瓦片匀掉两步 key 时才拆;每份写出带本行缩放最大值与和的未归一 fp32 输出,合并内核按最大最大值加权,因此抵消项与大 logit 都不丢精度。瓦片最后一步 key 中,若某 warpgroup 最后需要的 key 落在低 64 个 key 内,softmax 只对那 64 个做掩码与指数,warpgroup 0 的因果对角步必然命中。H200 1500 MHz 上 prefill-rope 的 f16 与 bf16 从 31.04 与 31.01 降到 17.09 与 17.09 微秒,对 FA3 是 2.324 与 2.328 倍;qwen3-8b 分块预填充从 31.30 降到 17.79。测试节点净增 14。

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

3.3 TileOPs:D128 持久化前向按线程重写并写出 LSE(#2774)

合于 10-26 18:22,2 文件 +656/-48。根因是 GQADenseWithLSEFwdOp 以前跑在非持久化的 warp 特化内核上,因为持久化内核拒绝 return_lse,夜间基准实测只有 FA3 的 0.73 至 0.78。本次用 TileLang 的逐线程原语(TMA 映射、mbarrier、WGMMA 描述符、warp shuffle、stmatrix)重写 D128 持久化稠密前向,直接写出 log-sum-exp,同一份实现同时服务 GQADenseFwdOp 与 GQADenseWithLSEFwdOp;当全部 key 与 value 超过 L2 的一半时,一个头的查询块相邻排布,使并发瓦片共享 L2 里的 K 与 V;每个瓦片的首次 QK 在上一个瓦片的尾声里启动。公共接口与清单未变。合并后 MHA 行仍为 FA3 的 0.996,追踪单标注调优继续,这与 #2777 属于两条独立的优化线。

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

3.4 TileOPs:融合加法的 LayerNorm 行跨 CTA 拆分,精度缺陷另开新单(#2768 / #2773)

#2768 合于 10-26 15:45,5 文件 +525/-24。起点是一次夜间基准的回退(llama-70b 解码行从 0.0037 升到 0.0044 毫秒),查下来不是代码变化而是调优重选:调优键改动后 FusedAddLayerNormKernel 重新调优,调优器从 L2 背靠背计时候选,选中了 1024 线程,冷启动比 256 线程慢 19%。新写的 FusedAddLayerNormHeldKernel 把一行(SM90 上是一个 cluster CTA 的份额)留在寄存器里,在最忙 SM 的成本更低时把一行拆到最多 8 个 CTA,块大小取整个 warp 数中能整除份额访问量的那个;派发按「排队行数」与「寄存器超过 255」两条互相拒绝来选,调优只提供受调构建份额能填满的宽度,使解码行不可能再被重选。llama-70b 解码 2.368 微秒对 flash-attn 的 4.800。同一笔改动暴露的精度问题另开 #2773:行远离零时 fp32 求和的均值与方差失去精度(常数 10000.1 的行归一化后约 0.7 而非偏置),靠近 1e36 时平方和溢出导致全 NaN,作者的建议是改用 Welford 或缩放两遍法。

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

3.5 TileOPs:逐元素与归约四条,clamp 的 bf16 行提速 5.9%(#2766 / #2767 / #2769 / #2770)

四条同日合入,都属「先找根因再改」。clamp(#2767)改为在自身精度内比较,根因是旧实现把每个元素加宽到 fp32、每个界都做一次 NaN 判定与选择,约 5 条指令一个元素;现在每对元素一条 HMNMX2.NAN,16384 见方 bf16 从 0.25704 毫秒降到 0.25024,对 torch-compile 1.0587 倍,并在全部 16 位模式、NaN、正负无穷与带符号零的边界上验证与 torch 的 fp32 比较逐位一致。累计扫描(#2766)把 dtype=float32 的 cumsum 与 cumprod 从「先加宽成一份再扫」改成一次启动直接写目标精度,hidden-state-f32 工作量下搬运从 112 MiB 降到 48 MiB,cumprod 47.36 降到 14.62 微秒、cumsum 47.86 降到 14.62。GELU 乘(#2769)改为经共享 gelu() 在 fp32 求值并只舍入一次,旧的 erf 路径带 NaN 选择分支、三次 16 位舍入与更长的依赖链;小门控网格的块数从 56 降到 14。二元运算(#2770)把二字节操作数的每线程元素数从 8 提到 16,保持两条 16 字节载入在飞,34 行默认配置随之变化,MinimumFwdOp 的 llama-7b 行从 8.42 降到 8.00 微秒。

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

3.6 TileOPs:调用张量的对齐判断上移到算子层,三笔收尾(#2772 / #2778 / #2780)

#2772 合于 10-26 15:28,55 文件 +373/-224:执行器在实现运行前把存储起点不在 16 字节边界的调用张量拷到对齐存储,成功后把被写的那几个拷回,并返回调用方的张量;若被拷贝的张量与其他实参重叠(且至少一方是被写的),在运行前直接拒绝,因为拷贝会掩盖内核自身的检查;无地址的存储(假张量或拒绝读取的包装)不拷贝;内核侧的对齐与暂存逻辑(含 MoE 分组 GEMM、因果卷积状态、融合 QK-norm RoPE 的 qkv)全部删除。代价是热路径上对齐调用的宿主时间涨 1 至 2 个百分点,MoEPostPermuteFwdOp 带 out 时 4.4 个百分点。#2778 把这项判断折进生成的观察器并让计划缓存决定拷贝,热路径只读该决定,执行器原来单独的地址遍历被删。#2780 于 10-26 开出:目前只有四个内核检查 out 与输入重叠,其余带 out 的算子会静默竞争,且每次检查都重复读 data_ptr(MoEPostPermuteFwdOp 一次调用多 12 次),方向是在清单里逐算子声明是否允许重叠,在算子层统一检查。

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

3.7 TileOPs:偏移存储三类修复,SM90 分组 GEMM 与输出 ABI(#2761 / #2765 / #2762)

#2761 让分组 GEMM 的 MMA 实现接受起点带偏移的连续张量,并对齐向量访问,把结果写回偏移输出时保持标识与未触及的专家行,拷贝前检查重叠;原实现遇到这类视图会以 CUDA 地址未对齐失败,新增 16 个回归覆盖 FP16 与 BF16 各自的 a、b、bias、out 偏移,配紧凑与掩码两种布局。#2765 把同一能力补到 SM90 的 MoE 分组 GEMM:存储起点不在 16 字节边界时拷 a、b 并暂存 out,vector_aligned 现在写明「被写的张量要暂存并在启动后拷回」这条规则,BF16 与 FP16 的 MMA 内核改用它,FP8 内核从 SM90 FP8 内核继承。#2762 修输出 ABI:当 TVM FFI 适配器声明由被调方分配输出时,交给 TileLang 课执行体自己分配并返回,其他适配器保持原有分配路径;报告的回归是旧实现在新构建上 27 个 Mamba 用例全败,候选实现在公开调用路径上通过。

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

3.8 TileOPs 在审:FP8 查询的分页预填充改按 e4m3 乘 value(#2775)

10-26 17:20 开出,10-27 07:00 仍更新。GQAPagedFP8WSKernel 把 FP8 查询的 value 乘积改在 e4m3 里做:负责转换的 warp 把每个 value 瓦片转置成 K 主序的 e4m3 操作数(ldmatrix 转置、按字做一次字节置换、stmatrix),在每个 16 内排好 key 的顺序,使每线程的分数正是它供给的概率,再用寄存器内打包把概率舍入到 e4m3。softmax 按 FA3 的 FP8 形式:2 的 8 次方偏移让小概率不进 e4m3 的次正规数,行最大值随每个更大的分数移动,key 瓦片从最后一个往回扫。H200 上 llama-8b-fp8-prefill 从 369.3 降到 339.8 微秒(相对 FA3 从 1.000 到 1.087),qwen3-fp8-ragged-odd 从 130.9 降到 123.9。作者如实标注精度容差从 2 的负 4 次方放宽到 2 的负 3 次方,原因是 e4m3 概率本身有半个 2 的负 4 次方量级的舍入,FA3 在同类输入上同样超出。

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

3.9 TileOPs 在审:稀疏长序列的 GDN 预填充拆分扫描(#2779)

10-26 开出。在 SM89 上对稀疏的长序列 GDN 预填充网格复用既有状态扫描与独立输出程序:适用条件是稠密、key 主序的 D128 输入、至少 16384 个 token、且状态扫描 CTA 不超过 SM 数的四分之一,其他调用保持原实现。分解方式写进程序缓存标识,并补了尾块与初始状态的常驻回归。实测(CUDA 13、PyTorch 2.13、TileLang 0.1.15)32768 长度 bf16 为 1.0384 毫秒对 FLA 的 1.0656,65536 为 2.1178 对 2.1786,16384 则仍是 FLA 更快(0.4695 对 0.4409)。作者明确说明收益只在这组四个 value head 上成立,八个与十六个 value head 的既有模型形状不选这条路径,候选也不普遍快于 FLA。

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

3.10 TileFoundry:动态共享内存按内核实际视图定尺寸(#271)

10-26 开出并在审,12 文件 +340/-32。合成的宿主入口以前给每个 TIR 内核都传 dynamic_smem=0,用整数 TensorView 偏移访问共享内存的内核因此在没有动态共享内存的情况下运行,超过 48 KiB 时也从不提升 cudaFuncAttributeMaxDynamicSharedMemorySize。作者选择按内核实际达到的最远字节数来定尺寸,而不是按目标容量,理由有两条:CUDA 把静态与动态共享内存算在同一个每块限额里,而该运行时自带若干静态共享内存(ProgramMetaData 或共享的分配张量),按满容量申请会直接起不来;每块申请多少还决定一个 SM 上能放几块,sm_89 上申请 99 KiB 时哪怕只用 512 字节也只能放一块。代码生成同时改为在首个语句前声明一次 smem_base,因为在 loop_outputs 里它原先落在第一个循环的块内,第二个循环看不见,nvcc 报未定义标识符。H200 与 sm_89 上两项 GPU 测试通过,未打补丁时前者非法访存、后者编译失败。

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

3.11 TileFoundry:整窗口按操作数本身下发,sm_89 的 GEMM 恢复可编译(#273)

10-26 开出并在审,5 文件 +54/-24。schedule finalize 给 sm_89 GEMM 写出的代码编不过:下降把整个累加器包进一个从零开始、覆盖其全长的窗口 T.ptr_of(acc[0:0+16, 0:0+8]),而 CUDA 代码生成对分片寄存器瓦片的指针切片报「需要普通源布局」;再往下 MMA 把寄存器分片读成 lhs_frag_dst[0:0+16, 0:0+32] 这类裸切片,同样没有处理分支。这些窗口其实都是整个操作数。修法放进下降:ConvertHIRToTIR 一次判定窗口是否覆盖其操作数(起始为静态零、尺寸等于操作数形状),三个构造窗口的地方共用该判定,最终写出 T.ptr_of(acc),MMA 直接读分片。49 个有 TIR 孪生的调度夹具里 29 个的 finalize 输出按此改变,sm_89 GEMM 与 mesh 视图恢复可编译,两个 wgmma 输出越过了指针切片拒绝、停在下一处视图,而那一处布局在 main 上同样被拒。

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

3.12 TileFoundry 新缺陷:分析侧把共享内存峰值算重、且不传到 TIR(#272)

10-26 开出,作者是在评审 #271 时发现的,两条缺口都指向同一种后果:分析能通过的程序在启动时被拒,或最终 TIR 无法自述需要多少共享内存。第一条,内存分析把 smem 与 facts.memory.shared.per_cta 比较(sm_89 上 101376 字节),但 CUDA 把这个限额同时作用于静态与动态共享内存,而运行时自身持有 12 字节的 __shared__ ProgramMetaData,这项占用只写在 runtime.cuh 里,没有告知分析,于是恰好差在 12 字节以内的放置能通过分析而在启动时被拒。第二条,分析把每 CTA 的共享内存高水位记在 RegionMemoryMetadata.peaks,但 ConvertHIRToTIR 不把它写进产出的 TIR,PrimFunction.metadata 既不打印也不解析,于是 finalize 之后必须重算一遍(#271 正是靠扫描所有整数地址的共享内存视图重算),而作者写的宿主入口申请过少也不会在运行前被发现。

来源:https://github.com/tile-ai/TileFoundry/issues/272


四、社区、教程与活动

4.1 语言服务:实验性 CUDA 调试上线,接 cuda-gdb 与 Nsight(10-26)

tilelang-lsp 在 10-26 17:11 推入首个调试提交(78a5761,标题为加入实验性 TileLang CUDA 调试)。内容四块:新增 tlsp.debug 启动器,做 Python 行号映射、加 -G 与 -lineinfo 编译、隔离缓存并抓取生成源;把 Nsight 的 cuda-gdb 接进 VS Code,支持 Python 断点、启动前检查与生成 CUDA 源的浏览;补单元、启动配置与可选 GPU 验收测试(覆盖待定断点、设备端单步、线程检视与数值);中英文文档与 CI 同步更新。这是该仓自 8 月以来第一次实质提交,此前最后一次更新是 8-04 的 CI 测试。

来源:https://github.com/tile-ai/tilelang-lsp/commit/78a5761404

4.2 文档站:单日八笔更新(10-26)

tilelang.github.io 在窗口内推送八次(北京时间 07:32、10:16、10:21、16:51、17:25、22:42、次日 00:08、00:16),提交标题统一是 Update docs,属于随主仓合入自动同步的文档构建。同日还有第三方本地化仓(tilelang-docs-l10n)一次推送与中文 RISC-V 社区枢纽仓的首次内容提交,说明文档同步链路正常运转。

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

4.3 贡献流程治理四连:PR 描述规范、纯文档跳过后端测试、Agent 标题确认、评审机器人降噪(#3671 / #3672 / #3673 / #3674)

四条同日合入,全部围绕「评审成本」。 #3671 在 CONTRIBUTING.md 里写明 PR 描述应有问题、方案与取舍,证据要贴在它支持的论断旁边,缺陷修复写前后行为或回归测试的意图与结果,性能改动要给可比测量与条件。#3672 让纯文档 PR 跳过 CUDA、ROCm、Metal、Ascend 与 CuTeDSL 的构建与测试,做法是用既有 lint 作业的检出看合并基差异,改动路径全是 Markdown、reStructuredText 或 docs/_static/ 下时跳过,白名单刻意收窄,可执行示例、构建文件与工作流不在此列。#3673 给提交技能加确认步:最终 diff、验证与 PR 描述就绪后,先把拟用标题原文给贡献者看,由其确认或替换。#3674 关掉评审机器人的自动摘要与文档字符串覆盖率检查,并加共享评审指示,要求把发现锚定在当前 diff 与实际调用路径上、区分回归与既有局限、先比较已有实现再提修法。

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

4.4 社区新仓:RISC-V 社区枢纽仓、CANN 算子仓、Hexagon 后端仓、调试工具仓(10-26)

窗口内出现若干个值得登记的外部推进。TongmingLAIC/tilelang-riscv 于 10-23 建立、10-26 推送,定位是 TileLang 在 RISC-V 上的社区枢纽,聚合编译器后端、AI 算子、开发工具与文档,目前仍是空骨架,但方向与国内 RISC-V AI 栈的诉求一致。Atlas-Ascend-Fork/ops-tilelang 是 CANN 社区面向昇腾 NPU 的 TileLang 算子仓,强调统一管理内核实现、自动化用例、标准负载与性能基线,10-26 有推送。xwhzz/tilelang-hexagon 做高通 Hexagon 后端,仓库 32 MB,10-26 有推送。superAngGao/tilelang-debugger 是 H200 上按源定位的调试工作区。另有 RuneFang/tvm_tilelang_cookbook(37 星)与 localizethedocs/tilelang-docs-l10n 的常规更新。

来源:https://github.com/TongmingLAIC/tilelang-riscv

4.5 产业侧:昇腾与 DeepSeek 生态差距收窄的报道(10-26)

10-26 有媒体转述 SemiAnalysis 10-22 的报告,称 DeepSeek 与昇腾的深度配合正在收窄与英伟达在芯片硬件与软件生态上的差距,并预计 DeepSeek 到 2027 或 2028 年将主要依赖昇腾技术栈;报道特别提到,这份判断出现在 DeepSeek 发布面向昇腾生态的软件工具数周之后,其中包含六个与英伟达平台一一对应的软件模块。该消息本身不涉及 TileLang 的功能变更,本期仅作生态背景登记:这六个模块里的高级语言编译工具即 tilelang-ascend,昇腾仓当日的安装器与 int4 矩阵乘修复正落在同一条链路上。

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


五、趋势观察

一、正确性缺陷的重心从「性能回退」移到「编译期不报、只在特定架构或存储域上错」。当日主线三笔都属此类:输出在错误的流上分配(#3649,只在非默认流且当前流有排队工作时损坏)、TMA 缓存提示少一个模板实参(#3683,只在带 eviction_policy 的批量加载上失败)、sm8x 稀疏元数据没有兜底(#3335,只有 torch 正规拼写才被认出,且套错 fp16 的布局约定)。这类缺陷的共同点是有硬件或参数前提,常规 CI 覆盖面之外,因此越来越多由模糊测试系列驱动——本章当日关闭的六条里有五条来自该系列。

二、同步与内存规划成为主战场。当日浮出的 #3680 三条 ThreadSync 缺陷、配 #3682 与 #3681 两笔修复、tvm 分叉仓 #88 的证明器修正、主仓 #3670 的区域化 kill 重排、TileFoundry #272 的共享内存峰值不落 TIR,指向同一件事:编译器的同步与内存判断正在从「启发式」转向「可证明」。这也解释了 #3666 为何重要——它把内置 TVM 换成独立打包的 TIRx,此后这类证明器修正才能以正常版本节奏进入主线,而不是停在分叉仓里。

三、后端适配进入「精确控制」阶段。昇腾的 T.annotate_buffer_addresses 让 UB、L1、L0A、L0B、L0C、SS 各池可以按字节定址并做重叠、对齐与容量校验,海光的寄存器流水线、warp 发散与 scale 缓冲双 bank 分区,都是把「能跑」推向「能按硬件结构摆」。这些能力集中在最近两周,说明国产后端已过可用性门槛,开始处理性能上限问题。

四、发布工程仍是短板。TileOPs 的首发在文档、清单、打包、API 稳定性、输入校验全部就绪之后,卡在 TestPyPI 试跑与首个 tag 上已多日,PyPI 制品停留在 2025-10-09。对一个 196 星、当日合入 14 笔的算子库来说,可安装性比再优化 1 个百分点更能决定外部采用速度;追踪单里唯一未完成的技术项是 MHA 行 0.996,管理层不应让它继续阻塞 tag。

五、外部生态在同步扩张。RISC-V 社区枢纽仓、CANN 算子仓、Hexagon 后端仓同日出现或更新,说明 TileLang 正在从「tile-ai 组织的项目」变成多平台各自建仓的生态。对国内读者,RISC-V 方向尤其值得跟踪:它与此前 AI 眼镜、端侧推理的硬件路线直接相关,但目前仍只有空骨架,实质性进展需要看后续是否有可运行的算子与后端落地。


附、素材与核查说明

信源核查表

信源 核查结果
GitHub tile-ai 组织(29 仓) 窗口内推送 7 仓:tilelang、TileOPs、TileOPs-nightly、tilelang-hygon、tilelang-metax、tilelang-lsp、tilelang.github.io
合入真值(is:merged) 30 笔,逐笔核对合入时间落在窗口内
在审真值(is:open 更新) 49 笔,按更新时间排序全量核对
新增 issue 11 条,逐条读正文
关闭 issue 12 条,其中 6 条为模糊测试系列
PyPI(tileops / tilelang) tileops 最新 0.0.1.dev2(2025-10-09);tilelang 最新 0.1.15(2026-09-30)
发布页 tag 核查 TileOPs 无 tag 无 release;tilelang 最新 v0.1.15(2026-09-30)
tilelang-ascend 仓 窗口内无推送(上次 10-23),但 PR 与 issue 活动密集
tilelang-metax 仓 组织视图有推送时间戳,六个分支均无窗口内提交,按无实质变更处理
采用方(TileKernels / FlashQLA) 均无窗口内推送(10-14)
Google News 中英文多组查询 24 小时内无 TileLang 组件精确命中;一条产业侧报道见 4.5
Hacker News 窗口内一条前端条目,与 TileLang 无关,剔除
arXiv 关键字检索无窗口内新论文,最近一篇为 2026-07-24
会员与采用方仓社区检索 六个外部仓窗口内有推送,见 4.4

完整信源清单