TileLang 每日動態報告(2026-11-02)
調研視窗:2026-10-30 07:00 ~ 2026-11-02 07:00(北京時間,72 小時,覆蓋週末;上一期截點為 10-30 07:00)。
信源:GitHub tile-ai 組織逐倉核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / TileOPs-nightly / 文件站)+ 合入與在審真值檢索(合入 35 筆、在審更新 82 筆)+ 新增 issue 全量核對(11 條)+ 主分支逐筆提交核查(tilelang 7、TileOPs 15、TileFoundry 3、tilelang-ascend 5、tilelang-hygon 2、tilelang-mlir-ascend 1、文件站 4)+ 釋出頁與 tag 核查(6 倉)+ 第三方生態倉核查(FlashQLA、TileKernels、RISC-V 研究鏈三倉、PTO-ISA/tilelangGym)+ Google News 中英文查詢(走代理)+ Hacker News + arXiv。
本期索引
- 今日重點:主線為 TIRX 遷移回退 ThreadSync 的精確執行緒計數(#3733),並新開跟蹤單承認正確性改進一併撤回;TileOPs 單日 13 筆效能合入;昇騰宿主固定開銷從約 450 微秒砍到約 106 微秒;72 小時合入 35 筆、在審更新 82 筆、新增 issue 11 條
- 一、核心專案進展
- 1.1 主線:ThreadSync 的精確執行緒計數回退,為 TIRX 遷移讓路(#3733)(11-01)
- 1.2 主線:CUDA 跨 warp 全歸約合併為一次共享記憶體交換,256 執行緒從 6 道柵欄降到 2 道(#3715)(10-30)
- 1.3 主線:PTO 後端的 T.Parallel 向量化落地第一里程碑,逐元素運算元走 VMI 指令(#3512)(10-30)
- 1.4 主線:昇騰 DeepGEMM 的 scale 緩衝改用寫計數多緩衝,資料載入可提到條件載入之前(#3629)(10-30)
- 1.5 主線:兩筆修復合入,CPU 標量 GEMM 只清目標區域、負索引合法化跳過不可達分支(#3709 / #3712)(10-30)
- 1.6 主線:語言層清掉 proxy 裡兩個無人使用的 Ref 佔位(#3722)(10-31)
- 1.7 新開缺陷與在審:TMEM 靜默丟結果、ROCm 符號找不到、TF32 就近偶數舍入(#3721 / #3711 / #3714)(10-30 ~ 10-31)
- 二、多後端適配
- 2.1 昇騰:Cython wrapper 惰性構造 TVM Analyzer,910B3 單次呼叫從約 450 微秒降到約 106 微秒(#1948)(10-30)
- 2.2 昇騰:RFC 提出把同步實現換成主線核心加獨立 Z3 證明服務(#1971)(10-30)
- 2.3 昇騰在審:把 TileLang-Ascend 整體遷到主線多後端架構(#1975)(10-31)
- 2.4 昇騰在審:GLM-5.3 kpool 的 fp8_mqa_logits 示例,雙 kernel 加 AIV/Cube 分工(#1966)(10-31)
- 2.5 昇騰:四筆文件與 API 事實核查合入,clamp 在 PTO 上真正生效(#1899 / #1893 / #1620 / #1600)(10-30)
- 2.6 昇騰:日測 2970 用例全部透過,11-01 有一次 workflow 失敗(#1977 / #1976)(10-31 ~ 11-02)
- 2.7 海光:MLS 邊界與持續 WASP 同步修復合入,gfx948 相容在審(#31 / #32)(10-30)
- 2.8 mlir-ascend:CI 清掉內部快取服務引用,Codex 工作流在審(#206 / #208)(10-31)
- 2.9 摩爾執行緒與沐曦視窗內無推送(11-02)
- 三、生態與採用方
- 3.1 TileOPs:72 小時 15 筆合入,13 筆壓在 11-01,幾乎每筆都帶實測表與清單迴歸門檻(11-01)
- 3.2 TileOPs:GQA 稠密前向用一條實測規則選 decode 或 prefill 程式,選擇遺憾從 1.818 降到 1.017(#2808)(11-01)
- 3.3 TileOPs:分頁 MHA decode 擴到共享記憶體裝得下的每個 head dim,GQA 分頁前向去掉舊 decode 核(#2807 / #2797)(11-01)
- 3.4 TileOPs:coop2 生產者收成單 warp、新增 192 寬轉置塊,W4A8 部分一波改成一 CTA 一 tile(#2806 / #2805)(11-01)
- 3.5 TileOPs:編譯快取按構建目標鍵控,運算元層統一檢查輸出與輸入重疊(#2802 / #2801)(11-01)
- 3.6 TileOPs 在修缺陷:Norm 四運算元在遠離零的行丟精度、約 1e36 溢位成 NaN(#2812)(11-01)
- 3.7 TileFoundry:lowering 把 issuing mesh 與視窗切分收進一處決策,83 個排程夾具位元組一致(#280 / #282)(10-31 ~ 11-01)
- 3.8 TileFoundry:站點發布 0.3 版排程文章《Scheduling, Designed for Agents》(#4 / #281)(10-31 ~ 11-01)
- 3.9 採用方:FlashQLA 合入 SM120 的 Dual-CP 前反向,2 張 RTX 5090 上對 32K 至 128K 序列與 Inter-CP 對比(#51)(10-30)
- 3.10 採用方:TileKernels 自 10-14 起無推送(11-02)
- 四、社群、教程與活動
- 4.1 RISC-V 研究鏈成形:三個倉把 Numba CUDA Python 轉 TileLang,再把受限 PrimFunc 落到單 hart RVV C(10-31)
- 4.2 PTO-ISA/til
elangGym 從 PTO-ISA/tilelang 獨立成倉,主迴歸門是昇騰 SimtVF 路徑(10-30)
- 4.3 新聞側:Google News 中英文查詢、Hacker News、arXiv 均無視窗內實質命中(11-02)
- 五、趨勢觀察
- 附、素材與核查說明
今日重點
這個視窗最重的動作是減法。主線把三週前合入的精確執行緒計數(#3291)整筆回退(#3733),理由寫得很直白:那套改動把昂貴的計數查詢搬到了滿塊柵欄檢查之前,即使所有執行緒都會到達柵欄也要跑證明,當前 TVM 的求解預算會被耗盡並返回部分計數,於是本來合法的 remap 與 sequence 輔助 kernel 直接以「Cannot prove a thread rectangle」編譯失敗。回退在兩套 ThreadSync 實現上同時進行,刪掉 #3291 引入的測試,補一個「迴圈對多數執行緒為空、256 執行緒仍全部到達柵欄」的歸約用例;商守衛、精確參與執行緒界、非矩形執行緒集這些正確性改進也一起撤回,新開的 #3732 留作遷移後重新最佳化的跟蹤單。這是明確標註為臨時措施的遷移讓路:TileLang 正在準備 TIRX 遷移以擺脫當前 TVM fork 依賴,配套的 TVM fork 最佳化提案已經關閉。
第二條線是昇騰在補地基。#1948 修的是宿主側一筆與運算元規模無關的固定損耗:Cython wrapper 每次呼叫都無條件構造 TVM Analyzer,實測在 910B3 上約 0.28 毫秒一次,而它只在形狀為符號表示式時才用到;改成惰性構造後,vec add 單塊從 442 至 467 微秒降到 105.1 微秒,32 塊從 441 至 493 微秒降到 108.0 微秒,kpool 壓縮從 447 至 450 微秒降到 111.3 微秒,作為下限的裸 ctypes 呼叫約 12 微秒沒變。同期 #1971 的 RFC 給出同步實現的替換路線:移植主線同步核心加獨立的 Z3 證明服務,用舊 TIR 適配層接進現有編譯流水,保留 TVM、排程、記憶體規劃、AscendC/PTO codegen 與執行 ABI,驗收透過後再切預設實現並刪除舊演算法——與 #1975 把整個 Ascend 倉遷到主線多後端架構是一件事的兩個落點。
第三條線在運算元庫。TileOPs 在 11-01 一天合入 13 筆,幾乎全部是效能補丁,且每筆都附 H200 實測表與 manifest 迴歸門檻。其中 #2808 把 GQA 稠密前向的程式選擇收成一條規則,相對最快候選程式的選擇遺憾從 1.818 降到 1.017(2 至 128 個 query 位置、586 次呼叫),並修掉主線一個真實錯誤:score scale 在乘上 log2(e) 後於 FP32 下消失,輸出變成 NaN 或錯值。
一、核心專案進展
1.1 主線:ThreadSync 的精確執行緒計數回退,為 TIRX 遷移讓路(#3733)
合於 11-01 14:50。被回退的 #3291 把精確執行緒計數移到滿塊柵欄檢查之前,並加入矩形認證,代價是即使全部執行緒都到柵欄也要跑一遍 CountSatisfyingValues;TVM 側求解預算耗盡後返回部分計數,remap 與 sequence 輔助 kernel 隨即以「[ThreadSync] Cannot prove a thread rectangle」失敗。迴歸在主線 cc903dcb9 加未改動的 TVM pin 4211874e9e 上覆現,新用例把它縮成「迴圈對多數執行緒為空、256 執行緒仍全部到達柵欄」。回退同時覆蓋共享與昇騰兩套實現,恢復滿塊快速路徑與舊的部分計數行為;原文明確這是 TIRX 遷移期間的臨時緩解,TVM fork 側的最佳化提案 tile-ai/tvm#89 已關閉,後續最佳化轉入 #3732。
來源:https://github.com/tile-ai/tilelang/pull/3733
1.2 主線:CUDA 跨 warp 全歸約合併為一次共享記憶體交換,256 執行緒從 6 道柵欄降到 2 道(#3715)
合於 10-30 20:26。舊實現在 CUDA 歸約標頭檔案 reduce.h 裡從 threads/2 遞迴到 scale,凡是 offset 不小於 32 的每一步都做一次完整的共享記憶體寫讀交換並配兩道柵欄:256 執行緒塊的三步(128、64、32)合計 3 次共享記憶體往返、6 道柵欄,512 執行緒塊是 4 次、8 道。改法是把所有跨 warp 步驟壓成一次交換(每元素一次 st.shared、共 2 道柵欄):先在 warp 內做 scale 小於 32 的那一段,再讓每個 warp 內組裡不超過 cross_radix 的若干 lane 協作摺疊,暫存器壓力有上界。標量與批次兩條路徑都改了。
來源:https://github.com/tile-ai/tilelang/pull/3715
1.3 主線:PTO 後端的 T.Parallel 向量化落地第一里程碑,逐元素運算元走 VMI 指令(#3512)
合於 10-30 01:29。目標是讓使用者在 T.SimdVF(lanes=L) 裡寫標量並行的逐元素 T.Parallel,由編譯器向量化成 PTO VMI 指令,再經 PTOAS 在昇騰 950 上執行。分工寫得很清楚:TileLang 側負責邏輯 lane 佈局推導、向量化 pass 流水線、VMI 指令生成與 Fragment 邏輯對映;PTOAS 側負責物理暫存器切分、掩碼具體化、Fragment 緩衝與向量暫存器分配、mem_bar 插入。路線圖分五個里程碑,本筆是 M1 核心框架(逐元素計算,覆蓋 extent 能否被 lanes 整除兩種情況),M2 控制流、M3 reduce、M4 UB 容量規劃與暫存器分配、M5 非連續訪問與更寬的 PACK 資料型別依次跟進。
來源:https://github.com/tile-ai/tilelang/pull/3512
1.4 主線:昇騰 DeepGEMM 的 scale 緩衝改用寫計數多緩衝,資料載入可提到條件載入之前(#3629)
合於 10-30 05:48。DeepGEMM 會在多個 K 塊之間複用 scale factor,本筆新增顯式的寫計數多緩衝模式,讓 scale 緩衝的版本按真實寫入推進:每個 owner 迴圈迭代裡只在第一個實際執行的寫任務處推進一次,後續寫入共享該版本,沒有寫入的迭代沿用舊值,支援多個不相交 owner 與共享時鐘。這樣資料載入可以排到條件化的 scale 載入之前;同時把 DeepGEMM 的 K 迴圈攤平,在 L1 與 L0 上都先取資料再取 scale,並補了守衛更新、巢狀更新、保留讀、flag 與緩衝複用以及 FP4 的 K 主序與 MN 主序帶尾塊和環形迴繞的用例。該模式要求顯式版本數,並假定每個 owner 迭代最多一次更新。
來源:https://github.com/tile-ai/tilelang/pull/3629
1.5 主線:兩筆修復合入,CPU 標量 GEMM 只清目標區域、負索引合法化跳過不可達分支(#3709 / #3712)
兩筆都在 10-30 合入,屬上一期在審條目的落地。#3709 讓 CPU 標量 GEMM 只清理目標區域,此前會連帶清掉不屬於本次輸出的範圍;#3712 讓 LegalizeNegativeIndex 跳過不可達分支,避免對不會執行的程式碼做負索引改寫。這兩條與上一期報告的 #3708、#3710 屬同一批缺陷清理,其中 #3708 的位寬改 int64、#3710 的關快取時繞過兩層程序內快取仍在審。
來源:https://github.com/tile-ai/tilelang/pull/3709、https://github.com/tile-ai/tilelang/pull/3712
1.6 主線:語言層清掉 proxy 裡兩個無人使用的 Ref 佔位(#3722)
合於 10-31 17:31。tilelang/language/proxy.py 裡同時存在兩個 Ref:一個給型別檢查用的 tirx.Var 樁,一個空的執行時類,兩者在樹內都沒有消費者;對外的 T.Ref 實際來自 eager.builder,那裡的實現才是宏引用引數的本體。本筆刪掉兩個佔位、_T 型別變數以及隨之無用的 Generic、TypeVar 匯入,保留 T.Ref 與宏引用行為,只移除 tilelang.language.proxy.Ref 這條無人使用的直接匯入路徑。前端構造測試 179 項透過、6 項跳過,執行時身份檢查確認 T.Ref 仍是 eager.builder.Ref。
來源:https://github.com/tile-ai/tilelang/pull/3722
1.7 新開缺陷與在審:TMEM 靜默丟結果、ROCm 符號找不到、TF32 就近偶數舍入(#3721 / #3711 / #3714)
#3721 是本視窗最值得警惕的一條:在 sm_100a 上對 TMEM 做 T.gemm,當 threads/32 不能整除 N 時,編譯能過、PTX 裡卻一條 tcgen05.mma 都沒有,T.copy(Ct, Cl) 讀的是沒人寫過的 TMEM,乘積被靜默丟棄;復現裡 threads=256 出一條,threads=384 出零條,作者懷疑 src/cuda/op/gemm.cc 給 TCGEN05 返回 {1, num_warps} 分割槽時缺了覆蓋性檢查。#3711 報告 ROCm 7.14 與 10.0 的 PyTorch wheel 下 kernel 報「named symbol not found」。#3714 在審,為 fp32 到 tf32 的轉換補上就近偶數舍入(cvt.rn.tf32.f32 與飽和形式),非飽和要求 SM90 以上、飽和要求 SM100 以上,關閉對應的功能請求 #3713。
來源:https://github.com/tile-ai/tilelang/issues/3721、https://github.com/tile-ai/tilelang/issues/3711、https://github.com/tile-ai/tilelang/pull/3714
二、多後端適配(昇騰 / 海光 / 摩爾執行緒 / 沐曦)
2.1 昇騰:Cython wrapper 惰性構造 TVM Analyzer,910B3 單次呼叫從約 450 微秒降到約 106 微秒(#1948)
合於 10-30 07:44。CythonKernelWrapper.forward() 此前每次呼叫都無條件構造一個 TVM Analyzer(),在 910B3 上是約 0.28 毫秒的固定宿主開銷,而該物件只在 result 或 workspace 形狀為符號表示式時才用到;靜態 shape 的 kernel 完全用不上,卻同樣付這筆錢。改法是把構造延遲到 tir.PrimExpr 分支內並按需複用。實測:vec add 單塊 442 至 467 微秒降到 105.1 微秒,32 塊 441 至 493 微秒降到 108.0 微秒,kpool 壓縮 447 至 450 微秒降到 111.3 微秒,裸 ctypes 呼叫約 12.5 微秒降到約 12 微秒。新增兩例迴歸,分別驗證靜態路徑不構造、符號路徑按需構造。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1948
2.2 昇騰:RFC 提出把同步實現換成主線核心加獨立 Z3 證明服務(#1971)
本 RFC 不預設提速結論,只固定比較基線:Ascend 倉在 TileLang 43b0b69 與 TVM c2921fd 上、IR 基礎是舊 tvm.tir;主線在 TileLang 1c0f402 與 TVM 4211874e9e 上、IR 基礎已是 tvm.tirx 與 tvm.s_tir,物件與註冊體系也換成新 FFI 與反射設施。因此遷移不能簡化為複製一個 InsertSync 檔案。建議路線是移植主線同步核心與獨立 Z3 證明服務,透過舊 TIR 適配層接入現有編譯流水,保留現有 TVM、排程、記憶體規劃、AscendC/PTO codegen 與執行 ABI,完成支援範圍、正確性、效能與安裝驗收後再切換預設實現並刪除舊同步演算法。配套 PR 為 #1965。
來源:https://github.com/tile-ai/tilelang-ascend/issues/1971
2.3 昇騰在審:把 TileLang-Ascend 整體遷到主線多後端架構(#1975)
10-31 提出並持續更新,目標是按提案 #1911 把 TileLang-Ascend 遷移到 TileLang 主線的多後端架構上;描述裡給出兩份差異對照,一份對 ascendc_pto 分支,一份對主線。這條與 2.2 的 RFC 是同一條收斂路徑的兩個層面:一條改同步演算法,一條改倉的骨架。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1975
2.4 昇騰在審:GLM-5.3 kpool 的 fp8_mqa_logits 示例,雙 kernel 加 AIV/Cube 分工(#1966)
對照主倉 kpool 目錄下的 GLM-5.3 fp8_mqa_logits 示例做昇騰版。運算元語義是 GLM-5.3-Flash 稀疏注意力 indexer 的池化歷史打分:對每個 query 行,在有效 pool 區間內按分頁表從 64-pool 頁讀 E4M3 key(uint8 承載位元組加 fp32 scale),反量化後算 32 個 indexer 頭的 MQA 打分,乘 per-pool scale,經 ReLU gate 與按頭加權歸約為每個 pool 一個 FP32 logit,範圍外或無效頁輸出負無窮。實現是雙 kernel:mask kernel 按 64-pool tile 生成逐 lane 的 0 或負無窮掩碼以替換宿主預填,main kernel 每行每頁一個 block、兩個 AIV 各處理 32-pool 半頁,Cube 讀整頁做一次 64×32×128 的 GEMM。頭數 32、維度 128、頁大小 64。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1966
2.5 昇騰:四筆文件與 API 事實核查合入,clamp 在 PTO 上真正生效(#1899 / #1893 / #1620 / #1600)
四筆都在 10-30 合入,屬持續進行的 API 事實核查批次:#1899 核對 T.tile.bitwise_xor 文件(#1681 的重提),#1893 核對 T.tile.bitwise_rshift(#1676 的重提),#1620 讓 T.tile.clamp 在 PTO 上真正生效並同步文件,同時補測試與 API 文件,#1600 更新 axpy 的 docstring、加強測試並補 API 文件。同批次另有 #1968、#1967、#1974 在審,分別是 sigmoid 與 round 計數折半、符號形狀解析快取、整塊行切分下 tile 尺寸被錯誤折半的修復。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1620、https://github.com/tile-ai/tilelang-ascend/pull/1899
2.6 昇騰:日測 2970 用例全部透過,11-01 有一次 workflow 失敗(#1977 / #1976)
11-02 06:42 的日測記錄為 2970 項、透過 2970 項、失敗 0 項,透過率 100%;10-31 06:56 的一期同樣是 2970 項全過。中間 11-01 05:08 有一期標記為 workflow test job failure,與本視窗內 #1975 的遷移改動、#1971 的 RFC 同期出現,屬需要留意的波動而非穩定迴歸。
來源:https://github.com/tile-ai/tilelang-ascend/issues/1977、https://github.com/tile-ai/tilelang-ascend/issues/1976
2.7 海光:MLS 邊界與持續 WASP 同步修復合入,gfx948 相容在審(#31 / #32)
#31 合於 10-30 10:56,三件事:改進動態 MLS 邊界推斷與 MatrixStore 的 stride 處理;修 MMAC 與 blockscaled GEMM 的 clear_accum 下沉;在持續 WASP 路徑改用 M 分割槽消費者與獨立的 A/B 柵欄,同時保留 slot 與 phase 輪轉並對尾部消費者加同步。#32 在審,把 gfx948 當作 gfx946 相容處理,讓既有指令能直接在新型號上走通。前一視窗報道的普通複製 LDS swizzle(#30)已在同期落地。
來源:https://github.com/tile-ai/tilelang-hygon/pull/31、https://github.com/tile-ai/tilelang-hygon/pull/32
2.8 mlir-ascend:CI 清掉內部快取服務引用,Codex 工作流在審(#206 / #208)
#206 合於 10-31 07:24,把 CI 裡指向內部快取服務的內容清掉,屬可復現性清理。#208 在審,為 Operator Harness 加首批 Codex 工作流:用迭代提問完成專案初始化、把已有運算元倉接入、以及從持久化上下文恢復開發、調優或分析;三個 skill 共用 .workspace/ 下的可編輯 Markdown 記錄,任務交接區分「當前候選」與「已驗證的最佳候選」,保留失敗與回退實驗。作者明確本次沒有真實 NPU 實測,NPU 執行適配與 npm 分發列為後續工作。
來源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/206、https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
2.9 摩爾執行緒與沐曦視窗內無推送(11-02)
tilelang-metax 最後推送停在 10-26,tilelang-musa 停在 09-30,兩倉在本視窗的 72 小時內都沒有提交。這意味著本期的後端動態集中在昇騰、海光與 mlir-ascend 三側。
來源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa
三、生態與採用方(TileOPs / TileFoundry / TileKernels / FlashQLA)
3.1 TileOPs:72 小時 15 筆合入,13 筆壓在 11-01,幾乎每筆都帶實測表與清單迴歸門檻(11-01)
v0.1.0 釋出後,TileOPs 在 11-01 一天合入 13 筆,加上 10-30 的 #2794 與 10-31 的 #2795,72 小時共 15 筆。改動幾乎全部是 perf 類,覆蓋注意力、GEMM、逐元素、歸約與 MoE,每一筆都給出 H200 上的 busy 與 back-to-back 兩組時間、或與 vLLM、FlashInfer 的對照,並宣告 manifest 裡沒有已測行回退。這種「帶表合入」的方式說明運算元庫已經進入按 shape 逐行校準的階段,而不是補功能。
來源:https://github.com/tile-ai/TileOPs/commits/main
3.2 TileOPs:GQA 稠密前向用一條實測規則選 decode 或 prefill 程式,選擇遺憾從 1.818 降到 1.017(#2808)
合於 11-01 13:04。GQADenseFwdOp 與 GQADenseWithLSEFwdOp 改由同一個規則 AttentionCall.dense_decode_program 決定走哪套程式,輸入是 query 與 key 長度、batch、頭分組、head dim、RoPE 與 SM 數:2 至 64 個 query 位置走 warp 專用 decode,一個 KV 頭組內每個位置打包進 64 行塊;單個 query 位置時按 32 段切分填滿一波、或走 1 序列 1 至 2 個 KV 頭的 batch-1 程式。相對最快可用程式的選擇遺憾(幾何均值與最大值):單位置 1.068 與 1.87 降到 1.018 與 1.69,2 至 128 位置 1.818 與 25.2 降到 1.017 與 1.81。同時修掉主線兩個錯誤:score scale 乘上 log2(e) 後在 FP32 下消失導致 NaN 或錯值;D256 多位置呼叫在 SM90 上因共享記憶體不足啟動失敗。
來源:https://github.com/tile-ai/TileOPs/pull/2808
3.3 TileOPs:分頁 MHA decode 擴到共享記憶體裝得下的每個 head dim,GQA 分頁前向去掉舊 decode 核(#2807 / #2797)
#2807 合於 11-01 12:30,把分頁 MHA decode 在「一個塊的共享記憶體裝得下」的前提下擴到每個 head dim,此前遇到較大 head dim 會直接不走這條路徑。#2797 合於同日 03:54,從 GQAPagedFwdOp 裡刪掉舊的 paged decode kernel,並把分頁 MHA decode 的上限壓到能編譯出來的 head dim 上,等於把前一版留下的越界能力收回。兩條合起來是一次「先擴能力、再削冗餘」的收尾。
來源:https://github.com/tile-ai/TileOPs/pull/2807、https://github.com/tile-ai/TileOPs/pull/2797
3.4 TileOPs:coop2 生產者收成單 warp、新增 192 寬轉置塊,W4A8 部分一波改成一 CTA 一 tile(#2806 / #2805)
#2806 合於 11-01 12:02,把 Hopper GEMM 的 coop2 生產者從多 warp 收成一個載入 warp,並在共享記憶體允許三級流水時提供 192 寬的均勻轉置塊;H200 上 wgrad-q-b 從 384.9 微秒降到 382.3,mlp-up fp16 從 317.9 微秒降到 311.9(back-to-back 分別 458.5 到 449.6、369.0 到 367.3),作者宣告沒有已測 GEMM 行回退,本筆屬 #2791 目標的一部分,雙卡確認仍未完成。#2805 合於同日 10:36,讓填滿至少四分之三 SM 的部分一波 W4A8 prefill 柵格改成一 CTA 一 tile 而非 stream-K;llama-8b-up-b33 從 23.52 微秒降到 17.98(vLLM 20.42),b128 從 31.33 降到 23.39(vLLM 24.05),清單外 M 33 至 128 的 44 組 N/K 掃描中改排程形狀快 5% 至 36%、中位 18%、無一變慢。
來源:https://github.com/tile-ai/TileOPs/pull/2806、https://github.com/tile-ai/TileOPs/pull/2805
3.5 TileOPs:編譯快取按構建目標鍵控,運算元層統一檢查輸出與輸入重疊(#2802 / #2801)
#2802 合於 11-01 11:31,修的是編譯程式快取跨裝置複用的問題:構建器與 Op 條目都按構建目標做鍵,懶編譯與自動調優留在預期裝置上,驗收為 99 專案標隔離、供應商與分組 GEMM 測試。#2801 合於同日 10:33,把「輸出不得與任何其他呼叫張量共享位元組」的檢查提到運算元層,由 manifest 的 overlaps 欄位按運算元宣告例外,刪掉了 MoE unpermute、MoE 分組 GEMM、持續 GEMM 模板等四處 kernel 側檢查;順帶修掉主線一個競態:IndexedExpertMLPFwdOp 允許 out 與 topk_weights 重疊,而後者在寫入時仍被歸約讀取。MoEPostPermuteFwdOp 帶 out 的宿主耗時從 51.48 微秒降到 39.33。
來源:https://github.com/tile-ai/TileOPs/pull/2802、https://github.com/tile-ai/TileOPs/pull/2801
3.6 TileOPs 在修缺陷:Norm 四運算元在遠離零的行丟精度、約 1e36 溢位成 NaN(#2812)
11-01 新開。LayerNorm、AdaLayerNorm(含 -Zero)、GroupNorm 與 BatchNorm 目前都用原始行的 fp32 直接求和取均值與方差(BatchNorm 用 E[x²] 減均值平方),這正是 #2773 為 FusedAddLayerNorm 修過、#2811 為 InstanceNorm 訓練核在修的同一類缺陷。表現是:常數 10000.1 的 fp32 行歸一化後輸出約等於正負 weight 而不是 bias,10000 加隨機數的行在 fp32 容差下對不上 float64 參考;接近 1e36 的行求和溢位,fp32 與 bf16 下全部輸出 NaN。受影響路徑覆蓋 fragment、共享記憶體走查、一行一 warp、分塊與流式等多種實現。
來源:https://github.com/tile-ai/TileOPs/issues/2812
3.7 TileFoundry:lowering 把 issuing mesh 與視窗切分收進一處決策,83 個排程夾具位元組一致(#280 / #282)
#280 合於 10-31 00:47,把「每次呼叫發射什麼」的決策收進一處;#282 合於 11-01 11:47,繼續把「由哪個 mesh 發射」與「視窗怎麼切」也收到一處:analysis/emission.py 新增 copy_mesh、issue_owner、issue_mesh,copy_mesh 與 issue_owner 共用一個助手,兩者唯一的差別(Partial 怎麼讀)只留在一處;三處重複的 Slice 切分改由 _cut 助手統一。驗收方式是位元組級對照:83 個排程夾具裡 75 個 finalize 輸出與主線完全一致,其餘 8 個在主線上同樣失敗;docs/spec/passes.md 第 7.1 節把這套 mesh 決策寫進規範。
來源:https://github.com/tile-ai/TileFoundry/pull/282、https://github.com/tile-ai/TileFoundry/pull/280
3.8 TileFoundry:站點發布 0.3 版排程文章《Scheduling, Designed for Agents》(#4 / #281)
站點倉在 10-30 合入 #4,釋出 0.3 版文章《Scheduling, Designed for Agents》,同時補了文章與配圖的目錄結構、公式的 MathJax 支援、圖注與表格樣式以及每個圖表錨點,首頁 Blog 卡片改為開啟部落格索引,索引列出兩篇文章;upstream-ref 未變,站點仍渲染 v0.0.3。倉庫側 #281 於 10-31 在 README 裡通告這篇文章。對 TileFoundry 這條線來說,這是把「面向 agent 的排程」從程式碼與釋出說明搬到公開敘述的一步。
來源:https://github.com/tile-ai/TileFoundry.github.io/pull/4、https://github.com/tile-ai/TileFoundry/pull/281
3.9 採用方:FlashQLA 合入 SM120 的 Dual-CP 前反向,2 張 RTX 5090 上對 32K 至 128K 序列與 Inter-CP 對比(#51)
10-30 03:02 合入。改動為 SM120 增加 Dual-CP 的前向與反向:前向複用已有的 state-scan kernel 做卡級狀態聚合,有序 M 乘積按 32 列一塊計算以控制在共享記憶體內;反向用轉置後的 M 逆序傳播狀態,複用前向快取下來的 M 矩陣;同時修掉反向 kernel 裡一處共享記憶體競態,辦法是讓第 7 階段與 dV 儲存 warp 同步。效能在 2 張 RTX 5090 上以 BF16、chunk 32、序列長 32K 至 128K 對比 Inter-CP,每例 3 次中位數、每次 25 次預熱加 100 次計時。
來源:https://github.com/QwenLM/FlashQLA/pull/51
3.10 採用方:TileKernels 自 10-14 起無推送(11-02)
深度求索的 TileKernels 最後推送停在 10-14,本視窗無變化。採用方這一側本期的有效動態全部來自通義 FlashQLA。
來源:https://github.com/deepseek-ai/TileKernels
四、社群、教程與活動
4.1 RISC-V 研究鏈成形:三個倉把 Numba CUDA Python 轉 TileLang,再把受限 PrimFunc 落到單 hart RVV C(10-31)
社群研究者 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 官方註冊的後端,也不是通用 GPU 模型模擬器,依賴 threadIdx 的定址、跨 hart 同步與未支援的 IR 一律拒絕,block 順序執行,shared 與 fragment 對映為單 hart 區域性陣列;研究環境為 TileLang 0.1.15、Clang 23.1.3、QEMU 11.1.0,不保證其他版本相容,且明確上游執行緒級產物不能自動視為該後端可接受的輸入。
來源:https://github.com/zishu0131-png/tilelang-riscv
4.2 PTO-ISA/tilelangGym 從 PTO-ISA/tilelang 獨立成倉,主迴歸門是昇騰 SimtVF 路徑(10-30)
這套敏感性與迴歸測試集從 PTO-ISA/tilelang 的 examples/ascend/t_parallel_st_suite 拆出來,理由是讓測試集能獨立釘住 TileLang 與 PTOAS 的版本。當前三條主路徑:昇騰 SoftLib SimtVF(target=ascend、cython,主迴歸門,在 remat 版本上全過)、VectorizeParallelToPTO 與 VL 的 vload/vstore,以及 benches 下用 @pto.jit 寫的手工 SIMD DSL 對照實現;每個用例目錄裡帶 config.toml 與 VERDICT.md,另有周度與 CI 迴歸指令碼。
來源:https://github.com/PTO-ISA/tilelangGym
4.3 新聞側:Google News 中英文查詢、Hacker News、arXiv 均無視窗內實質命中(11-02)
三天的 Google News 中英文查詢沒有返回任何落在視窗內的 TileLang 相關報道,Hacker News 與 arXiv 同樣沒有新條目;候選清單裡的 20 條全部來自 GitHub 倉庫推送,屬常規程式碼活動而非新聞。本期的資訊面因此完全由倉庫側與社群倉構成。
五、趨勢觀察
這一視窗最值得記錄的判斷來自上游自己:寧可回退正確性改進,也要把 ThreadSync 的複雜度從關鍵路徑上撤下來。回退說明裡寫清了代價——商守衛、精確參與執行緒界、非矩形執行緒集這三類改進一併撤回,只保留舊的保守路徑,#3732 留待 TIRX 遷移完成後再談最佳化。一個編譯器專案願意公開寫「我們的證明器預算不夠,先退回保守實現」,比再堆多少功能都更能說明當前階段的真實約束:TIRX 遷移是主線第一優先順序,其他最佳化都要給它讓路。
與之呼應的是昇騰側兩份「向主線收斂」的文字:#1971 的 RFC 要把同步演算法換成主線核心加獨立 Z3 證明服務,先調研、後替換、中間留舊 TIR 適配層;#1975 直接把整個 Ascend 倉遷到主線多後端架構。兩條路都要付出遷移成本,於是近期的可見工作會偏向地基類:#1948 把每次呼叫的固定宿主開銷從約 450 微秒壓到約 106 微秒,四筆文件事實核查,CI 清理內部快取引用,都是這類工作。
TileOPs 的節奏則從「發版」切到「調優」。72 小時 15 筆合入裡 13 筆壓在同一天,幾乎每筆都附 H200 實測表與「manifest 無回退」宣告,並且開始用選擇遺憾這類指標描述排程規則的好壞(#2808 從 1.818 到 1.017)。這意味著運算元庫的競爭點已經從「有沒有這個運算元」變成「每條 shape 上是否都挑對了程式」,而這恰好是 TileLang 這類 DSL 相對手寫 kernel 的長期價值所在。
生態側出現兩條新分叉。一條是硬體方言:海光把 gfx948 當 gfx946 相容、昇騰繼續補 PTO 文件與符號形狀。另一條是語言出口:RISC-V 研究鏈把受限 PrimFunc 降到單 hart RVV C,星標很低、作者也明確宣告不是官方後端,但它把 TileLang 的 IR 當作可移植資產來用,與昇騰、海光的倉級分支適配是同一類訊號。多條分叉同時存在,上游 TIRX 遷移能否給出穩定的 IR 邊界,將決定這些下游投入是沉澱還是返工。
附、素材與核查說明
信源核查表
| 信源 | 核查結果 |
|---|---|
| tile-ai 組織(29 倉庫) | 視窗內推送 10 個倉庫 |
| 主倉 tilelang | 主分支 7 筆合入;新開 issue 4 條;文件站 3 筆 |
| TileOPs | 主分支 15 筆合入(13 筆在 11-01);新開 issue 3 條;站點 1 筆 |
| TileFoundry | 3 筆合入;站點 1 筆 |
| tilelang-ascend | 5 筆合入;新開 issue 4 條(含日測記錄與 RFC) |
| tilelang-hygon | 2 筆合入;1 筆在審新增 |
| tilelang-mlir-ascend | 1 筆合入;1 筆在審新增 |
| tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack | 視窗內無推送 |
| 釋出頁(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / mlir-ascend) | 視窗內無新版本 |
| FlashQLA | 1 筆合入(#51) |
| TileKernels | 自 10-14 起無推送 |
| RISC-V 研究鏈三倉 / PTO-ISA/tilelangGym | 有推送,研究鏈三倉同日建鏈 |
| Google News / Hacker News / arXiv | 無視窗內實質命中 |
完整信源清單
- [1] tilelang ThreadSync 回退(#3733)— https://github.com/tile-ai/tilelang/pull/3733
- [2] ThreadSync 計數迴歸跟蹤(#3732)— https://github.com/tile-ai/tilelang/issues/3732
- [3] tilelang CUDA 跨 warp 全歸約(#3715)— https://github.com/tile-ai/tilelang/pull/3715
- [4] tilelang PTO 的 T.Parallel 向量化(#3512)— https://github.com/tile-ai/tilelang/pull/3512
- [5] tilelang 昇騰 DeepGEMM 寫計數多緩衝(#3629)— https://github.com/tile-ai/tilelang/pull/3629
- [6] tilelang CPU 標量 GEMM 清理範圍(#3709)— https://github.com/tile-ai/tilelang/pull/3709
- [7] tilelang 負索引合法化跳過不可達分支(#3712)— https://github.com/tile-ai/tilelang/pull/3712
- [8] tilelang 語言層清理 Ref 佔位(#3722)— https://github.com/tile-ai/tilelang/pull/3722
- [9] tilelang TMEM 靜默丟結果(#3721)— https://github.com/tile-ai/tilelang/issues/3721
- [10] tilelang ROCm 符號找不到(#3711)— https://github.com/tile-ai/tilelang/issues/3711
- [11] tilelang TF32 就近偶數舍入(#3714)— https://github.com/tile-ai/tilelang/pull/3714
- [12] tilelang 浮點轉 TF32 舍入功能請求(#3713)— https://github.com/tile-ai/tilelang/issues/3713
- [13] tilelang-ascend Cython wrapper 惰性構造 Analyzer(#1948)— https://github.com/tile-ai/tilelang-ascend/pull/1948
- [14] tilelang-ascend 自動同步與 Z3 遷移 RFC(#1971)— https://github.com/tile-ai/tilelang-ascend/issues/1971
- [15] tilelang-ascend 遷移到主線多後端架構(#1975)— https://github.com/tile-ai/tilelang-ascend/pull/1975
- [16] tilelang-ascend GLM-5.3 kpool fp8_mqa_logits 示例(#1966)— https://github.com/tile-ai/tilelang-ascend/pull/1966
- [17] tilelang-ascend clamp 在 PTO 生效(#1620)— https://github.com/tile-ai/tilelang-ascend/pull/1620
- [18] tilelang-ascend bitwise_xor 文件核查(#1899)— https://github.com/tile-ai/tilelang-ascend/pull/1899
- [19] tilelang-ascend bitwise_rshift 文件核查(#1893)— https://github.com/tile-ai/tilelang-ascend/pull/1893
- [20] tilelang-ascend axpy 文件與測試(#1600)— https://github.com/tile-ai/tilelang-ascend/pull/1600
- [21] tilelang-ascend 日測 100% 透過(#1977)— https://github.com/tile-ai/tilelang-ascend/issues/1977
- [22] tilelang-ascend workflow 失敗記錄(#1976)— https://github.com/tile-ai/tilelang-ascend/issues/1976
- [23] tilelang-hygon MLS 邊界與
WASP 同步(#31)— https://github.com/tile-ai/tilelang-hygon/pull/31
- [24] tilelang-hygon gfx948 相容(#32)— https://github.com/tile-ai/tilelang-hygon/pull/32
- [25] tilelang-mlir-ascend CI 清理內部快取引用(#206)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/206
- [26] tilelang-mlir-ascend Codex 工作流(#208)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
- [27] TileOPs GQA 稠密前向程式選擇(#2808)— https://github.com/tile-ai/TileOPs/pull/2808
- [28] TileOPs 分頁 MHA decode 擴 head dim(#2807)— https://github.com/tile-ai/TileOPs/pull/2807
- [29] TileOPs GQA 分頁前向去掉舊 decode 核(#2797)— https://github.com/tile-ai/TileOPs/pull/2797
- [30] TileOPs coop2 單 warp 與 192 寬塊(#2806)— https://github.com/tile-ai/TileOPs/pull/2806
- [31] TileOPs W4A8 部分一波一 CTA 一 tile(#2805)— https://github.com/tile-ai/TileOPs/pull/2805
- [32] TileOPs 編譯快取按目標鍵控(#2802)— https://github.com/tile-ai/TileOPs/pull/2802
- [33] TileOPs 運算元層重疊檢查(#2801)— https://github.com/tile-ai/TileOPs/pull/2801
- [34] TileOPs Norm 精度與溢位缺陷(#2812)— https://github.com/tile-ai/TileOPs/issues/2812
- [35] TileOPs 主分支提交記錄 — https://github.com/tile-ai/TileOPs/commits/main
- [36] TileFoundry issuing mesh 與視窗切分(#282)— https://github.com/tile-ai/TileFoundry/pull/282
- [37] TileFoundry 發射決策收攏(#280)— https://github.com/tile-ai/TileFoundry/pull/280
- [38] TileFoundry 0.3 排程文章釋出(#4)— https://github.com/tile-ai/TileFoundry.github.io/pull/4
- [39] TileFoundry README 通告排程文章(#281)— https://github.com/tile-ai/TileFoundry/pull/281
- [40] FlashQLA SM120 Dual-CP(#51)— https://github.com/QwenLM/FlashQLA/pull/51
- [41] 深度求索 TileKernels — https://github.com/deepseek-ai/TileKernels
- [42] RISC-V 研究鏈總入口 — https://github.com/zishu0131-png/tilelang-riscv
- [43] Numba 到 TileLang 轉換倉 — https://github.com/zishu0131-png/numba2tilelang
- [44] TileLang PrimFunc 到 RVV C 消費器 — https://github.com/zishu0131-png/tilelang2rvv
- [45] PTO-ISA/tilelangGym — https://github.com/PTO-ISA/tilelangGym
- [46] tilelang 釋出頁 — https://github.com/tile-ai/tilelang/releases
- [47] TileOPs 釋出頁 — https://github.com/tile-ai/TileOPs/releases
- **[48]
** TileFoundry 釋出頁 — https://github.com/tile-ai/TileFoundry/releases
- [49] tilelang-metax — https://github.com/tile-ai/tilelang-metax
- [50] tilelang-musa — https://github.com/tile-ai/tilelang-musa
- [51] Google News 檢索 — https://news.google.com/