TileLang 每日動態報告(2026-11-06)
調研視窗:2026-11-05 07:00 ~ 2026-11-06 07:00(北京時間,24 小時;上一期截點為 11-05 07:00)。
信源:GitHub tile-ai 組織逐倉核查(29 倉庫,視窗內推送 10 個)+ 合入與在審真值檢索(視窗內合入 18 筆、在審更新 57 筆)+ 新增 issue 5 條 + 主分支逐筆提交核查(tilelang 5、TileOPs 9、TileFoundry 1、tilelang-ascend 2、tilelang-musa 1)+ 非預設分支核查(tilelang-hygon 的 feat/hcu-scaled-cast 分支 3 筆)+ 釋出頁核查(tilelang / TileOPs / TileFoundry 三倉,視窗內無新版本)+ 採用方兩倉 + 文件站與夜間倉核查 + Google News 中英文四組查詢(走代理)+ Hacker News + arXiv。RISC-V 四源專項巡檢本視窗無變化。
本期索引
- 今日重點:TileOPs 一天九筆合入,把「廣播就地讀、單波次啟動、瓦片選擇」三件事同時推進;主倉五筆合入之外,一條從 0.1.12 起就存在的靜默錯誤 CUDA 缺陷被報出;昇騰修掉一個段錯誤,同時被指出 threads=2 下 sin/cos 有 97.1% 元素算錯
- 一、核心專案進展
- 1.1 主線:TF32 轉換支援 round-to-nearest-even,SM90 與 SM100 分開診斷(#3714)(11-05)
- 1.2 主線:host-mapped 緩衝可顯式註解,CPU 張量不再被裝置檢查拒絕(#3802)(11-05)
- 1.3 主線:標量條件裡的 cast 不再限制資料向量寬度(#3759)(11-05)
- 1.4 主線:Z3 的 rlimit 可按 kernel 配置,預設仍為十萬(#3752,依賴 tile-ai/tvm #91)(11-05)
- 1.5 主線:巢狀 T.macro 複用轉換模板,前端重複編譯開銷收斂(#3760)(11-05)
- 1.6 新 issue:條件標量化後缺 __syncthreads,0.1.12 至 0.1.15 靜默生成錯誤 CUDA(#3808)(11-05)
- 1.7 新 issue:無 warp specialization 時 T.Pipelined 在 tcgen05 等待前過載 stage(#3800)(11-05)
- 1.8 在審:遷到原生 Apache TIRx 的完整遷移草案(#3803);共享轉置改走 ldmatrix/stmatrix(#3809)(11-05)
- 二、多後端適配
- 2.1 昇騰:bitwise_rshift 的 64 位 dtype 改為顯式報錯,不再段錯誤(#1954)(11-05)
- 2.2 昇騰:A2/A3 效能技能對齊 Ascend950 工作流(#1932)(11-05)
- 2.3 昇騰:threads=2 下 tmp workspace 被無條件減半,sin/cos 在間接定址 tile 上 97.1% 元素算錯(#1991 / #1992)(11-05)
- 2.4 海光:分支三筆 LDS 相位修,並首次開出兩個 PR(#33 / #34)(11-05)
- 2.5 摩爾執行緒:README 文件更新,結束自 09-30 以來的靜默(11-05)
- 2.6 沐曦與 mlir-ascend:沐曦連續第三個視窗無推送;mlir-ascend 三項 TLIR 運算元在審(11-05)
- 三、生態與採用方
- 3.1 TileOPs:九筆合入,廣播運算元就地讀、GEMM 單波次填充、reduction 佈局快取(11-05)
- 3.2 TileOPs:roofline 計價規則重寫,三條不可能的頻寬行回到合理區間(#2830)(11-05)
- 3.3 TileOPs:兩筆合入相隔 2.5 分鐘導致 main 的 GPU Smoke 失敗,當日補修(#2837)(11-05)
- 3.4 TileOPs 在審:三階段瓦片因子、TN 單元素收縮、FP8 分組 GEMM 的 TMA 縮放載入(#2836 / #2838 / #2839)(11-05)
- 3.5 TileFoundry:WGMMA 同步計劃合入(#286),部分歸約的完成語義在審(#288)(11-05)
- 3.6 TileSight:H20 的 WGMMA 能力宣告修正,為倉庫首個 PR(#1)(11-05)
- 3.7 採用方與第三方:TileKernels、FlashQLA 無推送;兩個新的第三方倉出現(11-05)
- 四、社群、教程與活動
- 4.1 文件站三次更新;夜間倉無主分支提交;昇騰每日測試 3030 項全過(11-05)
- 4.2 新聞側:四組查詢、Hacker News 與 arXiv 均無視窗內命中(11-05)
- 五、趨勢觀察
- 附、素材與核查說明
今日重點
這一天最密集的產出來自運算元庫而不是編譯器。TileOPs 在 21:06 到 22:46 的一個多小時裡連續合入六筆,加上 14:01 與 21:31 的兩筆,一天九筆,全部是 perf、fix、ci 標籤。三件事同時推進:一是廣播運算元就地讀,lerp、where、clamp、masked_fill 五類運算元不再把 stride 0 的廣播張量先展開成輸出形狀,實測廣播行耗時降到主線的 0.13 到 0.44 倍(#2832);二是單波次啟動,一元逐元素在「體內最多兩條便宜指令、讀寫總量不超過 L2 的四分之三」時縮成一個 256 執行緒的波次,並且順手修正了 ReLU 對 NaN 返回 0 的老問題(#2829);三是瓦片選擇器,給小 M 的 NT GEMM 造了一個寬度能填滿一波的原生寬度瓦片,decode-down 從落後變成領先最快對比項 3%(#2834)。三件事的共同點不是「更快」,而是每筆都帶著可復現的 A/B 表、都標出自己仍然落後的那一行。
第二條是可信度問題,而且橫跨兩個倉。主倉新開 issue #3808 報告:從 0.1.12 到 0.1.15,條件相關的標量化之後,共享記憶體 epilogue 少了一道 __syncthreads(),編譯器不報錯、生成的 CUDA 是錯的;報告者給了完整復現指令碼與四個版本號,屬於「靜默錯程式碼」而非效能退化。同一視窗裡 #3800 報的是另一種靜默錯:不開 warp specialization 時,T.Pipelined 會在 T.tcgen05_gemm 的 barrier 等待之前過載下一 stage,於是 MMA 讀的 operand 已經被覆蓋。這兩條與主倉當天合入的 #3759 正好構成一組對照——#3759 修的是「標量條件裡的 int64 cast 把資料向量寬度從 16 lane 壓到 4 lane」,同一個標量化/向量化交界處,一邊在修寬度,一邊在漏同步。
第三條在昇騰側。tilelang-ascend 合入 #1954,把 T.tile.bitwise_rshift 在 int64/uint64 上的編譯期段錯誤改成顯式 ValueError,只允許 int16/uint16/int32/uint32;這不是新增 64 位右移支援,而是把「編譯器崩掉」變成「編譯器說不」。同一倉當天新開兩條同題 issue(#1991、#1992),報告在 910_9362 / CANN 9.0.0、threads=2(C:V=1:2)下,AscendVidReduction 把 kernel 內所有 UB 分配無條件減半,連注入的 tmp workspace 也減半,結果是 sin/cos 作用在間接定址的 skip-set buffer 上時 3978/4096(97.1%)個元素算錯,sigmoid/round 則違反 AscendC 的 calCount <= tmp 契約、越界寫落在佈局後方的 scratch 上。與之並置的是同日 #1988 的每日測試報告:3030 項全過、透過率 100%。這兩組數字放在一起說明一件事——現有測試覆蓋不到「threads=2 加間接定址」這種組合,透過率不能當作正確性的替代指標。
第四條是生態面的動靜。海光倉從「功能分支上攢了四筆、未開 PR」變成正式進入評審:feat/hcu-scaled-cast 分支當天又落三筆 LDS 相位修(預設 swizzle 適配 64 bank、MN 切片消費端恢復 swizzle 規劃、K 分割槽的佈局感知讀),並開出 #33(scaled cast)與 #34(reduction 的 lane swap 與約束傳播)兩個 PR。摩爾執行緒結束了自 09-30 以來的完全靜默,主分支落了一筆 README 的安裝與運算元開發說明更新。加上 TileSight 倉庫出現首個 PR(修正 H20 的 WGMMA 能力宣告)與主倉 #3803 的 TIRx 遷移草案,可以看到組織層面的三條線在同時展開:國產後端從自研分支進入公共評審、效能建模從論文走向倉庫、編譯器底座準備換掉 vendored TVM。
一、核心專案進展
1.1 主線:TF32 轉換支援 round-to-nearest-even,SM90 與 SM100 分開診斷(#3714)
合於 12:09。此前 T.cast 到 tfloat32 只有預設的 TF32 就近舍入路徑,想要 ties-to-even 得自己匯入 CUDA helper。這一筆給 round="rn" 加了原生 lowering:非飽和形式發 cvt.rn.tf32.f32,飽和形式發 cvt.rn.satfinite.tf32.f32,並且繞開預設的 TF32 舍入建構函式、直接保留轉換後的位模式。標量與向量兩條路徑都覆蓋,TF32 的 lane 訪問走打包載體。架構診斷按形式分開:非飽和形式要求 SM90 以上,飽和形式要求 SM100 以上,不滿足就在編譯期報錯;沒有顯式指定 rounding 的 cast 行為不變。報告者是關掉 #3713 的需求方,屬於「使用者不必再自帶 helper」這類介面收斂。
來源:https://github.com/tile-ai/tilelang/pull/3714
1.2 主線:host-mapped 緩衝可顯式註解,CPU 張量不再被裝置檢查拒絕(#3802)
當天建立、當天合入(12:08 提交、16:02 合入)。CUDA kernel 可以直接訪問對映進 GPU 地址空間的 host 記憶體,但 TileLang 原有的裝置檢查假定張量引數都住在目標裝置上,於是這類 CPU 張量在啟動前就被拒。新增的 T.annotate_host_mapped_buffers(*buffers) 把被註解引數的索引寫進 PrimFunc 的 tl.host_mapped_params 屬性,支援多緩衝、合併重複宣告,並且能穿過 IR 序列化與 kernel 快取恢復。Cython 與 TVM-FFI 兩個 launcher 對被註解引數接受 CPU 張量,其餘 dtype、形狀、秩、步長檢查照舊;裝置側引數的檢查仍對齊當前執行裝置,host 張量則不提供 CUDA 裝置號。執行時新增一個原生 helper,用 cuMemHostGetFlags 校驗註冊狀態、用 cuMemHostGetDevicePointer 解析裝置可見地址,普通可分頁記憶體仍會被拒。
來源:https://github.com/tile-ai/tilelang/pull/3802
1.3 主線:標量條件裡的 cast 不再限制資料向量寬度(#3759)
合於 19:27,標籤裡帶 ascend。起因是 #3734 在 DecoupleTypeCast 之後正確地收緊了資料 cast 的寬度限制,但把條件裡的 cast 也算進去了:一段 T.assume(n % 16 == 0) 保護的 BF16 複製,因為條件裡有 T.Cast("int64", ...),在 b5cb02633 上從 16 lane 被壓到 4 lane,即使條件在 16 個 lane 上一致。修法是把 planner 已有的條件一致性分析複用起來,把 cast 寬度豁免從緩衝索引擴充套件到條件:if 語句、tirx.if_then_else 與標量謂詞的 Select 的條件單獨訪問,分支仍然受限;在 lane 級向量謂詞的目標上仍保留條件 cast 的約束。stochastic rounding 的 rbits 分組在跳過寬度約束時也被保留。這一筆與視窗內的兩個新 issue 落在同一片區域,說明標量化與向量化的交界是當前缺陷的高發帶。
來源:https://github.com/tile-ai/tilelang/pull/3759
1.4 主線:Z3 的 rlimit 可按 kernel 配置,預設仍為十萬(#3752,依賴 tile-ai/tvm #91)
合於 23:05,前置是 fork 裡的 tvm #91(23:03 合入)。改動把 Z3 求解器的 rlimit 從寫死的 100,000 變成可按 kernel 定製:不配置時保持舊預設值。作者的表述很直接——「這間接修掉了很多與 Z3 證明器相關的問題」。這類改動的意義在於把「證明超預算」從不可調的環境行為變成可調的編譯引數,與昨天昇騰側收緊 AutoSchedule 週期下界(把 Z3 求解時間降 10.04%)是同一方向的兩種做法:一個改搜尋起點,一個改搜尋預算。
來源:https://github.com/tile-ai/tilelang/pull/3752、https://github.com/tile-ai/tvm/pull/91
1.5 主線:巢狀 T.macro 複用轉換模板,前端重複編譯開銷收斂(#3760)
合於 23:38。問題在外層函式每次執行時,巢狀的 T.macro 定義都會重新做原始碼解析、AST 重寫與 Python 編譯,即使 kernel 本身已經命中快取。改法是引入一個有界的程序內快取,只存與上下文無關的轉換結果(未繫結的程式碼與重寫後的原始碼),每個宏例項再繫結當前的 globals 與閉包值;依賴上下文的轉換走原路徑,被包裝過的函式繞開快取。宏呼叫點與 kernel 快取都不變。迴歸覆蓋新閉包、globals、捕獲的緩衝、註解、上下文相關回退與閉包生命週期;CPU 側前端檢查透過,GPU 執行本地未驗證——作者把這條邊界寫清楚了。
來源:https://github.com/tile-ai/tilelang/pull/3760
1.6 新 issue:條件標量化後缺 __syncthreads,0.1.12 至 0.1.15 靜默生成錯誤 CUDA(#3808)
21:14 開出,是本視窗最重的缺陷報告。復現是一個共享記憶體 epilogue 的 kernel:0.1.12、0.1.13、0.1.14、0.1.15 四個版本都會在條件相關的標量化之後丟掉一道 __syncthreads(),編譯器不報任何錯,生成的 CUDA 直接算錯。報告者給了完整可執行的 Python 復現指令碼與環境(Python 3.14.9、CUDA 2.15.0+cu134)。這類問題的破壞力在於它跨了四個釋出版本且無任何告警——凡是依賴共享記憶體在階段之間做交接的 kernel 都可能踩到,而使用者側看到的只是「結果不對」。
來源:https://github.com/tile-ai/tilelang/issues/3808
1.7 新 issue:無 warp specialization 時 T.Pipelined 在 tcgen05 等待前過載 stage(#3800)
08:44 開出。觸發條件是關掉 warp specialization(tl.disable_warp_specialized=True)後跑 T.Pipelined 加 T.tcgen05_gemm 的迴圈:stage 在 mbarrier_wait_parity 之前就被下一輪複製覆蓋,MMA 讀到的是被改寫的 operand。報告者把 lowering 出來的 kernel 原始碼貼了出來,指認 As 的寫入與 tcgen05mma_ss 的讀之間缺少順序約束。與 #3808 同屬「生成的程式碼看起來沒問題、實際不成立」這一類,區別在於這條發生在 Blackwell 的 tensor memory 路徑上。
來源:https://github.com/tile-ai/tilelang/issues/3800
1.8 在審:遷到原生 Apache TIRx 的完整遷移草案(#3803);共享轉置改走 ldmatrix/stmatrix(#3809)
#3803 是本期最有分量的在審項:把 vendored 的 TVM 編譯器換成獨立的 Apache TIRx,前端、IR 變換、程式碼生成、JIT 適配與打包全部遷到原生物件模型,並 pin 了 apache/tvm-tirx 的 7e49a148 與 apache/tvm-ffi 的 3e906bc6 兩個上游提交,依賴樹保持乾淨上游、不注入補丁。作者把它標成 draft,本身就是一次「底座換代」的公開準備。
#3809(01:27 提交)處理另一件具體的事:T.transpose 在兩個共享緩衝之間一直走通用 SIMT 迴圈,8 位與 16 位瓦片用不上共享記憶體頻寬,examples/deepseek_v32 裡的 FP8 稀疏 MLA 只能手寫 16x16 暫存器暫存來繞開;同時 T.Pipelined 給源區加了一維 stage 軸後,秩不匹配會讓 MakeIndices 直接失敗。改法是在支援 stmatrix 的目標(SM90 及以上、排除 CuTeDSL)先試 warp-matrix lowering:16 位瓦片每個 8x8 塊用一次 ldmatrix 加一次 stmatrix.trans,8 位瓦片按 16x16 位元組塊用兩次 ldmatrix.trans 加 __byte_perm 選擇拼出四行,其餘區域回退到原 SIMT 路徑。
來源:https://github.com/tile-ai/tilelang/pull/3803、https://github.com/tile-ai/tilelang/pull/3809
二、多後端適配
2.1 昇騰:bitwise_rshift 的 64 位 dtype 改為顯式報錯,不再段錯誤(#1954)
合於 16:36(ascendc_pto 分支)。A2/A3 上透過公開 API T.tile.bitwise_rshift 傳 int64/uint64 運算元時,編譯器會在 OptimizeForTarget 階段直接段錯誤。修法是在構造 intrinsic 之前檢查 dst 與 src0 的 dtype,只允許 int16/uint16/int32/uint32,其餘拋帶實際 dtype 與支援列表的 ValueError;位移量 scalarValue 不限表示式 dtype;Buffer 與 BufferRegion 統一走 _underlying_buffer()。新增直接 API 診斷測試與子程序崩潰迴歸測試。作者明確寫了邊界:本次不新增 64 位右移支援,也不改直接構造底層 intrinsic 的路徑。關聯 #1720。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1954
2.2 昇騰:A2/A3 效能技能對齊 Ascend950 工作流(#1932)
15:36 合入,純技能與流程重構:把 A2/A3 的效能調優技能與 A5(Ascend950)工作流結構對齊並統一,新增晶片型別識別與晶片資訊檢索兩項技能。單看沒有功能價值,但它意味著效能調優的「操作手冊」開始按晶片代際分叉——同一套 skill 要能識別自己跑在哪一代上,否則調優建議會跨代誤用。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1932
2.3 昇騰:threads=2 下 tmp workspace 被無條件減半,sin/cos 在間接定址 tile 上 97.1% 元素算錯(#1991 / #1992)
17:03 與 17:25 連開兩條同題 issue,標題與正文一致(內容上可視作同一條被提交了兩次)。報告環境是 910_9362 / CANN 9.0.0、threads=2(C:V=1:2)。根因在 AscendVidReduction:它把 kernel 內所有 UB 分配無條件減半,包括 InjectTmpBuffer 注入的 tmpspace,而 skip-set buffer(GM 索引含 UB load 的間接定址,如 T.copy(KV[b_i, idx_ub[bi_i], g_i, :D], kv_ub))的形狀保持整塊、本來不參與 vid 減半。後果分兩檔:sin/cos 直接算錯(upstream HEAD 上 3978/4096 個元素錯,97.1%);sigmoid/round 違反 AscendC 的 calCount <= tmp 契約,數值碰巧正確但越界寫落在佈局後方的 scratch 上,佈局一變就是靜默資料汙染。報告者說明這是 #1897 review 時發現的,而 #1897 的修法只覆蓋了 count 側語義(HasModifiedBuffer 跳過 workspace 槽位),workspace 側仍未處理。
來源:https://github.com/tile-ai/tilelang-ascend/issues/1991、https://github.com/tile-ai/tilelang-ascend/issues/1992
2.4 海光:分支三筆 LDS 相位修,並首次開出兩個 PR(#33 / #34)
主分支仍停在 11-02,當天三筆都落在 feat/hcu-scaled-cast 分支,作者是海光側(提交域 hygon.cn):
- 10:06
fix(hcu): adapt default LDS swizzle to 64-bank phases - 14:04
fix(hcu): restore LDS swizzle planning for MN slice consumers - 19:49
fix(hcu): support K-partitioned layout-aware LDS reads
更重要的是評審狀態的變化:09:08 開出 PR #33(scaled cast,即分支主體),19:52 開出 PR #34(reduction 的 XOR shuffle 用 DPP 與 lane swap 最佳化,外加「佈局推斷先傳播已知約束再分配自由候選」的修法)。上一期報告寫的是「四筆都在功能分支上、尚未開 PR」,本視窗這兩筆 PR 把海光從自研分支推進到公共評審,是國產後端裡第一個走完這條路徑的。
來源:https://github.com/tile-ai/tilelang-hygon/pull/33、https://github.com/tile-ai/tilelang-hygon/pull/34、https://github.com/tile-ai/tilelang-hygon/commit/07451160e5aa758419201adb598725cbef70e3a4
2.5 摩爾執行緒:README 文件更新,結束自 09-30 以來的靜默(11-05)
19:46 主分支落了一筆 docs: update README installation and operator development references,是本倉自 09-30 以來第一次提交。同期 6 個遺留 PR(#1 至 #7 中的依賴升級、TileOPs 子模組、MUSA 整合的證據門檻文件)時間戳被一併重新整理,但內容仍是 04 月到 08 月的老提案,本視窗沒有程式碼變更。判斷上宜保守:文件提交說明倉庫有人在維護,不足以證明後端開發恢復;真正的訊號要看後續是否出現運算元或 lowering 提交。
來源:https://github.com/tile-ai/tilelang-musa/commit/4cac45b859108a86ac9a0b1cbfdeb33da7711697
2.6 沐曦與 mlir-ascend:沐曦連續第三個視窗無推送;mlir-ascend 三項 TLIR 運算元在審(11-05)
tile-ai/tilelang-metax 最後推送仍是 10-26,連續第三個視窗靜默。tile-ai/tilelang-mlir-ascend 主分支無新提交,但三項基於 TLIR 方言的運算元 PR 在推進:#203 給 TLIR 加 tl.sub/mul/div/max/min 並補 TL 到 HIVM 的 lowering,01:42 更新;#211 加 tl.shl、#212 加 tl.shr,兩筆在 19:44 與 20:58 先後提交。這條線在做的事是把 TileLang 的語義一層層落進 MLIR 方言,節奏慢但方向單一。
來源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-mlir-ascend/pull/203、https://github.com/tile-ai/tilelang-mlir-ascend/pull/211、https://github.com/tile-ai/tilelang-mlir-ascend/pull/212
三、生態與採用方
3.1 TileOPs:九筆合入,廣播運算元就地讀、GEMM 單波次填充、reduction 佈局快取(11-05)
九筆集中在 14:01 與 21:06 至 22:46 兩個時段,逐筆看主幹:
- #2832 廣播運算元就地讀:
LerpTensorFwdOp、WhereFwdOp、ClampTensorFwdOp、MaskedFillScalarFwdOp、MaskedFillTensorFwdOp五類運算元不再先把 stride 0 的廣播張量展開成輸出形狀,而是在一個 kernel 裡按 stride 0 直接讀。一個共享 builder 服務全部五種,運算元形狀合併成一個迴圈巢狀,16 位元組以內的向量留在一行內層,每個運算元只寫自己的逐元素體。廣播行耗時是主線的 0.13 到 0.44 倍,與 torch-compile 打平或更好;其餘逐元素行與主線逐位元組一致。 - #2829 一元逐元素單波次:體內最多兩條各自降為一條指令的操作、讀寫總量不超過 L2 的四分之三、結果與運算元同寬時,縮成 256 執行緒的一波。順帶修了 ReLU 對 NaN 返回 0 的問題(改為返回 NaN,16 位走 torch 位模式、float32 用規範 NaN)。H200 上 327 項 GPU 迴歸全過,13 項新用例過了 Compute Sanitizer memcheck。
- #2834 小 M NT GEMM 的原生寬度瓦片:
ds-v3-decode-down(128x7168x2048,NT)在主線上只用 132 個 SM 裡的 112 個,112 寬的瓦片在 0.1.12 上被拆成七個 n16 WGMMA。改成每 16 個 K 發一次瓦片自身寬度的 WGMMA 之後,decode-down 領先最快對比項 3%,gate-up-b32 比主線快 1.4%(仍落後 cuBLAS 0.5 到 1%)。新增的fill瓦片只對 96 < M <= 128 且 N 能整除 8 的形狀啟用,結構因子 0.58 單獨擬合。 - #2835 reduction 佈局快取:歸約 kernel 每次呼叫都重建軸置換與輸出形狀,而
ReduceCall是固定的。改為每個 kernel 例項只算一次、與既有張量介面卡共享純形狀 helper,不引入全域性快取。六組 FP32/BF16VarMeanFwdOp配置的完整熱呼叫耗時下降 4.33% 到 6.17%,生成的 CUDA 與主機原始碼完全一致——省的是呼叫準備成本,不是 GPU 時間。 - #2833 CI 修:GitHub 廢棄 runner 2.336.0 之後
gpu-smoke全部排隊,entrypoint.sh又帶--disableupdate無法自升級,RUNNER_VERSION從 2.336.0 提到 2.338.2。
來源:https://github.com/tile-ai/TileOPs/pull/2832、https://github.com/tile-ai/TileOPs/pull/2829、https://github.com/tile-ai/TileOPs/pull/2834、https://github.com/tile-ai/TileOPs/pull/2835、https://github.com/tile-ai/TileOPs/pull/2833
3.2 TileOPs:roofline 計價規則重寫,三條不可能的頻寬行回到合理區間(#2830)
夜間任務 41873301494 抓出三行物理上不可能的 roofline:GQAVarlenFwdOp 的兩個 gpt-oss-serve-window 配置讀到 7.48 與 7.50 TB/s,GroupedGEMMFwdOp 的 qwen3-30b-ep8-empty 讀到 5.30 TB/s。這一筆把計價規則按「這次呼叫實際讀了多少」重寫:GQA 的稠密、varlen、分頁前向只對「有 query 能看到的 key 行」計 K/V、對「能看到 key 的 query 行」計 q、對 key 實際到達的截斷位置計 RoPE 表、對看到 key 的請求計 FP8 縮放;GroupedGEMM 透過新的 grouped_gemm_roofline 按 batch_sizes 給某行的組計價;NSA 壓縮、NSA 稀疏、DSA decode、MHA 分頁與 FP8 lightning indexer 都只在真正讀的位置計費,DSA 還認 is_causal=False。修正後三行變成 0.62、0.62、2.65 TB/s。附加收益是修掉了空 K/V 儲存與空分頁池的除零。測試節點加 16。
來源:https://github.com/tile-ai/TileOPs/pull/2830
3.3 TileOPs:兩筆合入相隔 2.5 分鐘導致 main 的 GPU Smoke 失敗,當日補修(#2837)
22:46 合入的這筆本身就是一條紀律記錄。#2824 在 21:06 合入時把 grouped GEMM 的 roofline 行改成必須宣告 sizes 並從它推導 G;#2830 在 21:08 合入,寫的還是隻有 G 沒有 sizes。兩筆相隔約 2.5 分鐘,各自 CI 都是綠的,合併後 main 的 GPU Smoke 立刻報 RowError: row misses ['sizes']。修法是把 _grouped 的 G: 2 換成 sizes: [1, 0],斷言的位元組數(24、32)不變。這類事故的成因不是誰寫錯了,而是並行合併缺少 post-merge 門禁:單筆 CI 透過無法覆蓋「兩筆語義相交」的情況,只能靠合入後立刻跑全量。
來源:https://github.com/tile-ai/TileOPs/pull/2837、https://github.com/tile-ai/TileOPs/pull/2824、https://github.com/tile-ai/TileOPs/pull/2830
3.4 TileOPs 在審:三階段瓦片因子、TN 單元素收縮、FP8 分組 GEMM 的 TMA 縮放載入(#2836 / #2838 / #2839)
#2836 給 coop2_swap 的三階段瓦片單獨擬合因子(1.36),因為按 FLOP 算它並不比四階段快,但按位元組與 K 步數被高估了約 10%,結果把 160 寬的瓦片擠掉、夜間成績從 1.005 到 1.011 倍落到 0.985 到 0.995 倍 QuACK。因子只在那 42 個 off-manifest 形狀上擬合(平臺區 1.344 到 1.392),44 形狀掃描上選擇器的 regret 從 0.89% 降到 0.06% 幾何均值、最大從 5.04% 降到 0.99%。#2838 修的是分派器把 K 當成連續維、於是否決 TN 佈局 K=1 的合法呼叫(A=[1,64]、B=[1,64] 求 AᵀB 本來可以用 cp.async 核心算),改為回退檢查兩側運算元的物理連續維是否各有至少 4 位元組行寬,新增 24 個測試節點。#2839(00:56 提交)讓 FP8 分組 GEMM 用 TMA 載入 A 的縮放(每盒四個 K 步、stage_scales 2、在三行上標定),並給 K <= 1536 的非融合緊分組 BF16/FP16 GEMM 換 64x256 的 ping-pong 瓦片;FP8 qwen3-30b-ep8-skew-up 從 0.973 升到 1.015 倍 DeepGEMM,BF16 skew-down 從 0.962 升到 1.099 倍 QuACK。#2840(06:51 提交)補 SM89 的 prefill sinks,並把 sink 的有無納入編譯例項身份。
來源:https://github.com/tile-ai/TileOPs/pull/2836、https://github.com/tile-ai/TileOPs/pull/2838、https://github.com/tile-ai/TileOPs/pull/2839、https://github.com/tile-ai/TileOPs/pull/2840
3.5 TileFoundry:WGMMA 同步計劃合入(#286),部分歸約的完成語義在審(#288)
13:40 合入的 #286 是上一期在審項,處理 WGMMA 的非同步性:指令發出後仍會繼續讀共享記憶體與暫存器,後續訪問必須排在它的效果之後。做法是給每個運算元效果宣告完成條件,把 fence、commit、wait 掛到需要它們的 HIR 節點上,lowering 在指令發射點消費這些步驟;若後續作用域放不下集體等待,就在發射作用域裡做完。同時修正暫存器 Reshard 的別名,讓每個元素在持有副本的每個執行緒上留在同一個暫存器,並用這個儲存身份追蹤暫存器依賴;SM90 上以 sm_90a 單獨編譯發出 WGMMA 組步驟的 CUDA 單元。作者保留了兩條風險宣告:矩陣指令發射器仍未接上,所以這筆不證明執行正確性、也不給效能;TMA 的完成語義另行處理。20:12 另開 #288,處理「已完成的部分歸約」在分析中的計價。
來源:https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/pull/288
3.6 TileSight:H20 的 WGMMA 能力宣告修正,為倉庫首個 PR(#1)
14:37 提交、17:26 更新,是 tile-ai/TileSight(24 星)的第一個 PR。該倉庫是 07-24 arXiv 論文《TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters》的開源實現:從瓦片形狀、資料移動、軟體流水與硬體檔案出發,在 CPU 上估計運算元代價、快取流量與執行時間,CUDA/HIP 探針單獨測目標 GPU。這筆改動很小但方向明確:H20 是 Hopper/GH100,帶 requires_arch=("wgmma",) 的計劃原先因為架構檔案沒有宣告 WGMMA 支援而拋 UnsupportedError,現在把 support_wgmma=True、並顯式保持 support_utcmma=False,於是這類計劃被接受、而 tcgen05 與 Tensor Memory 需求仍被拒。作者同時指出容量要按 SKU 分開處理(NVIDIA 同時列 96 GB 與 141 GB 的 H20),本次只動能力後設資料,容量、頻寬、頻率、吞吐與利用率數值全部不變。驗證是 CPU 上 847 項全過,明確宣告「不是 GPU 效能驗證」。
來源:https://github.com/tile-ai/TileSight/pull/1、https://arxiv.org/abs/2607.22432
3.7 採用方與第三方:TileKernels、FlashQLA 無推送;兩個新的第三方倉出現(11-05)
deepseek-ai/TileKernels 仍停在 10-14,QwenLM/FlashQLA 停在 10-30,本視窗都沒有提交,FlashQLA 最新發布仍是 09-30 的 v0.1.3。第三方側出現兩個新的小倉:Tianchen-ai/tilelang-weft-bench(「TileLang-Weft full CPU operator benchmarks using the shared Weft Bench engine」,0 星,建立 CPU 側運算元基準的統一引擎)、flytiger-eco/tilelang-wrapper-for-sail(0 星,暫無描述)。兩者都很早期,先記入觀察名單,不宜當技術進展寫。
來源:https://github.com/deepseek-ai/TileKernels、https://github.com/QwenLM/FlashQLA、https://github.com/Tianchen-ai/tilelang-weft-bench、https://github.com/flytiger-eco/tilelang-wrapper-for-sail
四、社群、教程與活動
4.1 文件站三次更新;夜間倉無主分支提交;昇騰每日測試 3030 項全過(11-05)
tile-ai/tilelang.github.io 在 16:07、19:33、23:43 各推一次 Update docs,對應主倉當天的介面變化(host-mapped 註解、TF32 rounding、Z3 配置)。TileOPs-nightly 只有分支推送、主分支無新提交;TileOPs.github.io 在 09:02 有一次推送。昇騰倉的每日測試 issue #1988 記錄 11-05 08:16 一輪 3030 項全部透過、透過率 100%,與同日新開的 #1991/#1992 形成對照(見 2.3)。
來源:https://github.com/tile-ai/tilelang.github.io、https://github.com/tile-ai/tilelang-ascend/issues/1988
4.2 新聞側:四組查詢、Hacker News 與 arXiv 均無視窗內命中(11-05)
Google News 四組查詢(英文 "TileLang" OR "tile-lang"、TileLang GPU kernel,中文 TileLang 運算元、TileLang 編譯器)裡,最新一條命中是 376 小時前的 SemiAnalysis《The Huawei-DeepSeek Symbiosis》,其餘集中在 09-30 那一波「DeepSeek 開源昇騰版 TileLang、DeepGEMM、DeepEP」的報道(約 880 小時前),全部早於本視窗。Hacker News 以 tilelang 檢索 24 小時內 0 條,抓到的三條 Show HN 是終端資料檢視器、工作流引擎與 macOS 檔案管理器,與主題無關。arXiv 的 all:tilelang 最新命中是 11-03 的《Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving》(視窗外,屬引用 TileLang 的工作),再往前是 07-24 的 TileSight 本體。報告主體因此仍全部來自程式碼側。
來源:https://news.google.com/、https://news.ycombinator.com/、https://arxiv.org/
五、趨勢觀察
第一,運算元庫的競爭點已經從「運算元數量」徹底轉到「選擇器與準備成本」。今天 TileOPs 九筆裡,真正改 GPU kernel 的只有 #2832(廣播就地讀)與 #2834(原生寬度瓦片),#2835 省的是呼叫準備(4.33% 到 6.17%,生成原始碼不變),#2836 改的是一個擬合因子,regret 從 0.89% 降到 0.06%。這說明效能差距的剩餘空間主要不在「怎麼算」,而在「選哪個瓦片、少做哪些準備工作」。這條路走到頭會得到一個用標定契約描述的代價模型,而不是一堆手調引數——今天 #2836 的「單獨擬合、只動一個常數」正是這個形態的雛形。
第二,編譯器側當前的主要缺陷型別是「靜默生成錯程式碼」,而不是編譯失敗或效能不足。#3808 跨 0.1.12 到 0.1.15 四個版本少一道 __syncthreads(),#3800 在 Blackwell 路徑上讓 MMA 讀到被覆蓋的 operand,#1991/#1992 在昇騰 threads=2 下讓 sin/cos 算錯 97.1% 的元素——三條都不報錯、都能透過常規測試、都要靠人寫出針對性的復現才能抓到。這類缺陷對使用者的傷害最大,因為它把「編譯器對不對」這個前提變成了需要驗證的假設。與之對應的工程動作應當是擴大形狀組合的差分測試(對照組用 reference 實現逐元素比),而不是繼續加單元測試數量。
第三,昇騰側的節奏從「架構改動後清迴歸」轉向「清正確性邊界」。#3512 那類結構性改動帶來的四類迴歸在上一個視窗剛清完,本視窗的兩筆合入(段錯誤改報錯、效能技能分代)和兩條新 issue(tmp workspace 減半、calCount <= tmp 違約)都屬於邊界收斂。值得注意的是 #1991 報告者是 #1897 的 review 方,而 #1897 的修法只覆蓋了 count 側——同一處佈局邏輯的另一個側面在兩週後才被指出,說明昇騰後端的佈局不變數還沒有形成可檢查的書面契約。
第四,國產後端的「公共化」在提速。海光從自研分支走到兩個正式 PR,摩爾執行緒恢復文件維護,mlir-ascend 用三個小 PR 把算術運算元一個個搬進 TLIR 方言。三條線的共同點是都不再等「一次性合入一個大後端」,而是把可評審的最小單元推到公共倉庫。這對上游是好事:評審者能早就佈局與 lowering 的設計提意見,而不是在三萬行一次性合入時才看。
最後一條風險面:主倉 #3803 的 TIRx 遷移如果推進,3rdparty/tvm 的 vendored 結構會被替換成兩個 pin 的上游提交,屆時所有依賴 TVM 內部 API 的後端(昇騰、海光、mlir-ascend 的膠水層)都要跟著改。草案裡明確寫了「不注入依賴補丁」,這既是承諾也是約束——今天 mlir-ascend 把 BishengIR 介面卡搬進 contrib、昇騰把效能技能按晶片代際分叉,都可以看作是為這類底座切換預留的可替換面。
附、素材與核查說明
信源核查表
| 信源 | 核查結果 |
|---|---|
| tile-ai 組織(29 倉庫) | 視窗內推送 10 個倉庫 |
| 主倉 tilelang | 主分支 5 筆合入;新開 issue 2 條(#3808、#3800);在審更新以標量化、CuTeDSL、Profiler、Ascend 四簇為主 |
| TileOPs | 主分支 9 筆合入;在審 4 筆新增或更新;無新 issue |
| TileFoundry | 主分支 1 筆合入(#286);在審 1 筆(#288) |
| TileSight | 首個 PR(#1),無主分支提交 |
| tilelang-ascend | 預設分支 ascendc_pto 2 筆合入(#1932、#1954);新開 issue 3 條(#1991、#1992 同題,#1988 為每日測試報告) |
| tilelang-mlir-ascend | 主分支無新提交;#203、#211、#212 三筆 TLIR 運算元 PR 推進 |
| tilelang-hygon | 主分支無新提交(停在 11-02);feat/hcu-scaled-cast 分支 3 筆;新開 PR #33、#34 |
| tilelang-musa | 主分支 1 筆文件提交(自 09-30 以來首次);6 個遺留 PR 時間戳重新整理,內容未變 |
| tilelang-metax | 視窗內無推送(連續第三個視窗) |
| TileOPs-nightly / TileOPs.github.io / tilelang.github.io | 夜間倉主分支無提交;TileOPs 文件站 1 次推送;tilelang 文件站 3 次 Update docs |
| 釋出頁(tilelang / TileOPs / TileFoundry) | 視窗內無新版本;tilelang 最新仍為 09-30 的 v0.1.15,TileOPs 為 10-27 的 v0.1.0,TileFoundry 為 10-25 的 v0.0.4rc1 |
| TileKernels / FlashQLA | 視窗內無推送 |
| TileLang×RISC-V 四源巡檢 | 無變化(四個上游的引用與 PR 狀態與上期快照一致) |
| Google News(中英文四組)/ Hacker News / arXiv | 無視窗內實質命中,最新命中為 376 小時前 |
完整信源清單
- [1] tilelang TF32 round-to-nearest-even 轉換(#3714)— https://github.com/tile-ai/tilelang/pull/3714
- [2] tilelang host-mapped 緩衝註解(#3802)— https://github.com/tile-ai/tilelang/pull/3802
- [3] tilelang 標量條件 cast 保留向量寬度(#3759)— https://github.com/tile-ai/tilelang/pull/3759
- [4] tilelang Z3 rlimit 可配置(#3752)— https://github.com/tile-ai/tilelang/pull/3752
- [5] tile-ai/tvm Z3 資源預算配置(#91)— https://github.com/tile-ai/tvm/pull/91
- [6] tilelang 宏轉換模板複用(#3760)— https://github.com/tile-ai/tilelang/pull/3760
- [7] tilelang 條件標量化後缺同步屏障(#3808)— https://github.com/tile-ai/tilelang/issues/3808
- [8] tilelang 無 warp specialization 時 stage 提前過載(#3800)— https://github.com/tile-ai/tilelang/issues/3800
- [9] tilelang 遷到原生 Apache TIRx(#3803)— https://github.com/tile-ai/tilelang/pull/3803
- [10] tilelang 共享轉置走 ldmatrix/stmatrix(#3809)— https://github.com/tile-ai/tilelang/pull/3809
- [11] tilelang 多快取目錄(#3807)— https://github.com/tile-ai/tilelang/pull/3807
- [12] tilelang-ascend bitwise_rshift dtype 校驗(#1954)— https://github.com/tile-ai/tilelang-ascend/pull/1954
- [13] tilelang-ascend A2/A3 效能技能對齊(#1932)— https://github.com/tile-ai/tilelang-ascend/pull/1932
- [14] tilelang-ascend tmp workspace 被減半(#1991)— https://github.com/tile-ai/tilelang-ascend/issues/1991
- [15] tilelang-ascend tmp workspace 被減半(#1992)— https://github.com/tile-ai/tilelang-ascend/issues/1992
- [16] tilelang-ascend 每日測試 3030 項全過(#1988)— https://github.com/tile-ai/tilelang-ascend/issues/1988
- [17] tilelang-hygon scaled cast PR(#33)— https://github.com/tile-ai/tilelang-hygon/pull/33
- [18] tilelang-hygon reduction lane swap PR(#34)— https://github.com/tile-ai/tilelang-hygon/pull/34
- [19] tilelang-hygon LDS 相位提交 — https://github.com/tile-ai/tilelang-hygon/commit/07451160e5aa758419201adb598725cbef70e3a4
- [20] tilelang-musa README 文件提交 — https://github.com/tile-ai/tilelang-musa/commit/4cac45b859108a86ac9a0b1cbfdeb33da7711697
- [21] tilelang-mlir-ascend TLIR 算術運算元(#203)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/203
- [22] tilelang-mlir-ascend tl.shl(#211)— <https://github.com/til
e-ai/tilelang-mlir-ascend/pull/211>
- [23] tilelang-mlir-ascend tl.shr(#212)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/212
- [24] TileOPs 一元逐元素單波次(#2829)— https://github.com/tile-ai/TileOPs/pull/2829
- [25] TileOPs 廣播運算元就地讀(#2832)— https://github.com/tile-ai/TileOPs/pull/2832
- [26] TileOPs 小 M NT GEMM 原生寬度瓦片(#2834)— https://github.com/tile-ai/TileOPs/pull/2834
- [27] TileOPs reduction 佈局快取(#2835)— https://github.com/tile-ai/TileOPs/pull/2835
- [28] TileOPs CI runner 升級(#2833)— https://github.com/tile-ai/TileOPs/pull/2833
- [29] TileOPs roofline 計價重寫(#2830)— https://github.com/tile-ai/TileOPs/pull/2830
- [30] TileOPs grouped GEMM 行宣告 sizes 修復(#2837)— https://github.com/tile-ai/TileOPs/pull/2837
- [31] TileOPs FP8 分組 GEMM 連續佈局(#2824)— https://github.com/tile-ai/TileOPs/pull/2824
- [32] TileOPs 三階段瓦片因子(#2836)— https://github.com/tile-ai/TileOPs/pull/2836
- [33] TileOPs TN 單元素收縮(#2838)— https://github.com/tile-ai/TileOPs/pull/2838
- [34] TileOPs FP8 分組 GEMM 的 TMA 縮放載入(#2839)— https://github.com/tile-ai/TileOPs/pull/2839
- [35] TileOPs SM89 prefill sinks(#2840)— https://github.com/tile-ai/TileOPs/pull/2840
- [36] TileFoundry WGMMA 同步計劃(#286)— https://github.com/tile-ai/TileFoundry/pull/286
- [37] TileFoundry 部分歸約的完成語義(#288)— https://github.com/tile-ai/TileFoundry/pull/288
- [38] TileSight H20 的 WGMMA 能力宣告(#1)— https://github.com/tile-ai/TileSight/pull/1
- [39] TileSight 論文(arXiv 2607.22432)— https://arxiv.org/abs/2607.22432
- [40] 第三方倉 tilelang-weft-bench — https://github.com/Tianchen-ai/tilelang-weft-bench
- [41] 第三方倉 tilelang-wrapper-for-sail — https://github.com/flytiger-eco/tilelang-wrapper-for-sail
- [42] TileKernels(採用方,10-14 後無推送)— https://github.com/deepseek-ai/TileKernels
- [43] FlashQLA(採用方,10-30 後無推送)— https://github.com/QwenLM/FlashQLA
- [44] tilelang 釋出頁(最新 v0.1.15,09-30)— https://github.com/tile-ai/tilelang/releases
- [45] TileOPs 釋出頁(最新 v0.1.0,10-27)— https://github.com/tile-ai/TileOPs/releases
核查說明
本期所有合入、在審、issue 與提交數量均來自 GitHub 檢索與原子源的實測輸出:合入與在審走 search API(合入 18 筆、在審更新 57 筆、新開 issue 5 條),提交標題走各分支的 commits atom(12 個分支、21 筆落在視窗內),釋出版本走 releases atom。組織級匿名 API 在採集後段觸及小時限額,主分支之外的組織清單以網頁與 atom 源交叉核對,未使用估算值。#1991 與 #1992 兩條 issue 標題與正文一致,按「同題兩條」記錄,未合併為一條。海光三筆與摩爾執行緒一筆屬未合入提交,按分支級動態表述,不代表主線已合入。TileSight 與 TileFoundry 的 WGMMA 相關改動均由作者自述限定為能力後設資料或同步計劃,未聲稱執行正確性或效能,報告沿用其限定。Google News、Hacker News 與 arXiv 三條新聞通道在本視窗均無命中。