TileLang 每日動態報告(2026-11-04)
調研視窗:2026-11-03 07:00 ~ 2026-11-04 07:00(北京時間,24 小時;上一期截點為 11-03 07:00)。
信源:GitHub tile-ai 組織逐倉核查(29 倉庫,視窗內推送 9 個)+ 合入與在審真值檢索(合入 8 筆、在審更新 48 筆)+ 新增 issue 全量核對(7 條)+ 主分支逐筆提交核查(tilelang 5、TileOPs 1、TileFoundry 1、VS Code 擴充套件 1、文件站 3)+ 釋出頁與 tag 核查(5 倉,視窗內無新版本)+ 採用方兩倉 + Google News 中英文查詢(走代理)+ Hacker News + arXiv。
本期索引
- 今日重點:ThreadSync 跨儲存域複用部分屏障,把 #3643 修正確性時引入的效能回退補回(#3761);清理 3384 行遺留 TVM 解析器分支(#3736);自動調優輸入供給改按當前裝置(#3546);TileFoundry 同步計劃正式落到 CUDA lowering(#285);24 小時合入 8 筆、在審更新 48 筆、新增 issue 7 條
- 一、核心專案進展
- 1.1 主線:ThreadSync 允許跨儲存域複用部分屏障,B300 上屏障佔用從 12 回到 16 塊/SM(#3761)(11-04)
- 1.2 主線:移除遺留的 TVMScript AST 與解析器分支,一次刪掉 3384 行(#3736)(11-03)
- 1.3 主線:自動調優的輸入供給改用當前裝置,昇騰與 MPS 上的調優恢復可用(#3546)(11-03)
- 1.4 在審:條件中的標量 cast 不再壓縮資料向量寬度,承接昨日 #3734(#3759)(11-03)
- 1.5 在審:兩處使用者可見錯誤改為明確報錯(#3765 / #3767)(11-03)
- 1.6 在審:帶 padding 的共享 tile 無法整體複製時回退,2-CTA 的 TMA 位元組按對等 leader 記賬(#3738 / #3766)(11-03)
- 1.7 在審:B300 的 sm_103a 適配兩條線,tcgen05 的 int8 選擇與 CuTeDSL 的 TMA 架構門(#3770 / #3772)(11-03)
- 二、多後端適配
- 2.1 昇騰:3D GM 到 L1 的批次 ND2NZ 複製合入,一次描述符覆蓋整批(#3689)(11-03)
- 2.2 昇騰在審:TileLang-Ascend 整倉遷入主線多後端架構,2887 個檔案(#1975)(11-03)
- 2.3 昇騰在審:copy_pipe_to_l1 的同步標註按執行引擎拆分,修 A2/A3 上 10/10 復現的錯結果(#1981)(11-03)
- 2.4 昇騰在審:首 Parallel 軸廣播走標量讀、Chunk O 前向遷移、MHA 示例(#1979 / #1866 / #1879)(11-03)
- 2.5 mlir-ascend:Codex 工作流與任務交接合入 harness 分支(#208)(11-03)
- 2.6 海光:Distributed 在審 PR 有更新,主分支無新提交(#19)(11-03)
- 2.7 沐曦、摩爾執行緒視窗內無推送(11-03)
- 三、生態與採用方
- 3.1 TileFoundry:同步計劃落成 CUDA lowering 並反向校驗,合入主線(#285)(11-04)
- 3.2 TileFoundry 在審:wgmma 同步計劃、FSDP 與 HSDP 訓練語義進 HIR 求值器(#286 / #284)(11-03)
- 3.3 TileOPs:方差合併從第一個部分量起算,常量輸入的方差不再出 NaN(#2821)(11-03)
- 3.4 TileOPs 在審:GEMM 加可選 bias、FP8 分組 GEMM 支援連續佈局(#2823 / #2824)(11-03)
- 3.5 TileOPs 在審:W4A16 與稠密 GEMM 的選擇器改成單一代價模型重擬合(#2825 / #2826)(11-03)
- 3.6 TileOPs 在審:短 varlen 按 KV 組打包、L2 flush 改為只讀(#2819 / #2822)(11-03)
- 3.7 採用方:FlashQLA 與 TileKernels 無推送,但 #3761 修的是 TileKernels 的 fused 核心(11-04)
- 四、社群、教程與活動
- 4.1 單日 7 條新 issue:一條生產級 CUDA graph 問題,四條來自同一外部報告人(11-03)
- 4.2 昇騰 950 的 vcvt 飽和語義與 mlir-ascend 日測(#3762 / #1980)(11-03)
- 4.3 文件站三次更新,VS Code 擴充套件改名 vsix(11-03)
- 4.4 新聞側:三通道無視窗內命中,外媒的昇騰訓練報道口徑不一(11-03)
- 五、趨勢觀察
- 附、素材與核查說明
今日重點
這一天的主線是一筆「把效能要回來」的修復。上一視窗合入的 #3643 解決了一個真實的正確性問題:分別呼叫的 ThreadSync("shared") 與 ThreadSync("shared.dyn") 可能把同一個命名屏障 ID 分配給互不相交的執行緒組。修法是讓後一次呼叫保留前一次已經用掉的 ID,但這條保留規則同時擋住了另一件事——同一組執行緒在兩個作用域之間本來可以複用同一個 ID。代價直接落在 TileKernels 的 mhc_pre_big_fuse 核心上:這個核心每塊 96 個執行緒,執行緒 32 到 95 跑兩段式 cp.async 流水線,非同步等待之後的同步在第一遍拿到 ID 3,同一批 64 個執行緒的緩衝依賴同步在第二遍拿到 ID 4,多出來的那個 ID 讓編譯產物宣告的屏障資源從 4 變成 5;在 B300 上跑隱藏維 1280 時,屏障佔用上限因此從 16 塊/SM 降到 12 塊/SM,暫存器、共享記憶體與同步指令數都沒有變化。#3761 的處理辦法是讓 ThreadSync 接受一個有序的作用域列表,CUDA 流水線改成 ThreadSync(["shared", "shared.dyn"]),同步仍按原來的作用域順序規劃與插入,但部分屏障的 ID 在兩個作用域都處理完之後才統一分配;跨作用域共用一個 ID 之前,用一個全新的分析器證明參與謂詞在整個候選執行緒範圍內成立,分支與 Bind 上下文由 ConstrVisitor 提供,執行緒相關的謂詞按證明義務處理而不是預設成立,已證密的密集組直接用包圍盒體積、跳過部分執行緒數的列舉,未證明的組則把作用域留在鍵裡並讓參與者計數一起入鍵,顯式編號的屏障(含更早一次 ThreadSync 分配的 ID)繼續保留。改動 401 行、涉及 4 個檔案。
作者在描述裡做了一個很乾淨的歸因實驗:單獨改編譯產物裡宣告的屏障資源數、不動任何可執行指令,就能復現同樣的變慢。這類實驗值得記一筆——它把「效能回退」從觀感變成了可指認的單一變數,也讓修復的驗收標準變得明確:不是籠統地跑分變快,而是資源宣告回到 4、佔用回到 16 塊/SM。
第二條是供給側的減法與口徑統一。#3736 把 tilelang/language/ast/、tilelang/language/parser/ 與 tilelang/language/tir/entry.py 一併刪掉,這些是早期本地 TVMScript 解析器分叉留下的模組,約 3380 行 Python,倉內已無呼叫方:語言層現在透過 tilelang.language.tir.common 使用上游的 tvm.tirx.script.parser 介面,AST 與語句覆寫註冊在 tilelang.language.overrides,T.prim_func 與 T.macro 由 tilelang.language.eager 匯出。實際改動是新增 1 行、刪除 3384 行、10 個檔案,同時移除了 docs/conf.py 裡針對 *language/ast* 的忽略規則;公開的語言層匯出由方言測試與前埠測試兜住。#3546 修的是自動調優的裝置口徑:set_autotune_inputs 安裝的供給閉包從一個呼叫點從不傳入的引數上解析基準裝置,實際傳過去的是核心引數列表,於是 CUDA 或 CPU 的兜底分支在沒有 CUDA 的機器上把每個捕獲張量搬到 CPU,每一輪調優都因啟動器裝置不匹配失敗;現在閉包改用 get_current_device(),與預設張量供給同一策略,CUDA 行為不變、昇騰與 MPS 改為在當前裝置上供給、其餘後端仍退到 CPU。驗收在雙卡 Ascend950 上:三筆昇騰自動調優迴歸修前全敗、修後全過,主倉 testing/python/autotune/ 12 項透過、28 項跳過,另有 3 項因該機沒有 CUDA 而失敗,作者明確留給 CUDA CI 驗證。
第三條在昇騰側:GM 到 L1 的複製從「一個描述符一個矩陣」擴到整批。#3689 允許把 [B, M, N] 區域寫成一個 T.copy,lowering 生成批次 ND2NZ 描述符,保留源側的批與行 stride,目的側的批間距按物理佈局推導,轉置批次走 DN2NZ;超出 GM 邊界的源批直接跳過並保持目的儲存不動,M/N 的尾部補零隻作用於合法批。既有二維與單矩陣行為不變,實現只針對 AscendC 後端、且要求 L1 目的端是規範矩陣佈局,PTO 不在本次範圍內。驗證在 Ascend950DT 與 CANN 9.2 上:50 個 AscendC 示例測試與 1733 個昇騰單測透過,其中 28 個裝置用例覆蓋批與 M/N 尾塊、整批越界、非連續源 stride、目的儲存更大與矩形轉置;CUDA 與 PTO 用例被排除,CI 映象用的 CANN 9.3 未在本地測到,這一點作者寫在了描述裡。改動 579 行、10 個檔案。
第四條是 TileFoundry 的一條暗線落地。上一期報告裡它還掛在「在審」,這一視窗合入:#285 讓 sync 分析規劃出的提交、等待與柵欄真正由 CUDA lowering 發射,並逐條反向校驗實際發射與計劃是否一致,不一致拋 LoweringError。語料上的結果是預設流程 14 個程式加上同步、拒絕 31 個,這 14 個的 TIR 與主線差別只有新增語句,另有 30 個與主線完全相同;單張 sm_89 上跑了 7 個能構建且帶同步的程式,6 個與參考完全一致、synccheck 與 memcheck 全乾淨,而同樣這 7 個按主線的方式 lower 有 4 個結果錯誤。作者在動機裡給的直接症狀是 sm_89 的 E4M3 GEMM 最大偏 659。這一筆 788 行新增、214 行刪除、25 個檔案,把「分析能證明」推進到「發射必須與證明一致」。
一、核心專案進展
1.1 主線:ThreadSync 允許跨儲存域複用部分屏障,B300 上屏障佔用從 12 回到 16 塊/SM(#3761)
合於 11-04 00:50,新增 401 行、刪除 72 行,4 個檔案。背景是上一個視窗的 #3643:ThreadSync("shared") 與 ThreadSync("shared.dyn") 分開呼叫時,可能給互不相交的執行緒組分配同一個命名屏障 ID,這是一個真實的正確性問題,修法是保留更早呼叫已經用掉的 ID。副作用是同一組執行緒跨作用域也不能複用,回到 TileKernels 的 mhc_pre_big_fuse 核心上就是每塊 96 執行緒的一次多餘分配:執行緒 32 到 95 的兩段式 cp.async 流水線在第一遍拿到 ID 3、第二遍拿到 ID 4,編譯產物宣告的屏障資源從 4 變 5,B300 隱藏維 1280 的屏障佔用上限從 16 塊/SM 掉到 12 塊/SM,而暫存器、共享記憶體與同步指令數完全沒變。本筆讓 ThreadSync 接受有序的作用域列表,CUDA 流水線改成 ThreadSync(["shared", "shared.dyn"]):同步仍按原順序規劃插入,部分的屏障 ID 在兩個作用域都處理完後統一分配;跨作用域共用一個 ID 之前用新分析器證明參與謂詞在候選執行緒範圍內恆成立,分支與 Bind 上下文由 ConstrVisitor 提供,執行緒相關謂詞作為證明義務而不是假設,已證密的密集組按包圍盒體積計算、跳過部分執行緒數列舉,未證明的組把作用域留在鍵裡並讓參與者計數入鍵,顯式編號的屏障(含更早呼叫分配的 ID)繼續保留,不相交執行緒組仍然隔離——#3643 的修復沒有被回退。字串形式的舊呼叫保持相容,CuTeDSL 也走同一條 CUDA pass 流水線。
來源:https://github.com/tile-ai/tilelang/pull/3761
1.2 主線:移除遺留的 TVMScript AST 與解析器分支,一次刪掉 3384 行(#3736)
合於 11-03 13:05,新增 1 行、刪除 3384 行,10 個檔案。被刪的是 tilelang/language/ast/、tilelang/language/parser/ 與 tilelang/language/tir/entry.py——早期本地 TVMScript 解析器分叉的殘留,倉內沒有任何呼叫方:語言層透過 tilelang.language.tir.common 使用上游的 tvm.tirx.script.parser 介面,TileLang 自己的 AST 與語句覆寫註冊在 tilelang.language.overrides,T.prim_func 與 T.macro 由 tilelang.language.eager 匯出,tilelang/language/tir/ 下只有 common、exports、ir、op 四個模組被消費。保留這個孤立的解析器分叉的風險寫得很具體:一旦有人直接 import tilelang.language.parser.operation,它會在 tirx.PrimExpr 與 tirx.IterVar 上註冊互相沖突的運算元。本筆同時刪掉 tilelang/language/tir/__init__.py 裡對 prim_func 的無用再匯出,以及 docs/conf.py 中針對 *language/ast* 的忽略規則(*parser* 保留,因為 tilelang/language/overrides/parser.py 還在用)。驗收口徑是公開匯出不變,由 test_tilelang_language_dialect.py 與語言前埠測試覆蓋。
來源:https://github.com/tile-ai/tilelang/pull/3736
1.3 主線:自動調優的輸入供給改用當前裝置,昇騰與 MPS 上的調優恢復可用(#3546)
合於 11-03 13:19,新增 153 行、刪除 12 行,3 個檔案,修 #3545。set_autotune_inputs 安裝的凍結供給閉包從「呼叫點從不傳入」的某個引數上解析基準裝置——按供給程式的契約,實際傳進來的是核心引數列表——於是在沒有 CUDA 的機器上,CUDA 或 CPU 的兜底分支把每個捕獲張量搬到 CPU,每一輪調優都因啟動器裝置不匹配而失敗。現在閉包改用 get_current_device(),與預設張量供給同一策略:CUDA 上行為不變,多卡基準 worker 仍按卡複製;昇騰與 MPS 上捕獲張量在各自當前裝置供給,調優因此可用;其餘後端仍退到 CPU。新增的昇騰用例放在 testing/ascend/autotune/,覆蓋裝飾器路徑、AutoTuner 的基準供給裝置、以及跨裝置捕獲的輸入三種情形,雙卡 Ascend950 上修前全敗、修後三項全過;主倉 testing/python/autotune/ 12 項透過、28 項跳過,另有 3 項(eager 模式、標量輸入、帶輸入)因該機無 CUDA 失敗,作者註明留待 CI 的 CUDA runner 驗證。
來源:https://github.com/tile-ai/tilelang/pull/3546
1.4 在審:條件中的標量 cast 不再壓縮資料向量寬度,承接昨日 #3734(#3759)
11-03 14:09 更新,新增 148 行、刪除 29 行,2 個檔案。這是對昨日合入的 #3734 的收尾:#3734 在 DecoupleTypeCast 之後正確收緊了資料 cast 的寬度限制,但也讓「標量條件裡的 cast」開始限制資料向量寬度——在 b5cb02633 上,一個受守衛的 BF16 複製可能因為條件裡有一個 int64 cast,從 16 個 lane 被壓到 4 個,哪怕該條件在 16 個 lane 上完全一致。修法是把規劃器已有的條件一致性分析複用到 cast 寬度豁免上,把 if 語句、tirx.if_then_else 與標量謂詞的 Select 條件從它們的子句裡分出來單獨訪問,條件裡的 cast 不再施加資料型別寬度限制,子句裡的 cast 仍然施加,並在帶逐 lane 向量謂詞的目標上保留條件 cast 的約束;#3579 的記憶體守衛行為與 #3734 修好的守衛 FP32 轉 BF16 限制都保持不變。迴歸覆蓋 8 lane 與 16 lane 的一致守衛、三種條件形式、子句內的 FP32 cast、向量謂詞與舍入位分組,九個 int64 守衛用例修前失敗,最終合法化與 cast 舍入兩組套件共 74 項透過。
來源:https://github.com/tile-ai/tilelang/pull/3759
1.5 在審:兩處使用者可見錯誤改為明確報錯(#3765 / #3767)
兩筆都在 11-03 17:10 前後更新,同一個報告人提交。#3765 對應 #3653:由單個執行緒索引邊界守衛的部分屏障,當該邊界是 warp 大小的整數倍時,整 warp 的執行緒對條件判斷一致;但通用算術證明漏掉了外層分支使用按位 & 時的 thread_binding < 448,ThreadSync 直接中止。本筆線上程維是整 warp 數時顯式接受這種情形(同時覆蓋邊界等於或超出該維的情況),其餘條件仍走既有證明路徑。 #3767 對應 #3005:fragment 跨執行緒分佈,同一 T.Parallel 體內用兩個結構不同的索引讀寫同一 fragment(例如 buf[i] 與 buf[(i + 1) % N] 這種跨 lane 模板)無法在不做跨執行緒交換的前提下 lower,此前 ParallelOpNode::RecordBufferAccess 撞上 StructuralEqual 的 ICHECK,暴露成編譯器內部斷言;現在改成 ValueError,把兩個索引表示式都寫進報錯並說明 fragment 不能這樣訪問。兩筆都屬於「把內部斷言換成使用者能讀懂的診斷」,對使用者體感的影響比改動規模大。
來源:https://github.com/tile-ai/tilelang/pull/3765、https://github.com/tile-ai/tilelang/pull/3767
1.6 在審:帶 padding 的共享 tile 無法整體複製時回退,2-CTA 的 TMA 位元組按對等 leader 記賬(#3738 / #3766)
兩筆都在 11-03 13:47 到 13:31 之間更新。#3738 修 #3704:從比 padding 後共享 tile 更小的全域性區域做 T.copy(例如 K=100 拷進 BK=112)會被 ClassifyWarpSpecializedCopy 與 SelectCopyInstForLowering 判成 TMA 整塊複製,然後在 lowering 硬失敗,報 s_extent == g_extent (112 vs. 100)。原因是選擇邏輯只檢查了最後一維對齊、dtype 配對與全域性 stride,沒有檢查 padding 後的範圍能否構成一次整塊複製;本筆補上這一判據,不滿足時回退到普通複製路徑。#3766 修的是 cluster 場景下 TMA 位元組計數:cta_group::2 的 TMA load 會清掉屏障地址的對等位,一對 CTA 的位元組都落在偶數 CTA 的屏障上,而帶 cluster 屏障的 tma_copy 原先只對 rank 0 期待 cluster_size × tile_bytes——在 (2,1,1) 與 (1,2,1) 上碰巧對,在 (4,1,1) 與 (2,2,1) 上 rank 0 只收到 2 份卻期待 4 份,另有位元組落到 rank 2 的屏障上。兩筆都不是新功能,屬於「邊界條件一破就報內部錯誤」的那類缺陷。
來源:https://github.com/tile-ai/tilelang/pull/3738、https://github.com/tile-ai/tilelang/pull/3766
1.7 在審:B300 的 sm_103a 適配兩條線,tcgen05 的 int8 選擇與 CuTeDSL 的 TMA 架構門(#3770 / #3772)
B300 報出的目標架構是 sm_103a,兩條在審 PR 分別處理它帶來的兩個卡點。#3770 對應 #3769:tcgen05.mma.kind::i8 只存在於架構專屬的 sm_100a 與 sm_110a,既不在 sm_103a(B300)上,也不屬於家族可移植的 sm_10xf,而 AllowTcgen5Mma 對 TargetIsSm100(arch 100 到 110)一律接受 int8/uint8 運算元,核心順利透過 lowering 後在 ptxas 報 Feature '.kind::i8' not supported on .target 'sm_103a';本筆讓 TileLang 在這些目標上不再選擇 kind::i8。#3772 處理 CuTeDSL 的 TMA 架構門:B300 的 sm_103a 在進入 DSL 之前就被拒,報錯列出 sm_90/90a/100a/110/120/120a/121/121a 而不含 103a;在主線 fa69f7fa 加 nvidia-cutlass-dsl 4.7.0 與 CUDA 13.3 的組合下,B300 上有六個 CuTeDSL 測試失敗,包括 tma_load 的多個 JIT 用例。同批還有 #3768 修 SM100 flash attention 反向示例:T.tcgen05_gemm 現在把 mbar 定為必填的僅關鍵字引數且不再隱式等待,而示例仍按同步 GEMM 的寫法呼叫,兩個指令碼在編譯前就報 missing 1 required keyword-only argument;作者同時指出 examples/flash_attention_sm100/ 目錄下沒有 test_*.py,這是問題一直沒被發現的原因。
來源:https://github.com/tile-ai/tilelang/pull/3770、https://github.com/tile-ai/tilelang/pull/3772、https://github.com/tile-ai/tilelang/pull/3768
二、多後端適配(昇騰 / 海光 / 摩爾執行緒 / 沐曦)
2.1 昇騰:3D GM 到 L1 的批次 ND2NZ 複製合入,一次描述符覆蓋整批(#3689)
合於 11-03 17:07,新增 579 行、刪除 39 行,10 個檔案。此前昇騰的 GM 到 L1 複製一個 DMA 描述符只處理一個矩陣,拷一批要寫 for b in T.serial(B) 逐批呼叫;本筆允許把 [B, M, N] 區域寫成一次 T.copy,lowering 生成批次 ND2NZ 描述符,保留源的批與行 stride,目的側的批間距按物理佈局推導,轉置批次走 DN2NZ,並複用既有的矩陣軸邊界處理:超出 GM 邊界的源批被跳過、目的儲存保持不動,M/N 的尾部補零隻作用於合法批。既有二維與單矩陣行為不變;實現只針對 AscendC 後端,且要求 L1 目的端是規範矩陣佈局,PTO 支援不在本次改動內。驗證環境是 Ascend950DT 加 CANN 9.2:50 個 AscendC 示例測試與 1733 個昇騰單測透過,其中 28 個 BF16/FP32 裝置用例覆蓋批與 M/N 尾塊、整批越界、非連續源 stride、目的儲存大於源、矩形轉置複製。作者註明 CUDA 與 PTO 用例被排除,CI 映象使用的 CANN 9.3 未在本地驗證。這條與上一期合入的昇騰 950 FlashMLA 核心是同一方向的配套:注意力與 GEMM 的批次搬運都需要一次描述符覆蓋整批,逐批呼叫在生產形狀上會攤出可觀的發射開銷。
來源:https://github.com/tile-ai/tilelang-ascend/pull/3689
2.2 昇騰在審:TileLang-Ascend 整倉遷入主線多後端架構,2887 個檔案(#1975)
11-03 16:34 更新,改動規模達到新增 607317 行、刪除 97595 行、2887 個檔案,是本期體量最大的一筆,也是整個組織在審 PR 裡最值得跟蹤的一條。目標是按提案 #1911 把 TileLang-Ascend 搬到主線的多後端架構上:原有昇騰元件移入主線對應目錄的 ascend/a2a3 下(tilelang/ 對 tilelang/ascend/a2a3,src/ 對 src/ascend/a2a3),原 Ascend C 與 PTO A2A3 後端在主線獨立註冊,與主線的 Ascend、PTO A5 正交,pass 與 codegen 互不干擾、只複用少量基礎設施;兩邊共用 target="ascend" 與 "pto",但 A2A3 後端附帶 arch="dav-2201" 標識(由 NPU 名稱自動識別)按此分流;A2A3 前端原語改用 import tilelang.ascend.a2a3.language as T,與 A5 的 import tilelang.ascend.language as T 並存,運算元寫法相對原倉不變;pass 與 codegen 直接用從 TileLang-Ascend 遷來的版本以保持原編譯流程,執行後端複用主線 cython 並補 workspace 索引等昇騰獨有適配,examples 與 testing 分別落到 examples/ascend_a2a3 與 testing/ascend_a2a3,CI 暫用原倉的 ci_cd.yml 並確認適配新路徑後仍能跑通。作者在描述裡把待辦寫得很實在:三個目錄下檔案與原檔案的差異尚待收斂、靜態檢查與格式化 CI 待恢復、文件與 Agent 相關待遷移,並註明「涉及檔案較多,Agent 批次操作難免有疏漏」,請求協助檢視。這種把遷移中的不確定直接攤在 PR 描述裡的做法,比等一個「已完成」的結論更利於評審。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1975
2.3 昇騰在審:copy_pipe_to_l1 的同步標註按執行引擎拆分,修 A2/A3 上 10/10 復現的錯結果(#1981)
11-03 12:36 更新,對應 #1780。根因寫得很清楚:A2/A3 的 PTO 跨核 V 到 C 通訊經 GM 中轉,三段搬運分屬三個引擎——AIV 的 MTE3 做 UB 到 GM 的傳送、AIC 的 MTE2 做 GM 到 L1 的接收(TPOP 內部 TLOAD)、AIC 的 MTE1 做 L1 到 L0A;而 copy_pipe_to_l1 在 operation_config.h 裡被標成 PIPE_MTE3,自動同步 pass 因此在它和 copy_l1_to_l0a 之間生成 set_flag(PIPE_MTE3, PIPE_MTE1, …) 加對應的 wait_flag——等的卻是 AIC 本核的 MTE3,而 AIC 的 MTE3 佇列為空、事件立刻置位,MTE1 讀 L1 於是可能在 MTE2 寫完之前放行,L0A 拿到舊資料、MMA 結果錯誤。錯誤標註隨 #1238 引入,同一張表裡 copy_pipe_to_ub 與 copy_l0c_to_pipe 都按執行側引擎標註,只有這一條沿用了傳送端視角。復現證據兩條:issue 作者把生成原始碼裡的 PIPE_MTE3 改成 PIPE_MTE2 後 10/10 全部恢復精確,本地在最新主幹上 10/10 次執行錯誤、單次最多 16384/16384 個元素不匹配。修法按 issue 的建議拆名字而不是直接替換:copy_pipe_to_l1 標註改為 PIPE_MTE2(A2/A3 的 GM 到 L1 載入),新增 copy_pipe_to_l1_bind 標註 PIPE_MTE1(A5 直連路徑上等待 L1 就緒並繫結地址),沿用 copy_pipe_to_ub 與 copy_pipe_to_ub_V 的既有先例,避免修好 A2/A3 卻讓 A5 出錯。同源問題也覆蓋了 WAR 方向:MTE1 仍在讀時,MTE2 的 TLOAD 可能覆寫同一 L1 緩衝。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1981
2.4 昇騰在審:首 Parallel 軸廣播走標量讀、Chunk O 前向遷移、MHA 示例(#1979 / #1866 / #1879)
三筆都在 11-03 上午集中更新,屬於昇騰運算元里程碑式的一批。#1979 對應 #1885:失敗族可以歸成一個模式——由第一根 T.Parallel 軸(或外層 serial 變數)索引的一維 UB 運算元出現在二元表示式左側,issue 的收縮模式 dk_i = sum_j v[j] * A[j, i] 正是這一形態,21 點二分矩陣裡 5 個 FAIL 全部屬於這一族,與歸約方向、廣播寫法無關;根因是 AscendLowerParallelToVector 把二維並行巢狀降成行迴圈時,入口 HandleSimpleCase 只對右側運算元做了標量訪問檢測,左側走向量路徑後撞上 UB 對齊陷阱;本筆讓首軸的廣播運算元也走標量讀。#1866 遷移 Chunk O 前向(基於 chunk 的門控線性注意力 O 前向,GDN)到昇騰 NPU,用開發者模式加四項 pass 配置(自動 C/V 合併、自動 C/V 同步、自動同步、記憶體規劃),單個 @tilelang.jit 核心加固定 20 核的 grid-stride 迴圈,做到片上融合、零 GM workspace、零宿主 ATen 運算元。#1879 新增 MHA 前向示例:BSHD 佈局、Expert 模式的三分派 MMA 流水,按 wrapper、共享層與三形態模組拆成 5 個 Python 檔案,wrapper 不含任何 T.* 程式碼、只做後設資料校驗與 JIT 分發;is_causal 採用右下角對齊掩碼,全掩碼行經 sumexp + 1e-30 保護輸出 0;cann-bench 20 個用例精度全過,對比 torch-npu 基線平均加速比 0.97 倍、綜合得分 73.13(本地 910B2 實測口徑)。0.97 倍意味著還沒有淨收益,但作為示例它把「昇騰上怎麼寫一個完整注意力運算元」的骨架給出來了。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1979、https://github.com/tile-ai/tilelang-ascend/pull/1866、https://github.com/tile-ai/tilelang-ascend/pull/1879
2.5 mlir-ascend:Codex 工作流與任務交接合入 harness 分支(#208)
合於 11-03 17:50,新增 2715 行、刪除 39 行,66 個檔案,合入目標分支是 harness 而不是 main(主線最新提交仍停在 10-31 的 ci: remove internal cache-service references),這一點在引用時不要寫成「合入主線」。內容是給運算元 harness 加第一套 Codex 工作流:專案初始化時用迭代提問、把已有運算元倉接入、以及從持久化上下文裡恢復開發、調優或分析;三套技能共享 .workspace/ 下可編輯的 Markdown 記錄,任務交接區分「當前候選」與「已驗證的最佳候選」,保留失敗與回退過的實驗,讓子代理拿到聚焦上下文而主代理維護共享狀態;本地安裝器負責註冊技能並把專案條目追加進 AGENTS.md(保留既有內容、忽略 .workspace/、把 Git 專案解析到 worktree 根)。README、依賴、複用說明與路線圖描述的都是 TileLang 先行、昇騰為範圍的設定;工作流參考了 GSD 並保留來源修訂與 MIT 宣告,NPU 執行介面卡與 npm 分發仍在待辦裡。這一筆與主倉、昇騰倉裡的 Agent 相關目錄是同一件事的三個面:把「智慧體寫運算元」的流程沉澱成倉庫內的可複用資產。
來源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
2.6 海光:Distributed 在審 PR 有更新,主分支無新提交(#19)
tilelang-hygon 在視窗內有推送記錄(11-03 11:06 北京時間),但主分支沒有新提交,最新一筆仍是 11-02 的 #32(把 gfx948 按 gfx946 相容處理,先讓既有指令可測)。推送來自在審的 #19(標題只有 Distributed,描述為空),11-03 11:06 更新,從時間與標題看是在推進分散式相關的適配,但 PR 描述沒有內容,本報告不對實現細節作推斷。該倉近期的常規節奏是每兩三天一筆 HCU(海光計算單元)側修正:10-30 的 #31 改 MLS 邊界與持久 WASP 同步、#30 補普通複製的 LDS swizzle、10-27 的 #29 做 scale 緩衝的 bank 劃分。
來源:https://github.com/tile-ai/tilelang-hygon/pull/19
2.7 沐曦、摩爾執行緒視窗內無推送(11-03)
tilelang-musa(摩爾執行緒)最後推送停在 09-30 09:50,tilelang-metax(沐曦)停在 10-26 07:08,兩倉本視窗都沒有提交,也沒有在審 PR 更新。本期的後端動態因此集中在昇騰與海光兩側:合入 8 筆裡有 1 筆屬於昇騰(#3689)、1 筆屬於 mlir-ascend(#208,合入 harness 分支),在審的 48 筆更新中昇騰倉佔 15 筆、主倉佔 23 筆。
來源:https://github.com/tile-ai/tilelang-musa、https://github.com/tile-ai/tilelang-metax
三、生態與採用方(TileOPs / TileFoundry / TileKernels / FlashQLA)
3.1 TileFoundry:同步計劃落成 CUDA lowering 並反向校驗,合入主線(#285)
合於 11-04 00:19,新增 788 行、刪除 214 行,25 個檔案。sync 分析規劃出共享記憶體訪問所需的提交、等待與柵欄,但 finalize 此前一條都不發射,經共享記憶體暫存的排程程式因此存在競爭,卡上表現為 sm_89 的 E4M3 GEMM 相對 torch 最大偏 659。本筆把分析步驟掛到它們所前導的 HIR 節點上,以 SyncMetadata 按標籤索引:節點自身語句的標籤為空,InsertSlice 加入的複製用 base 與 update,迴圈攜帶複製用 ("entry", 槽位) 與 ("backedge", 槽位);核心 lowering 在 lowering 某節點的語句之前呼叫新鉤子 before(owner, ctx, tag),由 CudaLowering 在那裡發射該節點的步驟;schedule finalize 負責跑 sync 分析。契約寫進了四份規範文件,其中分析規範 §1.2.5 描述 SyncMetadata、標籤、規則以及分析仍會拒絕的情形。風險面作者也標了:語料裡所有在主線能 finalize 的程式仍然能 finalize,只多出語句,例外是 hir/staged_dynamic_inner——它的 lane 在一個「行程數等於 lane」的迴圈裡讀寫同一塊共享 tile,沒有柵欄能把它們匯合;另外觸碰共享記憶體的保留呼叫按「未知位元組的一次訪問」處理,兩個這樣的呼叫在沒有執行緒網格的區域裡需要 TIR 無法放置的柵欄,分析會拒絕它們,即便位元組互不相交;CopyAsyncTensor、CopyReduceAsyncTensor 與 wgmma 還沒有發射器。
來源:https://github.com/tile-ai/TileFoundry/pull/285
3.2 TileFoundry 在審:wgmma 同步計劃、FSDP 與 HSDP 訓練語義進 HIR 求值器(#286 / #284)
#286 在 11-04 02:28 更新,新增 1907 行、刪除 310 行,52 個檔案,是 #285 的直接續筆:WGMMA 在發射之後仍可能繼續讀共享記憶體與暫存器,CUDA 同步計劃需要把後續訪問排在這些作用之後;本筆為每種運算元效應宣告完成語義,把 WGMMA 的柵欄、提交與等待掛到需要它們的 HIR 節點上,lowering 在指令的發射點消費這些步驟,若後續作用域放不下集合等待就在發射作用域裡完成該操作;同時修正暫存器 Reshard 別名,讓元素在持有副本的每個執行緒上留在同一個暫存器裡,計劃用這一儲存身份跟蹤暫存器依賴。作者明確寫了這條風險:矩陣指令發射器仍不可用,本筆不建立 WGMMA 的執行正確性或效能,TMA 完成語義仍單獨處理。#284 在 11-03 14:34 更新,處理訓練語義:此前的求值器把 reshard 當全域性張量,漏掉了不同參與者的貢獻;本筆求值參與者本地值、GPU 座標、呼叫與迴圈、gather/scatter 以及按軸的 Partial(sum) 解析,並讓 Reshard 的別名證明帶所有權語義,擴充套件 #260 與 #170 的結構化生命週期;配套加了一個受 modded-nanogpt 啟發的兩 block GQA 研究,含顯式前向反向、動量 SGD 與獨立的 PyTorch 參考,覆蓋 FSDP 與 2×2、4×4 兩種 HSDP 駐留策略,報告集合通訊組、載荷與儲存。驗收數字是重構後核心套件 775 項、後端套件 778 項(兩卡 sm90 上 13 項跳過),十組 HIR 配置在 CPU 上透過,兩種策略在四卡 sm90 上與真實 2×2 PyTorch HSDP 跑兩個更新步、本地梯度、引數與動量全部對齊。
來源:https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/pull/284
3.3 TileOPs:方差合併從第一個部分量起算,常量輸入的方差不再出 NaN(#2821)
合於 11-03 18:04,新增 43 行、刪除 4 行,3 個檔案。缺陷是方差合併從一個空的累加器起步:把一個有限的大均值合進空累加器時,平方差在乘以零權重之前就已經溢位,常量輸入的方差因此變成 NaN;受影響的是在保留軸前後各有一段歸約時對 var、std 與 var_mean 的呼叫。修法是讓合併從第一個真實的部分量初始化,後續部分量沿用既有的 Chan 遞推;行歸約、分發、輸出佈局與公開 API 都不變。驗收:三筆新增迴歸在原始核心上失敗、修復後透過,候選過 155 項迴歸測試與 24 項額外數值檢查,tests/ops/test_reduce.py 從 155 增到 158;三項自動調優測試被排除,原因是先前的基線嘗試超過了 60 秒測試上限。這類「空累加器 + 大偏移」的數值坑在歸約族裡並不罕見,值得記的是它只在保留軸前後同時歸約的形狀上才暴露,普通整軸歸約測不到。
來源:https://github.com/tile-ai/TileOPs/pull/2821
3.4 TileOPs 在審:GEMM 加可選 bias、FP8 分組 GEMM 支援連續佈局(#2823 / #2824)
#2823 給 GEMMFwdOp.forward 加了可選 bias[N]:bias 加進 float32 的累加結果、輸出只做一次舍入,與 torch.nn.functional.linear 的口徑一致;每種稠密結構都在自己的尾聲里加 bias——warp 專用(TMA 與行儲存)、coop2(含 176/192 寬與窄 tile)、coop2_swap、ping-pong、coop2s、simple、swap_ab、GEMV 各帶與 cp.async 核心,split-K 路徑在歸約裡只加一次。#2824 讓 MoEGroupedGEMMFP8FwdOp 除掩碼分塊外也接受連續專家行(緊湊打包或對齊打包,帶段末偏移 physical_psum 或逐行專家 id),語義對齊 deep_gemm.m_grouped_fp8_gemm_nt_contiguous;SM90 核心按段末遍歷每個專家的執行段,緊湊的逐行 id 在生產者與消費者拆分之前用二分轉成段末,對齊的逐行 id 按行 tile 逐塊走、從首行讀該 tile 的專家,緊湊專家最後一塊不足 block_m 行時逐行儲存,絕不寫到下一個專家的行上,SM89 拒絕連續佈局;掩碼呼叫仍構建與之前相同的程式,35 個掩碼構建的 TIR 與主線位元組一致。基準側新增三行 Qwen3-30B-A3B 的 EP8 形狀,用的是從真實路由測出的組大小(該模型 128 專家、top-8,取 48 層裡偏斜居中的第 44 層,選 EP8 中分到行數最多的那個 rank 的 16 個專家),prefill 行路由 4096 個 token(gate/up 與 down),decode 行取 128 條序列的末 token。
來源:https://github.com/tile-ai/TileOPs/pull/2823、https://github.com/tile-ai/TileOPs/pull/2824
3.5 TileOPs 在審:W4A16 與稠密 GEMM 的選擇器改成單一代價模型重擬合(#2825 / #2826)
兩筆同屬 #2133 的子任務 2,分別在 11-03 20:54 與 20:54 更新,一個管 W4A16、一個管稠密,且後者只改文件。#2825 把 _select_config 改成在 _legal_configs 上取 _config_cost 的 argmin,刪掉 tile 優先、最寬 K 的並列視窗與 joint_switch_margin 三條例外規則;_config_cost 新增 stage_fill 項——更深的環能藏住每個 K tile 的更多延遲,但填充本身更慢,沒有這一項時在 K=4096 上模型會把 256 寬四階段排到 512 寬兩階段之前,而後者實測最多快 11%;mma_big、iter_ws 與 launch_ms 被刪掉,因為資料不足以釘住它們。常數在基準的 SM 時鐘下按忙時重擬合,擬合形狀取清單之外的形狀、清單本身留作留出檢驗,_Calibration 的文件串寫明掃描、擬合與檢查三步,從中性常數重跑能復現出廠值。核心側還有一處實打實的省記憶體:split-K 的 CTA 只暫存自己那一片的 scale 與 zero 組,而不是整個 K,切片對齊時用向量化載入;_smem_bytes 與 _legal_configs 按每個 split 計共享記憶體,而主線上每個 split 的 CTA 都暫存了整個 K 的組。另一個細節是 _ConfigSpace.smem_bytes 放在比可選上限低 1 KB 的位置——正好壓在上限時,17 個被掃描的配置啟動失敗,因為分配器會把動態共享記憶體向上取整、而屏障是靜態的。#2826 則把稠密選擇器的常數如何重擬合寫進文件串(掃描、擬合、檢查三步,含軟退火與留出檢驗),不改程式碼。
來源:https://github.com/tile-ai/TileOPs/pull/2825、https://github.com/tile-ai/TileOPs/pull/2826
3.6 TileOPs 在審:短 varlen 按 KV 組打包、L2 flush 改為只讀(#2819 / #2822)
#2819 在 11-03 23:37 更新,是 #2496 的一部分,給 SM90 上的 GQAVarlenFwdOp 加了一個打包的 warp 專用程式:當一個請求用它的 KV 頭分組按位填滿一個 tile 更省 tile 時,鍵與值只按 KV 頭讀一次;打包核心的拒絕邏輯用一個代價模型衡量呼叫的批、打包後的 query 總數與組大小,只有當形狀保證的節省超過實測的 9% 單頭 tile 開銷時才接手;寬窄兩類 item 在同一次發射裡交替領取,不做宿主回讀、不額外加同步;當裝置上數出的未拆分 item 少於 SM 數時,單 tile 的打包請求把鍵掃描拆成多段、經 fp32 工作區合併。視窗側還有兩筆實測收益:視窗行省 1% 到 9%,交錯 RoPE 行省 3% 到 6%。同時修了主線上三個問題:帶 sink 且 score scale 大於約 2.4 時無可見鍵的行返回 NaN(單頭 warp 專用與 FP8 warp 專用兩個程式)、後備的 query 旋轉在請求首個鍵之前的行上先讀 RoPE 表、以及被拒的負控制用例讓呼叫方幀留在引用環裡導致一個測試的 CUDA graph 在另一個測試的捕獲期間被銷燬。測試節點從 173 增到 197。#2822 在 11-03 19:01 更新,把 bench_kernel 的 L2 flush 從「對 int8 緩衝做 zero_()」改成「讀一個私有 int32 緩衝,大小是報告 L2 的兩倍」,緩衝由固定種子的 CPU 生成器一次抽出;CUDA 上用 Triton 核心做這次讀(每元素一次載入、每 16 KB 一塊一次 atomic_max),不用 torch 自己的歸約是因為把它當 flush 會拖慢下一次 Marlin w4a16 呼叫 1.3 到 1.5 微秒。收益在 H200、1500 MHz、CUPTI 核心時間(3 次執行各 30 次呼叫的中位數)上:2048×11008 f32 的 Sum 從 27.81 降到 21.38 微秒,1024×11008 f32 的 Ge(torch-compile)從 31.31 降到 25.67,1024×4096 f32 從 6.30 降到 5.47;舊 flush 會留下約 120 MB 髒行,其回寫落進下一次計時呼叫,新 flush 自身只寫 71 到 445 個 L2 扇區(2 到 14 KB)。順帶修了一個主線 bug:L2 大小未知時日誌寫 256 MB 卻分配 512 MB。
來源:https://github.com/tile-ai/TileOPs/pull/2819、https://github.com/tile-ai/TileOPs/pull/2822
3.7 採用方:FlashQLA 與 TileKernels 無推送,但 #3761 修的是 TileKernels 的 fused 核心(11-04)
QwenLM/FlashQLA 最後推送停在 10-30 03:02,deepseek-ai/TileKernels 停在 10-14 05:49,兩倉在本視窗都沒有提交,採用方側沒有新程式碼。不過本期的主線頭條與採用方有直接關係:被 #3761 修掉的那筆效能回退,觸發點正是 TileKernels 的 mhc_pre_big_fuse 核心——同一個核心在 B300、隱藏維 1280 上丟掉了 25% 的屏障佔用預算。這說明採用方的核心已經成為編譯器側的迴歸探針:TileKernels 不推送程式碼的這幾周,它的核心仍透過主線的改動在被動地「被測試」。對跟蹤生態的人來說,這是一個值得留意的反向訊號——採用方靜默期裡,編譯器的改動仍然可能在它們的形狀上出問題。
來源:https://github.com/QwenLM/FlashQLA、https://github.com/deepseek-ai/TileKernels、https://github.com/tile-ai/tilelang/pull/3761
四、社群、教程與活動
4.1 單日 7 條新 issue:一條生產級 CUDA graph 問題,四條來自同一外部報告人(11-03)
視窗內主倉新增 7 條 issue(編號 #3762 到 #3775,另有一條昇騰倉的自動日測報告)。技術含量最高的是 #3775,報告人是帶生產負載的團隊:環境為從 pip 安裝的 tilelang 0.1.12、Python 3.12.3、兩張 H200 加 NVLink、驅動 580.178.04、CUDA 13 容器、torch 2.14.1+cu130,服務側是單程序 TP2(每個裝置一個工作執行緒、各自一個 CUDA 上下文)的 sglang,模型為 GLM 系 MoE 加 DSA 注意力,路徑上的 TileLang 核心是 mhc_pre_big_fuse_with_norm、mhc_post 與 DSA prefill 核心。把執行後端強制成 TILELANG_EXECUTION_BACKEND=nvrtc(預設 auto 會解析到 tvm_ffi)後,tilelang 的 JIT 核心在 CUDA graph 捕獲期間執行會讓捕獲損壞,而且損壞出現在 tilelang 發射之後的下游無關核心上,三次位元組一致的配置復現出三個不同的首個失敗點:一次是 cudaError700(無效的 TMA 描述符控制代碼)在 TMA 描述符初始化處丟擲、之後在無關發射點迴響,一次是 cublasGemmEx 報 CUBLAS_STATUS_EXECUTION_FAILED(純 BF16 進出),一次是宿主框架 JIT 出來的逐元素核心裡的非法記憶體訪問;三次都發生在 decode 的 CUDA graph 捕獲期間(批大小桶 8),都在服務啟動後約 100 秒,且 /proc/environ 與乾淨對照的差別只有那一個環境變數。這屬於最難定位的一類:現象隨機、失敗點漂移、誘因只有一個環境變數,報告人把三次執行的對照條件寫全了。
另四條來自同一位外部報告人,環境統一為主分支或 0.1.15 pip 包、WSL 下 Linux、Python 3.13.5、RTX 4060 Laptop(sm_89)、驅動 566.26、CUDA 12.6 執行時,且都附了可復現指令碼與完整輸出:#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 安裝」這一組合的集中報障面。剩下的兩條:#3769 是 B300 上 tcgen05 的 int8 GEMM 被 ptxas 拒絕(已有在審的 #3770 修),#3762 是昇騰 950 上 T.simd.vcvt 在 sat=False 時仍然飽和,詳見下條。
來源:https://github.com/tile-ai/tilelang/issues/3775、https://github.com/tile-ai/tilelang/issues/3774、https://github.com/tile-ai/tilelang/issues/3773
4.2 昇騰 950 的 vcvt 飽和語義與 mlir-ascend 日測(#3762 / #1980)
#3762 由昇騰側使用者提交,環境為 Ascend950PR(NpuArch 3510 / dav-3510)、CANN 9.2.0-beta.2、驅動 25.7.rc1、torch_npu 2.12.0、從原始碼裝的 TileLang(10-21 的 e84ce3de)。問題是 T.simd.vcvt(src, "uint8", sat=False) 對 int32 向量靜默忽略 sat=False,永遠飽和:負值變 0、超過 255 的變 255,既不報錯也不告警,而呼叫方以為自己在用迴繞語義。文件串對這種情形寫得很明確——sat=False 時越界值迴繞——期望值是 -1 變 255、-128 變 128、300 變 44,實測是 -1 變 0、-128 變 0、300 變 255。影響落在 int8 量化尾聲那類 f32 轉 int8 的路徑上:一旦出現負的中間值,語義就從迴繞變成截斷到零。#1980 是昇騰倉的自動日測,11-03 09:35 那輪 3030 項全部透過、透過率 100%,工作流執行編號 41553942797。這兩個數字放在一起看有意思:昇騰倉的日測是滿綠的,而日測覆蓋不到的語義差異(比如 sat 引數的實際行為)仍然要靠人工報告。
來源:https://github.com/tile-ai/tilelang/issues/3762、https://github.com/tile-ai/tilelang-ascend/issues/1980
4.3 文件站三次更新,VS Code 擴充套件改名 vsix(11-03)
文件站 tilelang.github.io 在視窗內有三次 Update docs 提交(11-03 13:11、13:28、11-04 00:59 北京時間),屬於隨主線合入自動構建的常規更新,內容對應本視窗合入的幾筆——#line 分級對映、ThreadSync 作用域列表與昇騰批次複製的文件相應位置。TileOPs.github.io 本視窗沒有推送,最新仍是 11-01 的「跟隨 TileOPs 主線:重疊與對齊保證、構建說明」。VS Code 擴充套件倉在 11-04 02:19 有一筆「更新文件並改名 vsix」,接在 11-02 的「把 CUDA 除錯對齊到 TileLang 行對映」之後:上一期主線合入的逐行錨定改變了生成程式碼的行號來源,擴充套件側的除錯配置需要跟著改,這一筆把產物名與文件一併調整。
來源:https://github.com/tile-ai/tilelang.github.io、https://github.com/tile-ai/tilelang-vscode-edition
4.4 新聞側:三通道無視窗內命中,外媒的昇騰訓練報道口徑不一(11-03)
Google News 的中英文組合查詢、Hacker News 與 arXiv 三個通道在本視窗都沒有 TileLang 相關的實質命中:候選清單 15 條裡 12 條來自 GitHub 倉庫發現,2 條是與主題無關的 Show HN(偵探遊戲與硬體設計提示詞工具),1 條是 Wccftech 關於 DeepSeek 與昇騰的報道。arXiv 上 all:tilelang 的最新命中仍是 2026-07-24 的 TileSight(以 tile 為中心的 GPU 效能分析模型),與 TileLang 的聯合查詢在視窗內無新論文。唯一值得記的行業側訊息有口徑衝突:Wccftech 稱 DeepSeek 將在昇騰 950DT 構成的 Atlas 950 SuperPoD 上從零訓練 V5,而同期 The Information 的報道是「最早 2026 年底開始在 950DT 上做預訓練」,Bloomberg 與金融時報的早期報道則稱 950DT 的 16 萬片訂單主要用於推理、此前在昇騰上訓練受挫後回退到 Nvidia。三種說法對「是否以及何時用於訓練」並不一致,且均未獲華為或 DeepSeek 官方確認。把它放在這裡只作背景:不論訓練是否落地,昇騰側的編譯器與運算元投入(本期 #3689、#1975、#1981)都是在為這條路徑鋪地基。
來源:https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods、https://arxiv.org/
五、趨勢觀察
第一,效能回退開始被當成一等公民,而且修法在方法論上比結論更值得記。#3643 為修正確性引入的 ID 保留規則,本身是必要且正確的;#3761 沒有把它回退,而是讓保留與複用在不同層級上並存——跨作用域共用一個 ID 之前先證明參與謂詞恆成立,證明不了就退回帶作用域的計數。配合作者那個「只改宣告的屏障資源數」的隔離實驗,這條修復鏈把「正確性」與「資源宣告」兩件事拆開了:後者是可以單獨復現、單獨驗收的變數。這種拆分能力,恰恰是編譯器專案從「能跑」走向「可維護」的標誌。
第二,昇騰的工程化在兩條腿上同時推進。一條是架構:#1975 用 2887 個檔案、60 萬行新增把整個 TileLang-Ascend 搬進主線多後端框架,靠 arch="dav-2201" 這樣的標識做後端分流,目標是讓 A2A3 與 A5 在同一個倉里正交共存;另一條是流水線:mlir-ascend 的 #208 把 Codex 的初始化、接入、恢復三套工作流和任務交接寫進 .workspace/,主倉與昇騰倉也各自有 Agent 目錄。把「遷移倉」與「用智慧體遷移倉」放在同一天看,能看出昇騰側的瓶頸已經從「缺運算元」轉向「缺一致的工程骨架與可複用的遷移流程」。
第三,運算元庫的自證文化在加深。TileOPs 本期在審的幾筆,幾乎都把「證據」寫在 PR 裡:真實路由測出的 Qwen3-30B-A3B EP8 組大小(取 48 層裡偏斜居中的第 44 層)、35 個掩碼構建的 TIR 與主線位元組一致、L2 flush 改動附 46 行實現的對照表並說明舊 flush 留下 120 MB 髒行、W4A16 的代價模型明確寫出「哪些項被刪掉是因為資料釘不住」。TileFoundry 走的是另一條路——把同步做成可證明的分析,再讓 lowering 反向校驗,分析拒絕的情形直接列出來(TMA 與 wgmma 的完成語義尚未宣告、跨執行緒邊界不等的迴圈)。兩條路合起來說明同一個取向:寧可承認「沒證明」,也不把「跑通了」當成「對了」。
第四,風險面在往安裝形態上集中。一天 7 條新 issue 裡,四條來自同一位使用 pip 包、跑 WSL、用消費級 sm_89 卡的外部報告人,全部附了可復現指令碼;#3775 則是生產環境裡「pip 包 + 多 CUDA 上下文 + nvrtc 後端 + CUDA graph 捕獲」的組合。這些不是編譯器的深層缺陷,卻決定了新使用者第一次嘗試時的體驗。對照昇騰倉日測 3030 項全綠而 sat=False 的語義偏差仍靠人工報告,可以看出測試覆蓋的形狀與真實使用形狀之間還有距離。
最後一點觀察:本期新聞側依舊空白,15 條候選裡 12 條是倉庫發現。TileLang 的對外聲量仍然繫結在「DeepSeek 與昇騰」這類事件上,而日常進展只發生在提交裡。對跟蹤這個專案的人來說,這意味著日報的主體只能是 PR 與 issue,而不是報道——這已經是連續多期的穩定形態。
附、素材與核查說明
信源核查表
| 信源 | 核查結果 |
|---|---|
| tile-ai 組織(29 倉庫) | 視窗內推送 9 個倉庫 |
| 主倉 tilelang | 主分支 5 筆合入;新開 issue 6 條;在審更新 23 筆;文件站 3 筆自動構建 |
| TileOPs | 主分支 1 筆合入;在審更新 7 筆 |
| TileFoundry | 主分支 1 筆合入;在審更新 2 筆 |
| tilelang-ascend | 預設分支 ascendc_pto 視窗內無新提交;在審更新 15 筆;日測 3030 項全過 |
| tilelang-mlir-ascend | 1 筆合入 harness 分支(非主線) |
| tilelang-hygon | 主分支無新提交;在審 #19 有更新 |
| tilelang-vscode-edition | 1 筆合入(文件與 vsix 改名) |
| tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / tilelang-puzzles | 視窗內無推送 |
| 釋出頁(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) | 視窗內無新版本,tilelang 最新仍為 09-30 的 v0.1.15 |
| FlashQLA / TileKernels | 視窗內無推送 |
| Google News(中英文六組查詢)/ Hacker News / arXiv | 無視窗內實質命中 |
完整信源清單
- [1] tilelang ThreadSync 跨儲存域複用部分屏障(#3761)— https://github.com/tile-ai/tilelang/pull/3761
- [2] tilelang 移除遺留 TVMScript AST 與解析器(#3736)— https://github.com/tile-ai/tilelang/pull/3736
- [3] tilelang 自動調優輸入供給裝置(#3546)— https://github.com/tile-ai/tilelang/pull/3546
- [4] tilelang 條件 cast 保留向量寬度(#3759)— https://github.com/tile-ai/tilelang/pull/3759
- [5] tilelang warp 對齊執行緒邊界的部分屏障(#3765)— https://github.com/tile-ai/tilelang/pull/3765
- [6] tilelang 多索引 fragment 訪問報錯(#3767)— https://github.com/tile-ai/tilelang/pull/3767
- [7] tilelang padded 共享範圍無法整塊複製時回退(#3738)— https://github.com/tile-ai/tilelang/pull/3738
- [8] tilelang 2-CTA TMA 位元組按對等 leader 記賬(#3766)— https://github.com/tile-ai/tilelang/pull/3766
- [9] tilelang tcgen05 kind::i8 不用於無此特性的目標(#3770)— https://github.com/tile-ai/tilelang/pull/3770
- [10] tilelang CuTeDSL 接受 sm_103a 與 sm_107a(#3772)— https://github.com/tile-ai/tilelang/pull/3772
- [11] tilelang SM100 flash attention 反向示例(#3768)— https://github.com/tile-ai/tilelang/pull/3768
- [12] tilelang CuTeDSL 示例缺口補齊(#3685)— https://github.com/tile-ai/tilelang/pull/3685
- [13] tilelang nvrtc 後端損壞 CUDA graph 捕獲(#3775)— https://github.com/tile-ai/tilelang/issues/3775
- [14] tilelang 共享記憶體置換丟列偏移(#3774)— https://github.com/tile-ai/tilelang/issues/3774
- [15] tilelang while 迴圈內共享記憶體 GEMM 不穩定(#3773)— https://github.com/tile-ai/tilelang/issues/3773
- [16] tilelang 巢狀 int16 向下取整 Let 衝突(#3763)— https://github.com/tile-ai/tilelang/issues/3763
- [17] tilelang B300 上 tcgen05 int8 被 ptxas 拒絕(#3769)— https://github.com/tile-ai/tilelang/issues/3769
- [18] tilelang 昇騰 950 上 vcvt 忽略 sat=False(#3762)— https://github.com/tile-ai/tilelang/issues/3762
- [19] tilelang-ascend 3D GM 到 L1 批次 ND2NZ 複製(#3689)— https://github.com/tile-ai/tilelang-ascend/pull/3689
- [20] tilelang-ascend 整倉遷入主線多後端架構(#1975)— https://github.com/tile-ai/tilelang-ascend/pull/1975
- [21] tilelang-ascend copy_pipe_to_l1 同步標註拆分(#1981)— https://github.com/tile-ai/tilelang-ascend/pull/1981
-
[22] tilelang-ascend 首 Parallel 軸廣播走標量讀(#1979)— https://github.com/tile-ai/tilelang-ascend/pull/1979
[23] tilelang-ascend Chunk O 前向遷移(#1866)— https://github.com/tile-ai/tilelang-ascend/pull/1866
- [24] tilelang-ascend MHA 前向示例(#1879)— https://github.com/tile-ai/tilelang-ascend/pull/1879
- [25] tilelang-ascend 每日測試 3030 項全過(#1980)— https://github.com/tile-ai/tilelang-ascend/issues/1980
- [26] tilelang-mlir-ascend Codex 工作流與任務交接(#208)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/208
- [27] tilelang-hygon 分散式在審 PR(#19)— https://github.com/tile-ai/tilelang-hygon/pull/19
- [28] TileFoundry CUDA lowering 發射同步計劃(#285)— https://github.com/tile-ai/TileFoundry/pull/285
- [29] TileFoundry wgmma 同步計劃(#286)— https://github.com/tile-ai/TileFoundry/pull/286
- [30] TileFoundry FSDP 與 HSDP 訓練語義(#284)— https://github.com/tile-ai/TileFoundry/pull/284
- [31] TileOPs 方差合併從首個部分量起算(#2821)— https://github.com/tile-ai/TileOPs/pull/2821
- [32] TileOPs GEMM 可選 bias(#2823)— https://github.com/tile-ai/TileOPs/pull/2823
- [33] TileOPs FP8 分組 GEMM 連續佈局(#2824)— https://github.com/tile-ai/TileOPs/pull/2824
- [34] TileOPs W4A16 選擇器代價模型(#2825)— https://github.com/tile-ai/TileOPs/pull/2825
- [35] TileOPs 稠密 GEMM 選擇器重擬合文件(#2826)— https://github.com/tile-ai/TileOPs/pull/2826
- [36] TileOPs 短 varlen 按 KV 組打包(#2819)— https://github.com/tile-ai/TileOPs/pull/2819
- [37] TileOPs L2 flush 改為只讀(#2822)— https://github.com/tile-ai/TileOPs/pull/2822
- [38] Wccftech 關於 DeepSeek 與昇騰 950DT 的報道 — https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods