報告週期: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):週一 #3733 把三週前合入的精確執行緒計數(#3291)整筆回退,理由寫得很直白——那套改動把昂貴的計數查詢搬到滿塊柵欄檢查之前,即使所有執行緒都會到達柵欄也要跑證明,TVM 的求解預算被耗盡後返回部分計數,合法的 remap 與 sequence 輔助 kernel 直接以「Cannot prove a thread rectangle」編譯失敗;商守衛、精確參與執行緒界、非矩形執行緒集三類正確性改進一併撤回,#3732 留作遷移後重新最佳化的跟蹤單。週四 #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 的修復沒有被回退。週五在審的 #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):週二 #3736 把 tilelang/language/ast/、tilelang/language/parser/ 與 tilelang/language/tir/entry.py 一併刪掉——早期本地 TVMScript 解析器分叉的殘留,新增 1 行、刪除 3384 行、10 個檔案,語言層改走上游 tvm.tirx.script.parser,AST 與語句覆寫註冊在 tilelang.language.overrides。週四 #3752 依賴的 tile-ai/tvm#91 與週五的 #3752 把 Z3 的 rlimit 從寫死的 100,000 變成可按 kernel 配置,作者的原話是「這間接修掉了很多與 Z3 證明器相關的問題」。週五 #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 完全相同。週四合入的 #285 是它的執行端:CUDA lowering 按同步計劃發射並逐條反向校驗,不一致拋 LoweringError;單張 sm_89 上 7 個能構建的程式裡 6 個與參考完全一致、synccheck 與 memcheck 全乾淨,而同樣這 7 個按 main 的方式 lower 有 4 個結果錯誤,直接症狀是 sm_89 的 E4M3 GEMM 最大偏 659。週五合入的 #286 補 WGMMA 的非同步完成語義,週三合入的 #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 無已測行回退」宣告;週四 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 的至多兩路證明、按物理讀寬度做混合檢查),週五又落三筆 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 筆。
  • 節奏上本週呈「周初清賬、週中探底、週末回升」的形態:週三只有 8 筆合入,是本週最低點,週五回到 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 除錯對齊到新的行對映;週四 VS Code 擴充套件倉再跟進一筆,更新文件並把產物改名 vsix。
  • 自動調優的口徑:#3546 把 set_autotune_inputs 的供給閉包從「呼叫點從不傳入的引數」改為 get_current_device(),此前在沒有 CUDA 的機器上兜底分支把每個捕獲張量搬到 CPU、每一輪調優都因啟動器裝置不匹配失敗;雙卡 Ascend950 上三筆昇騰自動調優迴歸修前全敗、修後全過。同族的 #37

84 在 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 筆後周五回升,與合入量的週末回升同步,屬正常的清賬—開單節律。

三、多後端適配(昇騰 / 海光 / 摩爾執行緒 / 沐曦)

昇騰(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 筆。週二那批出現了測試基建(#2816 的 tileops_test 目標,1281 行、59 個檔案,讓 tests/ops 的冒煙測試不再必須站在 CUDA 上)與 CI 契約(#2818)兩類非效能改動,說明運算元庫開始為多後端與可維護性做結構性準備;週四與週五回到純效能。
  • 注意力:#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)。
  • 但採用方不推送不等於不被影響:週四 #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 星,暫無描述)。前三者方向有意思——一條做工具鏈側的開發體驗,兩條做新目標架構與前沿模型結構的移植;後兩者很早期,先記入觀察名單,不宜當技術進展寫。

五、社群、教程與活動

  • 報障面在安裝形態上集中:週三一天 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 本體。報告主體因此全部來自程式碼側,這已是連續多期的穩定形態。
  • 產業側的口徑衝突值得留一筆:週三視窗裡,Wccftech 稱 DeepSeek 將在昇騰 950DT 構成的 Atlas 950 SuperPoD 上從零訓練 V5,而同期 The Information 的報道是「最早 2026 年底開始在 950DT 上做預訓練」,Bloomberg 與金融時報的早期報道則稱 950DT 的 16 萬片訂單主要用於推理、此前在昇騰上訓練受挫後回退到 Nvidia。三種說法對「是否以及何時用於訓練」並不一致,且均未獲華為或 DeepSeek 官方確認;本報告只作背景記錄,不據此推斷昇騰側投入的落地節奏。

六、趨勢觀察

第一,TIRx 遷移已經是主線事實上的第一優先順序,而且它開始以「減法」的形式出現。 一週之內三處證據:週一為遷移回退整套 ThreadSync 正確性改進(#3733,同時關掉 TVM fork 側的最佳化提案 #89),週二刪掉 3384 行遺留解析器分叉(#3736),週五提交完整遷移草案(#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/。把「遷移倉」與「用智慧體遷移倉」放在同一周看,能看出昇騰側的瓶頸已經從「缺運算元」轉向「缺一致的工程骨架與可複用的遷移流程」。

第四,運算元庫的競爭點徹底轉到選擇器質量與準備成本,效能差距的剩餘空間不在「怎麼算」。 週五 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 本體

主要引用

com/tile-ai/tilelang/pull/3783>

說明

  • 本報告事實、版本號與日期均取自週期內五期日報原文及其一手來源,未作推測性補充;五個採集視窗首尾相接(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 改動均由作者自述限定為能力後設資料或同步計劃,報告沿用其限定。