报告周期:2026 年 11 月 2 日(周一)至 11 月 6 日(周五),共 5 个工作日 素材来源:本刊每日 TileLang 社区动态日报 5 期(11-02、11-03、11-04、11-05、11-06),五期采集窗口首尾相接,合计覆盖 10-30 07:00 至 11-06 07:00 共 168 小时(周一那期为 72 小时窗口,覆盖上周末;周二至周五为 24 小时),与上一期周报(10-26 至 10-30)无重叠;另经 tile-ai 组织各仓库 168 小时窗口补充核查、各仓发布页与标签、TileLang×RISC-V 四源巡检复核(详见附录)


一、本周要闻

  • 昇腾 950 的持久化 FlashMLA 内核直接写进主线,用编译器可调度的写法而非另仓移植(11-02 合入,#3678):一笔新增 1997 行、删除 17 行、20 个文件、42 次提交,覆盖 BF16 prefill 与紧凑 FP8 decode 两条路径,并支持可选的 FP8 与 FP4 额外 KV 缓存、attention sink、非法索引与逐 query 独立长度。注意力循环、在线 softmax、反量化与调度全部用 TileLang 表达,靠 T.Stage 让编译器管理缓冲版本与同步;QK 与 PV 共用一块 L0B 存储环的两个不同 layout 视图,L1 落位由 T.annotate_buffer_addresses 直接指定;prefill 在一个累加器分配里分出四个不相交的输出切片,归一化算完一片就排空一片。为了让这套写法成立,编译器侧同期补了三处能力:L0C 到 UB 的拷贝保留 NZ 输出、每个 AIV 流在自身流水线之前选定、显式调度的阶段在物理发射相位相同时保持源码顺序。目标形状写得很窄:头数 64、Dqk 与 Dv 都是 512、topk 为 64 的倍数。来源:https://github.com/tile-ai/tilelang/pull/3678

  • ThreadSync 一周三改,先为迁移让路、再把性能与证明补回来(11-01、11-04、11-05):11-01 的 #3733 把三周前合入的精确线程计数(#3291)整笔回退,理由写得很直白——那套改动把昂贵的计数查询搬到满块栅栏检查之前,即使所有线程都会到达栅栏也要跑证明,TVM 的求解预算被耗尽后返回部分计数,合法的 remap 与 sequence 辅助 kernel 直接以「Cannot prove a thread rectangle」编译失败;商守卫、精确参与线程界、非矩形线程集三类正确性改进一并撤回,#3732 留作迁移后重新优化的跟踪单。11-04 的 #3761 则处理回退派生的副作用:为了修 #3643 的命名屏障重号,同一组线程跨作用域也不能复用 ID,TileKernels 的 mhc_pre_big_fuse 内核每块 96 线程因此在两段式 cp.async 流水线上多拿一个 ID,编译产物声明的屏障资源从 4 变 5,B300 上隐藏维 1280 的屏障占用上限从 16 块/SM 掉到 12 块/SM;修法是让 ThreadSync 接受有序作用域列表,跨作用域共用一个 ID 之前先用新分析器证明参与谓词恒成立,验收标准是资源声明回到 4、占用回到 16 块/SM,#3643 的修复没有被回退。11-05 在审的 #3787 继续补证明器盲区:192 线程下 tx < 160 and tx // 32 % 2 == 0 这种「边界项与 warp 索引取模项用 and 拼起来」的形态会误报,根因是 FloorDiv 与 FloorMod 的编码在「上界加取模」组合下建立不起两个除法的关系,修法是把 threadIdx.x 改写成 warp_index * warp_size + lane。来源:https://github.com/tile-ai/tilelang/pull/3733、https://github.com/tile-ai/tilelang/pull/3761、https://github.com/tile-ai/tilelang/pull/3787

  • TIRx 迁移从「口头理由」变成一周三处硬证据,底座换代进入公开准备(11-03 至 11-06):11-03 的 #3736 把 tilelang/language/ast/、tilelang/language/parser/ 与 tilelang/language/tir/entry.py 一并删掉——早期本地 TVMScript 解析器分叉的残留,新增 1 行、删除 3384 行、10 个文件,语言层改走上游 tvm.tirx.script.parser,AST 与语句覆写注册在 tilelang.language.overrides。11-06 合入的 #3752 与其前置的 tile-ai/tvm#91 把 Z3 的 rlimit 从写死的 100,000 变成可按 kernel 配置,作者的原话是「这间接修掉了很多与 Z3 证明器相关的问题」。11-06 的 #3803 提交完整迁移草案(draft):把 vendored 的 TVM 换成独立的 Apache TIRx,前端、IR 变换、代码生成、JIT 适配与打包全部迁到原生对象模型,并 pin 了 apache/tvm-tirx 的 7e49a148 与 apache/tvm-ffi 的 3e906bc6,依赖树保持干净上游、不注入补丁。三条线指向同一件事:主线这一阶段的第一优先级是换底座,其余优化都要给它让路。来源:https://github.com/tile-ai/tilelang/pull/3736、https://github.com/tile-ai/tilelang/pull/3752、https://github.com/tile-ai/tilelang/pull/3803

  • 「静默生成错代码」成为本周最主要的缺陷类型,且跨四个发布版本(11-03 至 11-05):最重的一条是 #3808——条件相关的标量化之后,共享内存 epilogue 少了一道 __syncthreads(),0.1.12、0.1.13、0.1.14、0.1.15 四个版本都会静默生成错误 CUDA,编译不报任何错,报告者给了完整可运行的复现脚本。同族还有 #3800(关掉 warp specialization 后,T.Pipelined 在 tcgen05 的 barrier 等待之前就重载下一 stage,MMA 读到被改写的 operand)、#3721(sm_100a 上对 TMEM 做 T.gemm,threads/32 不能整除 N 时 PTX 里一条 tcgen05.mma 都没有,乘积被静默丢弃)、#3775(生产环境 TP2 sglang 加 nvrtc 后端,CUDA graph 捕获被损坏,三次字节一致的配置复现出三个不同的首个失败点)、#1991 与 #1992(昇腾 910_9362 上 threads=2 时 AscendVidReduction 把注入的 tmp workspace 也无条件减半,sin/cos 作用在间接寻址 tile 上 3978/4096 个元素算错,即 97.1%)。这一族都不报错、都能通过常规测试,只能靠针对性复现抓到。来源:https://github.com/tile-ai/tilelang/issues/3808、https://github.com/tile-ai/tilelang/issues/3800、https://github.com/tile-ai/tilelang-ascend/issues/1991

  • TileFoundry 把同步从「分析能证明」推进到「发射必须与证明一致」(11-02 至 11-06):#283 把「共享内存访问应当如何排序」做成可证明的分析层——按 lowering 的发射顺序走每个函数,在冲突访问之前规划 commit、wait 与作用域内栅栏,83 个调度夹具里计划 52 个、拒绝 31 个(24 个用 TMA 拷贝、6 个用 wgmma,两者的完成语义尚未声明);因为默认没有调用方请求,75 个会 finalize 的夹具字节与 main 完全相同。11-04 合入的 #285 是它的执行端:CUDA lowering 按同步计划发射并逐条反向校验,不一致抛 LoweringError;单张 sm_89 上 7 个能构建的程序里 6 个与参考完全一致、synccheck 与 memcheck 全干净,而同样这 7 个按 main 的方式 lower 有 4 个结果错误,直接症状是 sm_89 的 E4M3 GEMM 最大偏 659。11-06 合入的 #286 补 WGMMA 的异步完成语义,11-05 合入的 #284 把 FSDP 与 HSDP 的训练语义做进 HIR 求值器并在四卡 sm90 上与真实的 2×2 PyTorch HSDP 对齐。作者对 #286 保留了明确的限定:矩阵指令发射器仍未接上,这笔不证明执行正确性、也不给性能。来源:https://github.com/tile-ai/TileFoundry/pull/283、https://github.com/tile-ai/TileFoundry/pull/285、https://github.com/tile-ai/TileFoundry/pull/286

  • TileOPs 五个窗口合入 15、7、1、7、9 笔,竞争点从「算子数量」彻底转到「选择器与准备成本」(全周):首窗 15 笔里 13 笔压在 11-01 一天,几乎每笔都带 H200 实测表与「manifest 无已测行回退」声明;随后的 11-04 与 11-05 两个窗口各有 7 笔与 9 笔,把节奏推回节后最密。三条硬数字可以概括这一周的方向:GQA 稠密前向用一条实测规则选 decode 或 prefill 程序,相对最快候选的选择遗憾从 1.818 降到 1.017(#2808);warp 专用 decode 把 softmax 与 QK、PV 重叠,受影响的 478 次扫描调用几何均值延迟为 main 的 0.952 倍(#2817);广播操作数就地读让 lerp、where、clamp、masked_fill 五类算子的广播行耗时降到主线的 0.13 到 0.44 倍(#2832)。同期还有 roofline 计价重写(三条物理不可能的带宽行从 7.48、7.50、5.30 TB/s 修正为 0.62、0.62、2.65 TB/s,#2830)与三阶段瓦片因子单独拟合(选择遗憾从 0.89% 降到 0.06%,#2836)。来源:https://github.com/tile-ai/TileOPs/pull/2808、https://github.com/tile-ai/TileOPs/pull/2832、https://github.com/tile-ai/TileOPs/pull/2830

  • 国产后端从「自研分支」走向「公共评审」,海光第一次开出正式 PR(11-05):海光仓主分支停在 11-02 的 #32(把 gfx948 按 gfx946 兼容处理),本周的动静先在 feat/hcu-scaled-cast 分支上攒了四笔(E8M0 缩放的 FP8/BF8/FP4 转换、64 bank 的 LDS 读相位、混合 LDS wrap 的至多两路证明、按物理读宽度做混合检查),11-05 又落三笔 LDS 相位修,并于 09:08 与 19:52 分别开出 #33(scaled cast)与 #34(reduction 的 XOR shuffle 用 DPP 与 lane swap)两个 PR——这是国产后端里第一个走完「自研分支到公共评审」这条路径的。mlir-ascend 同期把 BishengIR 工具链适配器搬进 contrib(#209),并用 #203、#211、#212 三个小 PR 把 tl.sub/mul/div/max/min、tl.shl、tl.shr 一个个搬进 TLIR 方言。摩尔线程在 11-05 落了一笔 README 文档提交,结束自 09-30 以来的完全静默。来源:https://github.com/tile-ai/tilelang-hygon/pull/33、https://github.com/tile-ai/tilelang-hygon/pull/34、https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

  • 硬件支持向两头同时扩:Blackwell 新 SKU 与国产加速器(全周):一头是新卡。B300 上报的 sm_103a 在两条路径上各卡一次——tcgen05.mma.kind::i8 只存在于 sm_100a 与 sm_110a,内核通过 lowering 后在 ptxas 报 Feature '.kind::i8' not supported(#3770);CuTeDSL 的 TMA 辅助函数维护一张架构白名单,B300 在进 DSL 之前就被拒(#3772,加上 sm_103a 与 gr100 的 sm_107a 后 B300 上 8 通过 1 跳过、gr100 上 13 通过 1 跳过)。SM120 上 FP8 加 fp32 累加的 GEMM 一直用 SM89 的 mma.sync 原子,换成带 ue8m0 单位缩放的 block-scaled MMA 后,RTX 5090 上 8192 立方从 463 升到 727 TFLOPS(在审 #3790)。另一头是国产卡:海光在 64 bank 的 LDS 相位上做微架构级调优,昇腾把 Cube 与 Vector 的核心比从固定 1:2 改成可配置(#3783)、把 AutoSchedule 的周期下界收紧使 Z3 求解时间降 10.04%(#3776)。来源:https://github.com/tile-ai/tilelang/pull/3772、https://github.com/tile-ai/tilelang/pull/3790、https://github.com/tile-ai/tilelang/pull/3783


二、核心项目进展(主仓 tilelang)

版本与发布

  • 主仓最新标签与发布仍为 v0.1.15(09-30 发布),全周无新标签、无新发布页条目,至本周五已满 37 天。
  • 组织内两个新版本都停在上周: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)、tilelang-musa v0.1.15+musa.1(09-30)、tilelang-mlir-ascend 停在 v0.1.2.020 系列。
  • 合入量与节奏:周一窗口为 72 小时(10-30 07:00 至 11-02 07:00),合入 35 笔;周二至周五四个 24 小时窗口分别为 18、8、12、18 笔,五窗合计 91 笔。在审更新依次为 82、44、48、42、57 笔,合计 273 笔;新开 issue 依次为 11、8、7、4、5 条,合计 35 条。主分支提交分布:tilelang 7、5、5、3、5 笔,TileOPs 15、7、1、7、9 笔,TileFoundry 3、1、1、1、1 笔,tilelang-ascend 5、4、0、1、2 笔,tilelang-mlir-ascend 1、0、1、2、0 笔,tilelang-hygon 主分支 2、1、0、0、0 笔(另有功能分支提交),tilelang-lsp 与 VS Code 扩展各 1 笔,tilelang-musa 1 笔。
  • 节奏上本周呈「周初清账、周中探底、周末回升」的形态:11-04 那个窗口只有 8 笔合入,是本周最低点,11-06 窗口回到 18 笔;在审队列则从周初的 82 笔一路降到 42 笔后于周五回升到 57 笔,说明周初积压的评审在本周被大量消化。

编译器、语言层与执行路径

  • 昇腾 950 FlashMLA 内核与配套的编译器能力(#3678,11-02 合入,20 文件 +1997/-17):见要闻第一条。值得注意的是这笔把「显式调度」所需的三项底层能力同时补进主线,而不是在昇腾分支里打补丁——L0C 到 UB 的拷贝保留 NZ 输出、每个 AIV 流在流水线之前选定、物理发射相位相同时保持源码顺序。
  • cast 解耦的向量宽度回归连修两笔(#3734 11-02 合入,#3759 11-05 合入):#3734 修 SM100 上 ViT 前向的 Cannot convert type float32x16 to CUDA type——向量化规划器为 256 位 BF16 存储选 16 个 lane 并推迟 cast 约束,而 #3579 让解耦跳过了读内存的守卫之后,最终合法化仍沿用推迟策略,残留的 FP32 表达式被撑成 512 位向量;修法是把 defer_cast_constraints 做成规划器的 C++ 参数而非 IR 属性,守卫的 FP32 转 BF16 降到 8 个 lane,成功解耦的 BF16 拷贝仍保留 16 个。#3759 收它的尾巴:条件里的 int64 cast 此前会把数据向量宽度从 16 lane 压到 4 lane,即使该条件在 16 个 lane 上完全一致;修法是把条件一致性分析复用到 cast 宽度豁免上,if 语句、tirx.if_then_else 与标量谓词的 Select 条件单独访问,子句里的 cast 仍受限。
  • ThreadSync 的三段式演进(#3733、#3761、#3787):见要闻第二条。#3733 是明确的临时措施,原文写明「为 TIRX 迁移让路」,配套的 TVM fork 优化提案 tile-ai/tvm#89 已关闭。#3761 的 401 行改动里最值得记的是归因实验:作者单独改编译产物里声明的屏障资源数、不动任何可执行指令,就复现了同样的变慢,把「性能回退」从观感变成了可指认的单一变量。
  • 语言层与前端做减法:#3736 删掉 3384 行遗留 TVMScript 解析器分叉(见要闻第三条),保留它的风险写得很具体——一旦有人直接 import tilelang.language.parser.operation,它会在 tirx.PrimExpr 与 tirx.IterVar 上注册互相冲突的算子;#3760 给嵌套 T.macro 引入有界的进程内转换缓存,只存与上下文无关的转换结果,宏实例再绑定当前的 globals 与闭包值,解决「kernel 已命中缓存但宏仍重复解析编译」的问题;#3802 新增 T.annotate_host_mapped_buffers,把被注解参数的索引写进 PrimFunc 的 tl.host_mapped_params 属性,运行时用 cuMemHostGetFlags 校验注册状态、用 cuMemHostGetDevicePointer 解析设备可见地址,普通可分页内存仍被拒。
  • 调试与可观测性:#3703 把 #line 指令从「每条 TIR 语句只锚定首行」改成分级发射(0 关闭、1 语句级、2 以上逐行重锚定),实现落在 CUDA、HIP、CPU C、C host、Metal、Ascend 共用的 C 系输出流里,新增 TILELANG_EMIT_LINE_DIRECTIVES 作为 Python 侧默认值;作者注明 CUDA 与 HIP 用例在本机被跳过、cuda-gdb 的实际行为未直接验证。同日 LSP 仓跟进一笔,把 CUDA 调试对齐到新的行映射;11-04 VS Code 扩展仓再跟进一笔,更新文档并把产物改名 vsix。
  • 自动调优的口径:#3546 把 set_autotune_inputs 的供给闭包从「调用点从不传入的参数」改为 get_current_device(),此前在没有 CUDA 的机器上兜底分支把每个捕获张量搬到 CPU、每一轮调优都因启动器设备不匹配失败;双卡 Ascend950 上三笔昇腾自动调优回归修前全败、修后全过。同族的 #3784 在 AscendC CI 上给跨设备拷贝探针前后各加一次设备同步,让异步错误落回 try ... except RuntimeError 的作用域内——此前它会晚一步浮出,看起来像是无关的 torch.arange() 报了 4 字节 memcpy 尺寸不符。
  • Z3 求解资源从写死变成可调(#3752,11-06 合入,依赖 tile-ai/tvm#91):rlimit 默认仍为 100,000,不配置时行为不变。这与同期昇腾侧收紧 AutoSchedule 周期下界(#3776,Z3 求解时间降 10.04%)是同一方向的两端:一个改搜索预算,一个改搜索起点。

正确性与缺陷修复

  • ThreadSync 与屏障:#3733 的整笔回退(见要闻)与 #3761 的跨域复用(见要闻)之外,在审的 #3765 把「单个线程索引边界守卫的部分屏障」在线程维是整 warp 数时显式接受,#3767 把多索引 fragment 访问的 StructuralEqual 内部断言换成带两个索引表达式的 ValueError,#3786 把 T.copy 目的区域越界改成编译期报错,#3788 修 ProveFragmentContains 的快速路径——它在大 fragment 完全复制时直接返回真,跳过了线程归属检查,于是「小 fragment 拥有线程 0 到 127、大 fragment 只拥有 0 到 95」被误判为包含成立。
  • 静默错代码一族:#3808、#3800、#3721、#3775、#1991/#1992(见要闻第四条)之外,本周还有 #3774(16×16 共享内存置换只做行位移、丢了列偏移,256 个元素错 223 个,而同样写法的 32×32 是正确的)、#3773(均匀 while 循环里做三次共享内存拷贝加 GEMM,串行对照 100 次 0 错、while 版本 100 次错 21 次并出现 10 种不同输出)、#3763(嵌套整数向下取整 (x // -2) // 4 在 int16、64 元素、32 线程下报 Let cannot bind the same var to two different values)。这三条与 #3775 构成「消费级卡 + WSL + pip 安装」这一组合的集中报障面。
  • 依赖分析与同步插入的根因:#3749 是本周技术含量最高的一条缺陷报告——两处连续 SIMD 加载丢掉了对更早一次 GM 到 UB 拷贝的依赖,根因在 MemoryAccessDetector::LogicalPointerRegion 重建出的逻辑区域漏掉了物理加载范围的一部分:最小复现里 work.shape 为 (2, 128),生成的读描述符覆盖物理 [64, 192),而重建出的区域是 [0, 128),于是依赖分析认为输入 DMA 与加载无关,AutoSchedule 把输入拷贝挪到加载之后、InsertSync 也不发 MTE2 到 V 的 set/wait。同一窗口主仓还新增 8 条指向昇腾 AscendC 与 PTO 两个后端代码生成与同步插入的 issue(#3740 到 #3749),其中 7 条由同一报告人提交,都附了可复现的 prim_func 与逐行对照的生成代码。
  • 边界条件一破就报内部错误:#3738 修「从比 padding 后共享 tile 更小的全局区域做 T.copy(例如 K=100 拷进 BK=112)被判成 TMA 整块拷贝后硬失败」;#3766 修 cta_group::2 的 TMA 字节计数——一对 CTA 的字节都落在偶数 CTA 的屏障上,而带 cluster 屏障的 tma_copy 原先只对 rank 0 期待 cluster_size × tile_bytes,在 (2,1,1) 与 (1,2,1) 上碰巧对,在 (4,1,1) 与 (2,2,1) 上错。
  • 算子与示例的边界:#3709 让 CPU 标量 GEMM 只清理目标区域,#3712 让 LegalizeNegativeIndex 跳过不可达分支;#3714 为 fp32 到 tf32 的转换补上就近偶数舍入(cvt.rn.tf32.f32,非饱和要求 SM90 以上、饱和要求 SM100 以上);#3702 让多个注意力示例从 shared_memory_per_block_optin 推 tile,RTX PRO 6000(sm_120)上 CUDA 9 项由「2 过 7 败」变为 9 全过,H100 有 227 KiB 可选上限、仍走原 tile。
  • 昇腾侧的两个语义偏差:#3762 报 T.simd.vcvt(src, "uint8", sat=False) 对 int32 向量静默忽略 sat=False 并永远饱和,文档写的回绕语义(-1 变 255、-128 变 128、300 变 44)实测成截断(-1 变 0、-128 变 0、300 变 255),影响落在 int8 量化尾声那类 f32 转 int8 的路径上;#3789 报 PTO 后端为整数 T.atomic_add 发射带符号的 pto.si32/si64 而 PTO 原子接口只接受无符号载体,当天就有 #3792 与 #3794 两个方向的修法。

运行时、CI 与工具链

  • CI 的两类稳定性问题都有处理:#3784 修 AscendC CI 的自动调优供给用例(见上),#2833 把 TileOPs 的 RUNNER_VERSION 从 2.336.0 提到 2.338.2(GitHub 废弃旧 runner 后 gpu-smoke 全部排队,而 entrypoint.sh 带 --disableupdate 无法自升级),#2827 把 TileOPs 的 off-CUDA 套件挪到夜间作业(原来 preflight 里这两步在 ubuntu-latest 上要 13 到 18 分钟,挪走后 preflight 只保留收集检查、丢弃测试报告与 cuda_only 必失败检查,约 3 分钟)。
  • 并行合并缺少 post-merge 门禁的实证:#2824 在 21:06 合入时把 grouped GEMM 的 roofline 行改成必须声明 sizes 并从它推导 G,#2830 在 21:08 合入、写的还是只有 G 没有 sizes;两笔相隔约 2.5 分钟、各自 CI 都是绿的,合并后 main 的 GPU Smoke 立刻报 RowError: row misses ['sizes'],当日以 #2837 补修。这是本周最值得记的一条流程记录:单笔 CI 通过无法覆盖「两笔语义相交」。
  • 文档与本地化:tilelang.github.io 在本周四个采集窗口内分别有 3、5、3、3 笔 Update docs 自动构建,内容对应 #line 分级映射、ThreadSync 作用域列表、host-mapped 注解、TF32 rounding 与 Z3 配置等接口变化;TileOPs.github.io 有一次推送,TileOPs-nightly 只有分支推送、主分支无新提交。
  • 昇腾侧的工具链细节:#1956 让 AI Core 异常转储兼容 CANN 9.2.0(补 acldumpTensorInfo 的必填字段、跳过空地址与零大小张量、把注册接口从已废弃的 aclrtSetExceptionInfoCallback 换成 aclrtExceptionInfoCallbackRegister,版本门槛从 9.3 下放到 9.2);#1932 把 A2/A3 的性能调优技能与 Ascend950 工作流结构对齐,并新增芯片类型识别与芯片信息检索两项技能。

评审队列观察

  • 体量最大的一笔仍是昇腾整仓迁移(#1975):11-04 00:34 更新时达到新增 607317 行、删除 97595 行、2887 个文件,把 TileLang-Ascend 搬进主线多后端架构,原有昇腾组件移入 ascend/a2a3 下,靠 arch="dav-2201" 与主线的 A5 路径分流。作者在描述里把待办写得很实在:三个目录下文件与原文件的差异尚待收敛、静态检查与格式化 CI 待恢复、文档与 Agent 相关待迁移。
  • TIRx 迁移草案(#3803)是全周最有分量的在审项:它决定后续 TVM 侧修复以何种节奏进入主线,也决定昇腾、海光与 mlir-ascend 的胶水层要不要跟着改。
  • TileOPs 的在审队列集中在选择器与布局:#2836 给 coop2_swap 的三阶段瓦片单独拟合因子(1.36)、#2838 修「把 K 当成连续维、于是否决 TN 布局 K=1 的合法调用」、#2839 让 FP8 分组 GEMM 用 TMA 加载 A 的缩放并给 K ≤ 1536 的非融合紧分组换 64×256 ping-pong 瓦片、#2840 补 SM89 的 prefill sinks 并把 sink 的有无纳入编译实例身份。
  • 在审数的走势:82、44、48、42、57 笔。周初的高位来自上周积压,周期中段降到 42 笔后于最后一个窗口回升到 57 笔,与合入量的回升同步,属正常的清账—开单节律。

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

昇腾(Ascend):从「适配目标」变成「首发目标」

  • 首发侧:FlashMLA 内核(#3678)以昇腾 950 为目标直接写在主线里,用 T.Stage、T.annotate_buffer_addresses 这类为显式调度准备的底层能力,而不是另开一个仓做移植;目标形状头数 64、Dqk 与 Dv 512、topk 为 64 的倍数。配套的批量搬运能力在 11-03 补齐:#3689 让 [B, M, N] 区域写成一次 T.copy、lowering 生成批量 ND2NZ 描述符,保留源的批与行 stride,转置批量走 DN2NZ,超出 GM 边界的源批被跳过;验证在 Ascend950DT 加 CANN 9.2 上,50 个 AscendC 示例测试与 1733 个昇腾单测通过,其中 28 个 BF16/FP32 设备用例覆盖批与 M/N 尾块、整批越界、非连续源 stride、目的存储大于源与矩形转置。
  • 同步正确性:#1981 修 copy_pipe_to_l1 的同步标注——A2/A3 的 PTO 跨核 V 到 C 通信经 GM 中转,三段搬运分属 AIV 的 MTE3、AIC 的 MTE2 与 AIC 的 MTE1,而它在 operation_config.h 里被标成 PIPE_MTE3,自动同步于是等的是 AIC 本核的 MTE3、而该队列为空、事件立刻置位,MTE1 读 L1 可能在 MTE2 写完之前放行;复现证据两条:把生成源码里的 PIPE_MTE3 改成 PIPE_MTE2 后 10/10 全部恢复精确,本地在最新主干上 10/10 次运行错误、单次最多 16384/16384 个元素不匹配。修法是拆名字而不是直接替换:copy_pipe_to_l1 改标 PIPE_MTE2,新增 copy_pipe_to_l1_bind 标 PIPE_MTE1(A5 直连路径),避免修好 A2/A3 却让 A5 出错。
  • 收敛到主线:#1971 的 RFC 给出同步实现的替换路线——移植主线同步核心与独立 Z3 证明服务,通过旧 TIR 适配层接入现有编译流水,保留 TVM、调度、内存规划、AscendC/PTO codegen 与运行 ABI,验收通过后再切默认实现并删除旧算法;配套的 #1965 把主线的 z3 版 insert sync 迁到 Ascend 仓;#1975 则是仓级骨架的迁移。三笔是同一条路径的三个层面。同期 #3752 让 Z3 的 rlimit 可按 kernel 配置、#3735 提出在 thread sync 里证明模块化索引的不相交性,正是从「配置求解资源」与「减少需要证明的形状」两头处理上周回退留下的证明预算问题。
  • 宿主开销这条线:#1948 把 CythonKernelWrapper.forward() 里每次调用都无条件构造的 TVM Analyzer() 改成惰性构造并按需复用,910B3 上 vec add 单块从 442 至 467 微秒降到 105.1 微秒、32 块从 441 至 493 降到 108.0、kpool 压缩从 447 至 450 降到 111.3,裸 ctypes 调用约 12.5 微秒基本不变;承接的 #1967 在 cython wrapper 里缓存符号形状的解析结果,继续压符号路径上的成本,同族还有在审的 #3757(按 Var 对象而不是打印名做键,避免同名不同变量被当成同一符号)。
  • 算子供给:#1827 新增 WY Fast 前向、#1828 迁移并补上 wy_fast_bwd_split 的分段反向(两笔在 11-02 02:33 前后相差不到一分钟合入),把线性注意力的这条算子链在昇腾上补齐到可训练;#1866 迁移 Chunk O 前向(chunk 门控线性注意力),用开发者模式加四项 pass 配置、单个 @tilelang.jit 内核加固定 20 核的 grid-stride 循环,做到片上融合、零 GM workspace、零宿主 ATen 算子;#1879 新增 MHA 前向示例(BSHD 布局、Expert 模式的三分派 MMA 流水,wrapper 不含任何 T.* 代码),cann-bench 20 个用例精度全过、对比 torch-npu 基线平均加速比 0.97 倍、综合得分 73.13(本地 910B2 口径)——0.97 倍意味着还没有净收益,但骨架给出来了;#1966 在审的 GLM-5.3 kpool fp8_mqa_logits 示例是双 kernel 加 AIV/Cube 分工,头数 32、维度 128、页大小 64。
  • 工程口径:#1954 把 T.tile.bitwise_rshift 在 int64/uint64 上的编译期段错误改成带实际 dtype 与支持列表的 ValueError(只允许 int16/uint16/int32/uint32),明确「不新增 64 位右移支持」;#3783 让 Cube 与 Vector 的核心比可显式配置或按设备推导(vector_core_num / cube_core_num,取不到退回 2),并把选定的组数一路带到 sid 域、任务分配、同步与代码生成;#3776 给 AutoSchedule 的启动间隔补单任务界、资源团界与自依赖界三条必要条件,用三者最大值初始化二分,昇腾 950PR 加 CANN 9.3.0 周更版上全部示例编译通过且调度结果不变、Z3 求解阶段累计耗时降 10.04%。
  • 回归状态:周一窗口日测 2970 项全过、11-01 有一次 workflow 失败(#1976);周二与周四的日测记录为 3030 项全部通过、通过率 100%(#1980、#1988)。用例规模一周内从 2970 涨到 3030,说明覆盖面仍在扩,但 #1991/#1992 与 #3762 这两个靠人工报告才发现的问题说明「通过率」不能当作正确性的替代指标。

海光(Hygon):从自研分支进入公共评审

  • 主分支:#31 合于 10-30,改进动态 MLS 边界推断与 MatrixStore 的 stride 处理、修 MMAC 与 blockscaled GEMM 的 clear_accum 下沉、在持续 WASP 路径改用 M 分区消费者与独立的 A/B 栅栏并保留 slot 与 phase 轮转;#32 合于 11-02 10:46,把 gfx948 当作 gfx946 兼容、复用既有的 MLS atoms、MMAC、barrier 与编译选项,目的是在接入 gather load 之前先让现有 builtins 在新型号上跑通。这是典型的两步走:先把新架构号挂进现有路径拿到可测基线,再补真正需要新指令的部分。
  • 功能分支:feat/hcu-scaled-cast 在 11-04 的八小时内落了四笔,分别是 E8M0 缩放的 FP8/BF8/FP4 转换(标量与打包两种 lowering,新增 223 行的 HCU scaled_cast.h 与 lower_scaled_cast.cc,主变换文件改动 650 行)、按 bank-word 几何推导 plain-copy 的 swizzle 并为窄 b16 tile 提供显式逆变换(12 文件 +1536/-112,新增 5 个测试文件)、混合 LDS wrap 的「至多两路」证明、混合 wrap 检查改用物理读宽度;11-05 又落三笔(默认 swizzle 适配 64 bank、MN 切片消费端恢复 swizzle 规划、K 分区的布局感知读)。
  • 评审状态:11-05 09:08 开出 #33(scaled cast,即分支主体)、19:52 开出 #34(reduction 的 XOR shuffle 用 DPP 与 lane swap 优化,外加「布局推断先传播已知约束再分配自由候选」的修法)。从「功能分支上攒了四笔、未开 PR」到两个正式 PR,是国产后端里第一次走完这条路径;主题集中在共享内存 bank 冲突与相位对齐,说明该后端已经从「能编译」进入「按微架构调优」的阶段。
  • 另有在审的 #19(标题只有 Distributed、描述为空,11-03 11:06 更新),本报告不对其实现细节作推断。

摩尔线程(MUSA)与沐曦(MetaX):一个恢复文档维护,一个仍只有上游同步

  • 摩尔线程:11-05 19:46 主分支落了一笔 docs: update README installation and operator development references,是本仓自 09-30 以来第一次提交;同期 6 个遗留 PR(依赖升级、TileOPs 子模块、MUSA 集成的证据门槛文档)时间戳被一并刷新,但内容仍是 04 月到 08 月的老提案。判断上宜保守:文档提交说明仓库有人在维护,不足以证明后端开发恢复,真正的信号要看后续是否出现算子或 lowering 提交。
  • 沐曦:日报在五个窗口里都记为「无推送」,168 小时窗口复核时确认其主分支本周确有提交,但逐笔看全部是上游 tilelang 的同步(#3760、#3752、#3759、#3802、#3714、#3784、#3772、#3534、#3761、#3689、#3546、#3736、#3758、#3748、#3678、#3734、#3702、#3703 等),作者均为上游维护者,未见沐曦侧自有提交。因此准确的表述是:主分支在上游同步,自有开发仍处于静默期,已连续第三个窗口。dev 分支最新仍停在 10-16。

mlir-ascend 与 TLIR:把语义一层层落进方言

  • 结构整理:#206 合于 10-31,把 CI 里指向内部缓存服务的内容清掉,属可复现性清理;#209 合于 11-04 16:35,把 BishengIR 工具链的适配代码搬进 contrib,把「随上游工具链演进而需要频繁改动的胶水层」与核心 lowering 分开,后续升级 BishengIR 时改动面收敛在 contrib 内。
  • Agent 工作流:#208 合入 harness 分支(不是主线,引用时不宜写成「合入主线」),新增 2715 行、66 个文件,给算子 harness 加第一套 Codex 工作流:项目初始化时用迭代提问、把已有算子仓接入、从持久化上下文里恢复开发、调优或分析;三套技能共享 .workspace/ 下可编辑的 Markdown 记录,任务交接区分「当前候选」与「已验证的最佳候选」,保留失败与回退过的实验。作者明确本次没有真实 NPU 实测,NPU 执行适配与 npm 分发列为后续工作。
  • 方言算子:#203 给 TLIR 加 tl.sub/mul/div/max/min 并补 TL 到 HIVM 的 lowering,#211 加 tl.shl、#212 加 tl.shr,三笔在 11-05 的 01:42、19:44、20:58 先后推进。这条线节奏慢但方向单一:把 TileLang 的语义一层层落进 MLIR 方言。

四、生态与采用方

TileOPs:0.1.0 之后的第三周,从调优转入选择器工程

  • 合入节奏:首窗 15 笔(13 笔压在 11-01),随后四个窗口为 7、1、7、9 笔,全周合计 39 笔。11-02 那批出现了测试基建(#2816 的 tileops_test 目标,1281 行、59 个文件,让 tests/ops 的冒烟测试不再必须站在 CUDA 上)与 CI 契约(#2818)两类非性能改动,说明算子库开始为多后端与可维护性做结构性准备;11-04 与 11-05 两个窗口回到纯性能。
  • 注意力:#2808 把 GQADenseFwdOp 与 GQADenseWithLSEFwdOp 的程序选择收成一条规则 AttentionCall.dense_decode_program,输入是 query 与 key 长度、batch、头分组、head dim、RoPE 与 SM 数,2 至 64 个 query 位置走 warp 专用 decode;相对最快可用程序的选择遗憾(几何均值与最大值)从 1.818 与 25.2 降到 1.017 与 1.81,并修掉主线上「score scale 乘上 log2(e) 后在 FP32 下消失导致 NaN」与「D256 多位置调用在 SM90 上因共享内存不足启动失败」两个真实错误。#2817 把 warp 专用 decode 的 softmax 与 QK、PV 重叠(概率留寄存器、K 与 V 用独立 barrier、共享三槽环形缓冲),H200 上 (16, 8, 4096, 64, 8, 128) 的 float16 为 0.074610 毫秒对 fa3 的 0.077490(1.039 倍,3.654 TB/s),bfloat16 为 0.074260 对 0.077220(1.040 倍,3.671 TB/s)。#2807 把分页 MHA decode 扩到「一个块的共享内存装得下」的每个 head dim,#2797 从 GQAPagedFwdOp 删掉旧的 paged decode kernel。
  • MoE 与 GEMM:#2815 让 tight 布局的 MoE 输入按 token 收集、同一条 hidden 行在它的多条路由之间复用,(4096, 256, 8, 7168, 2048) 上 bfloat16 为 7.4417 毫秒对 flashinfer-cutlass 的 7.6982(1.034 倍,387.844 TFLOPS),float16 为 7.5040 对 7.9869(1.064 倍,384.624 TFLOPS)。#2805 让填满至少四分之三 SM 的部分一波 W4A8 prefill 栅格改成一 CTA 一 tile 而非 stream-K,llama-8b-up-b33 从 23.52 微秒降到 17.98(vLLM 20.42),清单外 M 33 至 128 的 44 组 N/K 扫描中改调度形状快 5% 至 36%、中位 18%、无一变慢。#2834 给小 M 的 NT GEMM 造了一个宽度能填满一波的原生宽度瓦片,decode-down(128×7168×2048)在主线上只用 132 个 SM 里的 112 个,改后领先最快对比项 3%、gate-up-b32 比主线快 1.4%。
  • 逐元素与归约:#2814 让 bfloat16 的 Sigmoid 与 16 位 SiLUAndMul 改用四分之一缩放的倒数、绕开带范围检查的除法,#2813 让 8 字节操作数在长启动下按一线程一向量处理(先用 TileFoundry 的 roofline 分析确认这类按位或算子全部是访存受限);#2832 让 lerp、where、clamp、masked_fill 五类算子不再把 stride 0 的广播张量先展开成输出形状,广播行耗时降到主线的 0.13 到 0.44 倍;#2829 让「体内最多两条便宜指令、读写总量不超过 L2 的四分之三」的一元逐元素缩成 256 线程的一波,并顺手修掉 ReLU 对 NaN 输入返回 0 的老问题;#2828 给按行分块的 ArgMax/ArgMin 在「一波驻留」前提下让每线程抱八个 16 字节向量,(2048, 4096) 的 bf16 ArgMin 从 8.704 微秒降到 7.393(FlagGems 8.543)。#2835 让归约 kernel 的轴置换与输出形状每个实例只算一次,六组 FP32/BF16 VarMeanFwdOp 配置的完整热调用耗时下降 4.33% 到 6.17%,生成的 CUDA 与主机源码完全一致——省的是调用准备成本,不是 GPU 时间。
  • 精度这一类缺陷被正面修复:#2811 给每个由行构成的 fp32 量加一条不变式「只被行的扩散度约束、不被偏移量约束」,实现上取行内某个元素作基准再相减(warp 主导的核用 warp 首元素、分片核用行的首元素),消掉靠近 1e36 的溢出;#2821 让方差合并从第一个真实的部分量初始化而非空累加器,常量输入的方差不再出 NaN(tests/ops/test_reduce.py 从 155 增到 158)。这类「空累加器 + 大偏移」的数值坑只在保留轴前后同时归约的形状上暴露,普通整轴归约测不到。
  • 度量口径成为产品特性:#2830 把 roofline 按「这次调用实际读了多少」重写——GQA 的稠密、varlen、分页前向只对「有 query 能看到的 key 行」计 K/V、对 key 实际到达的截断位置计 RoPE 表,GroupedGEMM 通过新的 grouped_gemm_roofline 按 batch_sizes 给某行的组计价,修正后三条物理不可能的行从 7.48、7.50、5.30 TB/s 变成 0.62、0.62、2.65 TB/s;#2822 把 bench_kernel 的 L2 刷写从「对 int8 缓冲做 zero_()」改成「读一个私有 int32 缓冲,大小是报告 L2 的两倍」,因为旧 flush 会留下约 120 MB 脏行、其回写落进下一次计时调用(2048×11008 f32 的 Sum 从 27.81 降到 21.38 微秒)。
  • 选择器与代价模型:#2825 把 W4A16 的 _select_config 改成在 _legal_configs 上取 _config_cost 的 argmin,删掉 tile 优先、最宽 K 的并列窗口与 joint_switch_margin 三条例外规则,新增 stage_fill 项;#2826 把稠密选择器的标定契约写进文档串(扫描、拟合、检查三步,含软退火与留出检验);#2836 给 coop2_swap 的三阶段瓦片单独拟合因子(1.36),regret 从 0.89% 降到 0.06% 几何均值、最大从 5.04% 降到 0.99%。
  • 发布形态与测试口径:#2816 的 tileops_test 把冒烟测试从 CUDA 上摘下来,作者同时把边界写清楚——比对共享参考只能证明后端无关的契约与可移植性,不能证明真实后端的数值正确性,那条线仍挂在 #2611;#2827 把 off-CUDA 套件挪到夜间;#2833 升级 runner。

TileFoundry:同步分析、训练语义与命名治理

  • 同步:#283 的分析层与 #285 的 lowering 执行端(见要闻第五条)之外,#286 在 11-06 13:40 合入,为每种操作数效应声明完成条件,把 WGMMA 的 fence、commit、wait 挂到需要它们的 HIR 节点上,lowering 在指令的发射点消费这些步骤,若后续作用域放不下集体等待就在发射作用域里完成;同时修正寄存器 Reshard 的别名,让元素在持有副本的每个线程上留在同一个寄存器。作者保留两条风险声明:矩阵指令发射器仍未接上,这笔不证明执行正确性、也不给性能;TMA 的完成语义另行处理。11-06 20:12 另开 #288,处理「已完成的部分归约」在分析中的计价。
  • 训练语义:#284 在 11-05 16:36 合入,此前的求值器把 reshard 当全局张量、漏掉不同参与者的贡献差异;本笔按参与者求局部值、GPU 坐标、调用与循环、gather/scatter 以及沿轴的 Partial(sum) 归约,让 Reshard 的别名证明对归属敏感。验收:核心套件 775 项、后端套件 778 项(两卡 sm90 上 13 项跳过),十组 HIR 配置在 CPU 上通过,两种驻留策略在四卡 sm90 上与真实的 2×2 PyTorch HSDP 跑两个更新步、本地梯度、参数与动量全部对齐。
  • 调度与发布:#280 与 #282 把「每次调用发射什么」「由哪个 mesh 发射」「窗口怎么切」的决策收进 analysis/emission.py 一处,验收是字节级对照——83 个调度夹具里 75 个 finalize 输出与主线完全一致,其余 8 个在主线上同样失败;站点仓在 10-30 发布 0.3 版调度文章《Scheduling, Designed for Agents》,把「面向 agent 的调度」从代码与发布说明搬到公开叙述。
  • 命名治理:#287 是一条文档向 issue,指出代码库里「一个概念五个名字」——mesh 的一个成员被叫作 placement 133 处、program 110 处、instance 92 处、participant 44 处,而 unit 这一个词又同时指空类型、硬件计算单元与 mesh 成员,Python 与 CUDA 运行时对 local_layout 的签名还不一致。这类 issue 通常不会立刻改代码,但它是编译器项目从「能写」到「能被人长期读」的必经一步。

TileSight:性能模型从论文走向仓库

  • tile-ai/TileSight(24 星)是 07-24 arXiv 论文《TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters》的开源实现:从瓦片形状、数据移动、软件流水与硬件档案出发,在 CPU 上估计算子代价、缓存流量与执行时间。11-05 出现仓库首个 PR(#1):H20 是 Hopper/GH100,带 requires_arch=("wgmma",) 的计划原先因为架构档案没有声明 WGMMA 支持而抛 UnsupportedError,本笔把 support_wgmma=True、并显式保持 support_utcmma=False,于是这类计划被接受、而 tcgen05 与 Tensor Memory 需求仍被拒。验证是 CPU 上 847 项全过,作者明确声明「不是 GPU 性能验证」。
  • 这条线的意义在于与 TileOPs 的 roofline 计价、TileFoundry 的分析层形成三角:一个给解析模型,一个给逐行计价,一个给可证明的同步计划。三者都在本周有实质动作。

采用方:TileKernels 与 FlashQLA

  • 两个采用方仓全周都没有提交:deepseek-ai/TileKernels 停在 10-14 05:49,QwenLM/FlashQLA 停在 10-30 03:02(最新发布仍是 09-30 的 v0.1.3)。
  • 但采用方不推送不等于不被影响:11-04 的 #3761 修掉的那笔屏障占用回退,触发点正是 TileKernels 的 mhc_pre_big_fuse 内核——同一个内核在 B300、隐藏维 1280 上丢掉了 25% 的屏障占用预算。这说明采用方的内核已经成为编译器侧的回归探针:TileKernels 不推送代码的这几周,它的内核仍通过主线的改动在被动地「被测试」。对跟踪生态的人来说,这是一个值得留意的反向信号。

社区分支与第三方仓

  • RISC-V 研究链成形:社区研究者 zishu0131-png 在 10-31 同日整理出三个仓——tilelang-riscv 作总入口(说明研究目标、两段编译链路与版本组合),numba2tilelang 把 Numba CUDA Python 静态解析后转成 TileLang 源码,tilelang2rvv 把受支持的真实 tvm.tirx.PrimFunc 转成 RVV intrinsic C、再经 Clang 编成 RV64GCV ELF 在 QEMU virt 上以单 hart 运行。作者把边界写得很清楚:不是 TileLang 官方注册的后端,依赖 threadIdx 的寻址、跨 hart 同步与未支持的 IR 一律拒绝,block 顺序执行,研究环境为 TileLang 0.1.15、Clang 23.1.3、QEMU 11.1.0。168 小时窗口的四源巡检在本周无新增变化。
  • PTO-ISA/tilelangGym 独立成仓:这套敏感性与回归测试集从 PTO-ISA/tilelang 的 examples/ascend/t_parallel_st_suite 拆出,理由是让测试集能独立钉住 TileLang 与 PTOAS 的版本;当前三条主路径是昇腾 SoftLib SimtVF(target=ascend、cython,主回归门)、VectorizeParallelToPTO 与 VL 的 vload/vstore、以及 benches 下用 @pto.jit 写的手工 SIMD DSL 对照实现,每个用例目录带 config.toml 与 VERDICT.md。
  • 第三方小仓:sablin39/tilelang-cuda-skills(149 星,给编写 TileLang 与用 CUDA 工具链调试提供技能集)、xwhzz/tilelang-hexagon(2 星,把 TileLang 往高通 Hexagon DSP 上带)、Index4NekoMeteor/Tilelang-AttnRes(1 星,用 TileLang 实现 Kimi K3 的注意力残差结构)、Tianchen-ai/tilelang-weft-bench(0 星,CPU 侧算子基准的统一引擎)、flytiger-eco/tilelang-wrapper-for-sail(0 星,暂无描述)。前三者方向有意思——一条做工具链侧的开发体验,两条做新目标架构与前沿模型结构的移植;后两者很早期,先记入观察名单,不宜当技术进展写。

五、社区、教程与活动

  • 报障面在安装形态上集中:11-03 一天 7 条新 issue 里,四条来自同一位使用 pip 包、跑 WSL、用消费级 sm_89 卡(RTX 4060 Laptop)的外部报告人,全部附了可复现脚本与完整输出;技术含量最高的 #3775 则是生产环境里「pip 包 + 双 H200 NVLink + 多 CUDA 上下文 + nvrtc 后端 + CUDA graph 捕获」的组合,报告人把三次运行的对照条件写全了——三次字节一致的配置复现出三个不同的首个失败点,/proc/environ 与干净对照的差别只有 TILELANG_EXECUTION_BACKEND=nvrtc 一个环境变量。这些不是编译器的深层缺陷,却决定了新用户第一次尝试时的体验。
  • 昇腾侧出现「报告人即评审人」的闭环:#1991/#1992 的报告者是 #1897 的 review 方,而 #1897 的修法只覆盖了 count 侧语义(HasModifiedBuffer 跳过 workspace 槽位),workspace 侧仍未处理——同一处布局逻辑的另一个侧面在两周后才被指出。这说明昇腾后端的布局不变量还没有形成可检查的书面契约。
  • 文档与本地化链路正常:tilelang.github.io 全周多笔 Update docs 自动同步,TileOPs.github.io 一次推送;LSP 仓把 CUDA 调试对齐到新的行映射(71e762e0),VS Code 扩展仓更新文档并把产物改名 vsix,接在 11-02 的「把 CUDA 调试对齐到 TileLang 行映射」之后——上一期主线合入的逐行锚定改变了生成代码的行号来源,扩展侧的调试配置需要跟着改。
  • 教程与知识资产:本周没有新的教程仓出现,社区投入集中在既有仓的维护与上面提到的两个 RISC-V 方向的小仓;PTO-ISA/tilelangGym 的独立成仓是测试资产第一次被单独版本化。
  • 新闻侧全周零实质命中:五个采集窗口的 Google News 中英文查询(覆盖精确词、GPU kernel、中文算子与编译器、昇腾与沐曦/海光/摩尔线程)里,最新一条相关报道仍是 10-21 的 SemiAnalysis《The Huawei-DeepSeek Symbiosis》,其余集中在 09-30 那一波「DeepSeek 开源昇腾版 TileLang、DeepGEMM、DeepEP」的报道,全部早于本周期;Hacker News 以 tilelang 检索 24 小时为 0 条;arXiv 最新命中是 11-03 的《Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving》(引用 TileLang 的工作,窗口外)与 07-24 的 TileSight 本体。报告主体因此全部来自代码侧,这已是连续多期的稳定形态。
  • 产业侧的口径冲突值得留一笔:11-03 的窗口里,Wccftech 称 DeepSeek 将在昇腾 950DT 构成的 Atlas 950 SuperPoD 上从零训练 V5,而同期 The Information 的报道是「最早 2026 年底开始在 950DT 上做预训练」,Bloomberg 与金融时报的早期报道则称 950DT 的 16 万片订单主要用于推理、此前在昇腾上训练受挫后回退到 Nvidia。三种说法对「是否以及何时用于训练」并不一致,且均未获华为或 DeepSeek 官方确认;本报告只作背景记录,不据此推断昇腾侧投入的落地节奏。

六、趋势观察

第一,TIRx 迁移已经是主线事实上的第一优先级,而且它开始以「减法」的形式出现。 一周之内三处证据:11-01 为迁移回退整套 ThreadSync 正确性改进(#3733,同时关掉 TVM fork 侧的优化提案 #89),11-03 删掉 3384 行遗留解析器分叉(#3736),11-06 提交完整迁移草案(#3803,pin 两个上游提交、明确不注入补丁)。中间还夹着 Z3 资源从写死到可配置(#3752 与 tvm #91)。这套动作的共同点是把复杂度从关键路径上撤下来:宁可暂时退回保守实现,也不让 vendored 依赖的内部结构继续长在语言层里。代价是主干在迁移窗口里会有一批「已知不最优」的路径,收益是底座换代完成后所有下游后端能对着一个稳定的 IR 边界做适配。

第二,正确性风险已经换了形态:从「编译不过」变成「编译通过、运行不报错、结果错」,而且跨版本潜伏。 本周至少有七条属于这一族:#3808 跨 0.1.12 到 0.1.15 四个版本少一道 __syncthreads(),#3800 在 Blackwell 的 tensor memory 路径上让 MMA 读到被覆盖的 operand,#3721 在 sm_100a 上静默丢弃整个 GEMM 乘积,#3775 在生产环境损坏 CUDA graph 捕获,#3749 让依赖分析看不见物理加载范围,#1991/#1992 在昇腾 threads=2 下让 sin/cos 算错 97.1% 的元素,#3774 与 #3773 在消费级卡上给出错结果。这一族对用户的伤害最大,因为它把「编译器对不对」这个前提变成了需要验证的假设。本周的应对方向是把正确性做成可拒绝的证明——TileFoundry 的 #283/#285 是先证明再发射、发射后反向校验,#3786 与 #3787 是把默默生成错代码改成编译期报错或明确拒绝——而不是继续加单元测试数量。

第三,昇腾的角色从「适配目标」变成「首发目标」,工程化在两条腿上同时推进。 一条是架构:FlashMLA 这种 DeepSeek 自家解码路径上的核心内核直接以昇腾 950 为目标写在主线里,用的还是 T.Stage、annotate_buffer_addresses 这类为显式调度准备的底层能力;配合编译器侧三处配套改动,昇腾路径正在从「能跑」转向「能按调度意图跑」。另一条是骨架:#1975 用 2887 个文件、60 万行新增把整个 TileLang-Ascend 搬进主线多后端框架,靠 arch="dav-2201" 做后端分流;mlir-ascend 的 #208 把 Codex 的初始化、接入、恢复三套工作流和任务交接写进 .workspace/。把「迁移仓」与「用智能体迁移仓」放在同一周看,能看出昇腾侧的瓶颈已经从「缺算子」转向「缺一致的工程骨架与可复用的迁移流程」。

第四,算子库的竞争点彻底转到选择器质量与准备成本,性能差距的剩余空间不在「怎么算」。 11-05 的九笔 TileOPs 合入里,真正改 GPU kernel 的只有广播就地读与原生宽度瓦片两笔;reduction 布局缓存省的是调用准备(4.33% 到 6.17%,生成源码不变),三阶段瓦片因子改的是一个常数(regret 从 0.89% 降到 0.06%)。再加上 roofline 计价重写与 L2 刷写改只读,可以看到一条清晰的路线:先把每个形状的合法配置枚举出来,再用统一的代价模型排序,标定过程写成可复现的契约(#2826 已经把「扫描、拟合、检查」三步写进文档串)。这条路做对了,性能就不再依赖个别内核的手工调优;做不对,就会退化成「某些形状特别快、其他形状靠运气」。

第五,国产后端从「一次性合入一个大后端」转向「把可评审的最小单元推到公共仓库」。 海光从自研分支走到 #33 与 #34 两个正式 PR,摩尔线程以一笔 README 结束两个月静默,mlir-ascend 用三个小 PR 把 tl.sub/mul、tl.shl、tl.shr 一个个搬进 TLIR 方言,昇腾则把核心比、异常转储、性能技能这类工程细节逐项补齐。三条线的共同点是不再等一个大合入:这对上游是好事,评审者能早就布局与 lowering 的设计提意见,而不是在三万行一次性合入时才看;对下游也有代价——主干在这类窗口里会处于「已知有几处不工作」的状态,需要更强的发布前回归纪律。

最后一条风险面:底座切换会把所有胶水层的成本一次性暴露出来。 如果 #3803 推进,3rdparty/tvm 的 vendored 结构会被替换成两个 pin 的上游提交,届时所有依赖 TVM 内部 API 的后端(昇腾、海光、mlir-ascend 的胶水层)都要跟着改。草案里明确写了「不注入依赖补丁」,这既是承诺也是约束——本周 mlir-ascend 把 BishengIR 适配器搬进 contrib、昇腾把性能技能按芯片代际分叉,都可以看作是为这类底座切换预留的可替换面。判断这条线是否健康,看的是迁移草案合入后两周内,这几个后端能不能在不动核心 lowering 的前提下完成适配。


附录:素材与核查说明

信源核查表

信源 核查结果
GitHub tile-ai 组织逐仓(29 仓库) 五个采集窗口依次有 10、9、9、8、10 个仓库推送;168 小时窗口复核为 14 个仓库有推送
合入真值检索(is:pr is:merged) 五个窗口 35、18、8、12、18 笔,合计 91 笔;首窗为 72 小时(10-30 07:00 至 11-02 07:00)
在审真值检索(is:pr is:open updated) 五个窗口 82、44、48、42、57 笔,合计 273 笔
新增 issue 核对 五个窗口 11、8、7、4、5 条,合计 35 条
主分支提交级核查 tilelang 7、5、5、3、5 笔;TileOPs 15、7、1、7、9 笔;TileFoundry 3、1、1、1、1 笔;tilelang-ascend 5、4、0、1、2 笔;tilelang-mlir-ascend 1、0、1、2、0 笔;tilelang-hygon 主分支 2、1、0、0、0 笔;tilelang-lsp 与 VS Code 扩展各 1 笔;tilelang-musa 1 笔
发布页与标签 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)
采用方 TileKernels 停在 10-14;FlashQLA 停在 10-30,最新发布仍为 09-30 的 v0.1.3
TileLang×RISC-V 四源巡检 168 小时窗口无变化
第三方与生态仓 研究链三仓、tilelangGym、cuda-skills、hexagon、AttnRes、weft-bench、wrapper-for-sail、TileSight 均有推送或新建
Google News(中英四至六组) 五个窗口均无窗口内实质命中,最新相关报道为 10-21 的 SemiAnalysis 一文
Hacker News 以 tilelang 检索 24 小时为 0 条,抓到的是与主题无关的 Show HN
arXiv 最新命中为 11-03 的 KV 压缩论文(引用 TileLang,窗口外)与 07-24 的 TileSight 本体

主要引用

说明

  • 本报告事实、版本号与日期均取自周期内五期日报原文及其一手来源,未作推测性补充;五个采集窗口首尾相接(10-30 07:00 至 11-06 07:00 共 168 小时),与上一期周报无重叠,故五窗的合入、在审与新增 issue 数直接相加。
  • 合入数、在审数与 issue 数来自 GitHub 检索的实测输出;主分支提交数为逐仓 atom 核查的结果,两者口径不同(一次合入可能对应多个提交,非默认分支的提交不计入合入数),附录按两类分别列出、不加总。
  • TileOPs 的「15 笔」出现在周一那个 72 小时窗口内,其中 13 笔压在 11-01,引用时不宜写成「周一单日 15 笔」。
  • #1991 与 #1992 标题与正文一致,按「同题两条」记录,未合并为一条;海光功能分支的七笔与摩尔线程的文档提交按分支级动态表述,不代表主线已合入。
  • 沐曦在五个窗口均记为「无推送」,168 小时复核发现主分支有提交,但逐笔核对全是上游 tilelang 的同步(作者为上游维护者),故表述为「主分支在上游同步、自有开发仍处静默期」。TileSight 与 TileFoundry 的 WGMMA 改动均由作者自述限定为能力元数据或同步计划,报告沿用其限定。