調研視窗：2026-11-04 07:00 ~ 2026-11-05 07:00（北京時間，24 小時；上一期截點為 11-04 07:00）。

信源：GitHub tile-ai 組織逐倉核查（29 倉庫，視窗內推送 8 個）+ 合入與在審真值檢索（合入 12 筆、在審更新 42 筆）+ 新增 issue 4 條 + 主分支逐筆提交核查（tilelang 3、TileOPs 7、TileFoundry 1、tilelang-ascend 1、mlir-ascend 2）+ 非預設分支核查（tilelang-hygon 的 feat/hcu-scaled-cast 分支 4 筆）+ 釋出頁與 tag 核查（5 倉，視窗內無新版本）+ 採用方兩倉 + Google News 中英文四組查詢（走代理）+ Hacker News + arXiv。

本期索引

  • 今日重點：一筆 10-14 合入的昇騰 PTO 框架改動（#3512）進入集中清賬期，一天冒出五個互不依賴的補丁；TileOPs 單日合入 7 筆效能改動，創節後最密的一天；主倉三筆合入全部與「度量口徑」有關，L2 刷寫從寫零改成只讀
  • 一、核心專案進展
    • 1.1 主線：CuTeDSL 的 TMA 架構門接受 sm_103a 與 sm_107a，B300 與 gr100 上六個用例恢復（#3772）（11-04）
    • 1.2 主線：L2 刷寫由寫零改為只讀歸約，do_bench 去掉 fast_flush 引數（#3534）（11-04）
    • 1.3 主線：AscendC CI 的自動調優供給用例顯式同步兩臺裝置，非同步複製錯誤不再偽裝成 arange 失敗（#3784）（11-04）
    • 1.4 在審：ThreadSync 的 warp 一致性證明補上由 warp 索引拼出的組合謂詞（#3787）（11-04）
    • 1.5 在審：T.copy 目的區域越界改為編譯期報錯，全複製 fragment 的包含證明補執行緒歸屬校驗（#3786 / #3788）（11-04）
    • 1.6 在審：SM120 的 FP8 GEMM 換用 block-scaled MMA，RTX 5090 上 8192 立方從 463 到 727 TFLOPS（#3790）（11-04）
    • 1.7 在審：Profiler 與自動調優的口徑校驗成組提交，六筆同日出（#3797 / #3798 / #3799 / #3716 / #3717 / #3718）（11-04）
  • 二、多後端適配
    • 2.1 昇騰 PTO：一筆架構改動引出四類迴歸，五個互不依賴的補丁同日提交（#3792 / #3793 / #3794 / #3795 / #3796）（11-04）
    • 2.2 昇騰：Cube 與 Vector 的核心比可配置，不再固定 1:2（#3783）（11-04）
    • 2.3 昇騰：AutoSchedule 的週期下界收緊，Z3 求解時間降 10.04%（#3776）（11-04）
    • 2.4 昇騰：AI Core 異常轉儲相容 CANN 9.2.0，版本門檻從 9.3 下放（#1956）（11-04）
    • 2.5 海光：feat/hcu-scaled-cast 分支四筆，E8M0 縮放轉換與 64 bank 的 LDS 讀規劃（11-04）
    • 2.6 mlir-ascend：BishengIR 工具鏈介面卡遷入 contrib（#209）（11-04）
    • 2.7 沐曦與摩爾執行緒視窗內無推送（11-04）
  • 三、生態與採用方
    • 3.1 TileOPs：單日 7 筆合入，GQA varlen 打包、ArgMax/ArgMin 每執行緒八向量、預設快取策略（11-04）
    • 3.2 TileOPs 在審：一元逐元素單波次啟動、FP8 分組 GEMM 連續佈局、GEMM 可選 bias（#2829 / #2824 / #2823）（11-04）
    • 3.3 TileFoundry：FSDP 與 HSDP 的訓練語義進 HIR 求值器（#284）（11-04）
    • 3.4 TileFoundry 在審：wgmma 同步計劃（#286），另開一條命名統一的 issue（#287）（11-04）
    • 3.5 採用方與第三方：TileKernels 與 FlashQLA 無推送，三個第三方倉在 24 小時內活躍（11-04）
  • 四、社群、教程與活動
    • 4.1 新增 issue 4 條：warp 尺寸 64 的證明、PTO 有符號原子、CuTeDSL 的 NVFP4 量化器（11-04）
    • 4.2 新聞側：四組查詢、HN 與 arXiv 均無視窗內命中（11-04）
  • 五、趨勢觀察
  • 附、素材與核查說明

今日重點

這一天的主線不是「又合了什麼」，而是「昨天的地基開始被驗收」。10-14 合入的 #3512 把 T.Parallel 向量化到 PTO 的 VMI 流水線上，是昇騰後端一次結構性改動：整數型別從無符號語義的 i* 改成顯式帶符號的 si*，SIMD 的 lowering 改由新的 PTO 校驗、合法化、向量化三段流程接管，顯式 SIMD 並行迴圈的判定也重寫了一遍。三週之後，這套改動在真實運算元的形狀上被逐個踩出四類迴歸：整數原子操作因為 pto.si32 不被接受而編譯失敗（#3789 報告、#3792 修復），T.clear 在 T.SimdVF 裡走到只支援浮點的廣播路徑（#3795），標量 cast 的邊界事實在 VMI 記憶體合法化裡被過早繫結（#3793），顯式 SIMD 的並行迴圈被當成標量 lane 訪問而拒絕（#3796），另有一筆把整數原子適配到無符號載體（#3794）。五個補丁全部基於主線、彼此不依賴，作者在描述裡反覆寫同一句「This PR is based directly on main and does not depend on the other PTO regression fixes」——這是有意的：讓五筆各自可獨立合入，避免像過去那樣攢成一個巨型回滾包。驗證也不再是「跑通了」，而是帶具體用例數的硬體實測：47、18、16、37 項，四個獨立運算元（SwiGLU、Top-k、Cache、Engram 反向）作為代表用例。

第二條是度量口徑的自我修正，而且兩個倉在同一天做了同一件事。主倉 #3534 把 do_bench 的 L2 刷寫從「往緩衝裡寫零」改成「讀一遍緩衝」：寫零會在 L2 裡留下整塊髒行，下一輪計時期間這些行要寫回，等於把刷快取的代價攤進了被測核心。#3534 用 cache.view(-1, 4096).sum(dim=1) 分塊讀，並順手刪掉了 fast_flush 這個公開引數——那個引數本來只是選擇用 int32 還是 int8 建緩衝。昇騰側的 msprof 後端改得更徹底：原來用一個 SimtVF 迴圈把零寫進 UB 再 WTS_FV 儲存，現在改成按 NORMAL_FV 從全域性讀入 UB，型別也固定為 int32。TileOPs 的 #2822 是同一思路的孿生改動，還額外記錄刷寫方法、只在方法一致時才與歷史讀數比較。三筆改動放在一起看，是一個訊號：效能數字開始要求「可跨版本比較」，而不只是「這一次比上一次快」。

第三條是硬體前沿的補票。B300 上報的目標架構是 sm_103a，而 CuTeDSL 的 TMA 輔助函式用一張白名單把架構擋在門外，於是 B300 上六個 CuTeDSL 用例在還沒進 DSL 之前就拋 ValueError；#3772 把 sm_103a 與 gr100（cc 10.7）上報的 sm_107a 一起加進白名單，作者在 B300 節點上覆現了 8 透過 1 跳過、在 gr100 上 13 透過 1 跳過。同一視窗裡 #3790 走另一條路：SM120 上 FP8 加 fp32 累加的 GEMM 一直用 SM89 的 mma.sync 原子，在 RTX 5090 上只有 cuBLASLt 的 68%，換成帶 ue8m0 單位縮放的 block-scaled MMA 後，8192 立方從 463 升到 727 TFLOPS，且輸出與舊原子逐位一致。兩筆的性質不同——一筆是「讓新卡能跑」，一筆是「讓新卡跑滿」——但都指向同一件事：架構支援的邊界正在從 Hopper 往 Blackwell 與國產加速器兩頭同時擴。

第四條留給 TileOPs 的提交節奏。單日 7 筆合入裡，5 筆是效能、1 筆 CI、1 筆文件：GQA varlen 在 SM90 上按 KV 組打包請求，把短請求的 key/value 讀取次數降到每個 KV 頭一次；ArgMax/ArgMin 在能一波駐留的前提下讓每執行緒抱八個 16 位元組向量，(2048, 4096) 的 bf16 ArgMin 從 8.704 微秒降到 7.393 微秒，快過 FlagGems 的 8.543；一元逐元素與歸一化類核心把「只讀一次」的運算元改回預設快取策略，只在真正重複讀的地方保留 evict-last/evict-first。這些改動的共同點不是「更快」，而是「知道為什麼快」——每筆都帶對照表、都標出自己仍然落後於對比項的那一行（Softmax 的 (32768, 4096) 就從 121.792 退到 131.648 微秒）。

一、核心專案進展

1.1 主線：CuTeDSL 的 TMA 架構門接受 sm_103a 與 sm_107a，B300 與 gr100 上六個用例恢復（#3772）

合於 11-04 12:46（北京時間，下同），是本期最早的一筆合入。B300 報告的目標架構是 sm_103a，而 CuTeDSL 的 TMA 輔助函式維護一張架構白名單，遇到不在表裡的值直接在進 DSL 之前拋錯，錯誤資訊把允許的取值列成 sm_90、sm_90a、sm_100a、sm_110、sm_120、sm_120a、sm_121、sm_121a。在 fa69f7fa 的主線上、配 nvidia-cutlass-dsl 4.7.0 與 CUDA 13.3，B300 上有六個用例因此失敗：test_tilelang_jit_cutedsl.py 裡 tma_load 路徑三個（GEMM JIT、多流核心、動態形狀各一），以及另一個 tcgen05 的 CuTeDSL 測試檔案裡 tma_store 路徑的三個 tmem_copy_roundtrip 變體。DSL 本身是接受 sm_103a 的——它的 tcgen05 操作接受整個 sm_100f 家族——所以問題純粹出在這道門禁。加上這一項之後，兩份檔案在 B300（cc 10.3、CUDA 13.3、cutlass-dsl 4.8.1）上是 8 透過 1 跳過。同一筆還加了 sm_107a，因為 gr100 的 bring-up 板（cc 10.7）上報這個值、會撞同一道門；兩項都加上後，同樣的檔案在 gr100 上是 13 透過 1 跳過，其中含 #3685 引入的 gather/scatter 檔案。

值得記的是作者交代的驗證邊界：手上沒有 sm_103 硬體，改動是靠另一位報告者（0z5a，在 #3685 裡首次報告並 bisect）提供的 B300 機器驗證的，CI 不覆蓋 sm_103a。這類「程式碼合入、驗證藉機器」的情形在架構追趕期會越來越多，讀者需要知道哪些能力是有實測背書的、哪些只是加入了白名單。

來源：https://github.com/tile-ai/tilelang/pull/3772

1.2 主線：L2 刷寫由寫零改為只讀歸約，do_bench 去掉 fast_flush 引數（#3534）

合於 11-04 11:27，改動 7 個檔案、新增 36 行、刪除 49 行，其中三處是真正的實現，四處是跟著改的測試。原來是 cache.zero_()：用一個 cache_size MB 的緩衝寫零來擠掉 L2，問題在於寫零之後緩衝裡全是髒行，下一輪被測核心執行時這些行要寫回記憶體，寫回的頻寬被算進了被測核心的時間裡——緩衝區越大，這個偏差越明顯。現在改成 cache.view(-1, 4096).sum(dim=1, dtype=cache.dtype)：按 4096 列分塊讀，把行讀進 L2 併產生歸約輸出，髒行的問題不存在；分塊是為了避免整塊歸約帶來的全域性同步；保留 dtype 是為了不觸發一次全緩衝的 int64 轉換。這一改動同時落在 bench.py 的 estimate 迴圈、torch_bench.py 的 event、cupti、cudagraph 三條路徑，以及 msprof.py 的昇騰專用刷寫核心上——後者原來用 SimtVF 把零寫進 UB 再以 WTS_FV 存回全域性，屬於典型的寫式刷寫，現在改成以 NORMAL_FV 從全域性讀入 UB，元素型別也從 int8/int32 二選一固定成 int32。

公開介面的變化需要留意：do_bench 的 fast_flush 引數被刪除。它原本只是在 int32 與 int8 兩種緩衝元素型別之間切換（int32 建緩衝更快），現在緩衝固定為 int32，引數沒有意義。對直接呼叫 do_bench(..., fast_flush=False) 的指令碼，這是破壞性改動；倉庫內的測試已經同步修改，對外部使用者則要靠 release note 提醒。

來源：https://github.com/tile-ai/tilelang/pull/3534

1.3 主線：AscendC CI 的自動調優供給用例顯式同步兩臺裝置，非同步複製錯誤不再偽裝成 arange 失敗（#3784）

合於 11-04 17:01。現象是 AscendC 的 CI 作業裡，自動調優的 set-inputs 用例中 [other_npu] 這個引數化項在下一次呼叫 torch.arange() 時報 srclen=4, dstlen=4——一次 4 位元組的 memcpy 尺寸不符，而這正好是跨裝置複製探針的形狀。真正的原因是跨裝置複製失敗被非同步上報，錯誤浮出的位置比發生的位置晚一步，於是報錯點看起來在一個完全無關的 arange 上。修法很樸素：在探針前後各加一次裝置同步（當前裝置與 npu:other），讓非同步錯誤落回 try ... except RuntimeError 的作用域內。驗證是單機雙卡跑該用例檔案，四項全過，且當跨裝置複製真的非同步失敗時，[other_npu] 這一引數化項被跳過而不是讓整個用例失敗。

這筆本身很小，但它接的是上一期 #3546 的尾巴——那筆把自動調優的輸入供給改成按當前裝置解析，驗收時提到昇騰雙卡上的迴歸，這裡補的是同一場景在 CI 上的穩定性。同一功能連著兩個視窗各修一次，說明昇騰側的自動調優仍是「能跑但邊界不穩」的狀態。

來源：https://github.com/tile-ai/tilelang/pull/3784

1.4 在審：ThreadSync 的 warp 一致性證明補上由 warp 索引拼出的組合謂詞（#3787）

這筆接手的是上一期 #3761、#3765 留下的同一片區域：ThreadSync 要證明「屏障參與執行緒在整個 warp 內一致」，證明不了就拒絕 lowering。觸發形狀是 192 個執行緒下一個組合謂詞 tx < 160 and tx // 32 % 2 == 0——warp 0、2、4 整組參與，warp 1、3、5 整組不參與，沒有任何 warp 被劈開，本該是可證明的，但主線直接拋內部錯誤「參與執行緒不是 warp 一致」。作者做了一張比對的表，把缺陷範圍壓得很準：單項謂詞 tx // 32 % 2 == 0、tx % 64 < 32、tx < 160 在主線上都能 lowering，只有「邊界項與 warp 索引取模項用 and 拼起來」這一種形態會掛；而 tx < 168（真的劈開了第 5 個 warp）在兩邊都正確拒絕。根因在證明器的編碼：它把執行緒號重新命名成兩個假想執行緒，讓算術求解器證明 linear1 / warp_size == linear2 / warp_size 能推出兩側謂詞相等，而 FloorDiv 與 FloorMod 的編碼在「上界 + 取模」這種組合下建立不起兩個除法的關係。修法是把 threadIdx.x 改寫成 warp_index * warp_size + lane，兩個假想執行緒共享 warp 索引、各自一個自由 lane，正好就是「同一個 warp」這個關係；於是只透過 warp 索引讀執行緒號的謂詞在兩邊塌成同一個表示式，一致性可證，而真正依賴 lane 的謂詞仍然不等，#3786 裡列出的 tx < 168 依舊被拒——證明沒有放鬆，只是補上了漏掉的一類。

同一作者同日還提交了 #3786：T.copy 的目的區域超出緩衝時改為編譯期報錯，而不是生成越界訪問。兩筆合起來是同一取向——把「編譯器默默生成錯程式碼」變成「編譯器說不」。

來源：https://github.com/tile-ai/tilelang/pull/3787、https://github.com/tile-ai/tilelang/pull/3786

1.5 在審：全複製 fragment 的包含證明補執行緒歸屬校驗（#3788）

ProveFragmentContains 在判斷「大 fragment 是否覆蓋小 fragment」時有一條快速路徑：只要大 fragment 是完全複製的，直接返回真。這條路徑在檢查執行緒歸屬之前就返回了，於是出現這樣的誤判：小 fragment 擁有執行緒 0 到 127，大 fragment 完全複製但只擁有執行緒 0 到 95，96 到 127 根本沒有被覆蓋，函式卻說包含成立；開啟 check_forward_index 時，兩個完全複製但物理槽位不同的 fragment 也會被接受。後果是佈局推斷把不相容的對映當成相容，屬於「錯誤被靜默放過」的型別，報告者是在查 #2399 時順手發現的。修法是先看物理索引，再把小 fragment 的執行緒對映換算到大 fragment 的座標系裡（要考慮有界的執行緒範圍偏移），用算術分析器證明落在複製範圍與執行緒範圍內，證不出來就返回假。驗證在 A100 上對上游主線新構建，透過一個測試期編譯的橋接直接呼叫真實的 C++ 實現，覆蓋執行緒數 128 對 96 的拒絕、96 對 128 與 96 對 96 的接受、物理槽位、非連續歸屬、有界範圍、符號索引與部分複製。

來源：https://github.com/tile-ai/tilelang/pull/3788

1.6 在審：SM120 的 FP8 GEMM 換用 block-scaled MMA，RTX 5090 上 8192 立方從 463 到 727 TFLOPS（#3790）

這筆是本期效能側最有分量的一條。問題出在指令選擇：SM120 上 FP8 加 fp32 累加的 T.gemm 在 RTX 5090、8192 立方的形狀上只有約 465 TFLOPS，而 cuBLASLt 是 681；調 tile 與流水級數都沒用，因為瓶頸是指令本身——TileLang 發射的是 SM89 時代的 mma.sync ... f32.e4m3.e4m3.f32，這種原子在這張卡上只有半速。作者做了一組只測暫存器的探針：e4m3 到 f32 走 mma.sync 是 512 TFLOPS，而走 kind::mxf8f6f4.block_scale（單位縮放）是 1021，走 f16 累加的 mma.sync 是 1022——差距全在指令上。既然 ue8m0 縮放取 0x7f（即 2 的 0 次方）時 block-scaled 與普通 MMA 結果相同，就可以在 TL_SM120_BLOCK_SCALED_MMA_ENABLED 開啟且目標為 sm_120a/f、sm_121a/f 時，讓四個 fp32 累加的 FP8 派發路徑改走 block-scaled 原子，其餘目標、f16 累加與顯式非 a 字尾的 sm_120 保持原樣。逐位一致性在指令級上驗過：e4m3 與 e5m2 的四種組合，隨機暫存器內容含 NaN 與 Inf 編碼。

數字：e4m3 GEMM、bf16 輸出、128×128×128 tile、3 級流水、256 執行緒的條件下，4096 立方從 408 升到 673（cuBLASLt 530），8192 立方從 463 升到 727（681），4096×14336×4096 從 449 升到 696（498），256×8192×8192 從 348 升到 606（457）。唯一變慢的是 M=16 配 128 行 tile 的形狀，31.5 微秒退到 33.5 微秒，作者說明該形狀受頻寬限制、換 16 或 64 行 tile 會更快。新增用例直接檢查 PTX：sm_120a 必須拿到 block-scaled 原子、sm_89 必須拿到普通原子。

來源：https://github.com/tile-ai/tilelang/pull/3790

1.7 在審：Profiler 與自動調優的口徑校驗成組提交，六筆同日出（#3797 / #3798 / #3799 / #3716 / #3717 / #3718）

六筆都來自同一位作者（Afloat16），主題是把「驗證器與被測物件的口徑」對齊。assert_allclose 現在會拒絕核心在該給張量的位置返回 None——原來的實現只要有一側是 None 就跳過比較，於是一個返回空的核心也能「透過」（#3797）；assert_consistent 原來直接對返回值做 zip，標量張量會拋 TypeError，缺一個或多一個輸出會被靜默忽略，形狀不同則可能靠截行或廣播蒙過去，現在先歸一化成輸出列表再核對數量與形狀，空張量也納入（#3798）；CUPTI 基準測試裡的 suppress_stdout_stderr 原來對 Python 流呼叫 fileno()，筆記本與 StringIO 沒有檔案描述符就崩，現在分別重定向原生描述符 1/2 與 Python 流，並用 ExitStack 保證異常路徑也能還原（#3799）。另外三筆在自動調優側：目標字典作為快取鍵時用字串表示，插入順序不同的等價字典會產生不同鍵，改成按排序後的 JSON 序列化（#3716）；Tensor Core 檢測優先用顯式算力、再用顯式或當前目標架構，最後才查本地裝置，否則無 GPU 的交叉編譯會誤報不支援（#3718）；Profiler 的動態形狀替換補上 m + 1 這類複合維度，並按缺失變數給出 ValueError 而不是 TypeError（#3717）。

這一組的技術含量不在改動本身，而在它們暴露的形狀：驗證器長期是「透過即正確」的薄弱環節，而自動調優的快取鍵、目標解析這類基礎設施一旦有偏差，會以「結果不可復現」的形式出現，很難歸因。

來源：https://github.com/tile-ai/tilelang/pull/3797、https://github.com/tile-ai/tilelang/pull/3798、https://github.com/tile-ai/tilelang/pull/3799

二、多後端適配

2.1 昇騰 PTO：一筆架構改動引出四類迴歸，五個互不依賴的補丁同日提交（#3792 / #3793 / #3794 / #3795 / #3796）

先說清背景。10-14 合入的 #3512（[Ascend][PTO] Vectorize T.Parallel to PTO VMI）把並行迴圈的向量化從舊的 SIMD 通道遷到 PTO 的 VMI 流水線：整數型別從無符號的 i* 改成顯式帶符號的 si*，Lowering 走 Verify、Legalize、Vectorize 三段，顯式 SIMD 迴圈的判定也重新實現。這五個補丁各自對應它在真實運算元上踩出的一處：

  • #3792 整數原子編譯失敗：codegen_pto 為整數發射 pto.si32 / pto.si64，而 PTO 的原子介面只接受無符號的 i32 / i64，於是 T.atomic_add、T.atomic_max、T.atomic_min 在整數緩衝上直接編譯失敗（對應 issue #3789）。修法是在 tl.atomic_* 外面套一層相容包裝：指標元素型別與運算元重寫為同寬無符號，整數 min/max 傳正確的符號性，返回值再轉回原型別。用例從 15 項擴到 18 項，覆蓋有符號與無符號 32/64 位的加減取極值、負運算元、前值返回、超過 32 位的計數器。
  • #3794 有符號載體適配：與 #3792 同源但走另一條路——不改 tl.atomic_* 的語義，而是在生成的 PTODSL 呼叫邊界上重寫指標檢視與運算元，保留地址空間，浮點原子保持原路徑。代表用例是 SwiGLU 的 clamp 計數（bf16、hidden 2048、clamp 0.5、psum 佈局），18 項原子用例全過。
  • #3795 SIMD 填充丟失：T.SimdVF 裡的 int32 T.clear 走到了只支援浮點的廣播路徑，原因是新 PTO 流水線不再執行舊的 SIMD pass，而那個 pass 同時負責填充與並行迴圈。修法是抽出一個只做填充的入口，複用原有的 LowerFill，放在緩衝版本歸一化之後、PTO 並行校驗之前，仍發射原來的 vdup/vsts 序列，並保留連續區域、對齊、整向量以及多緩衝版本切片的檢查。16 項用例透過，代表用例是 Top-k（512 token、1 專家、top-1）。
  • #3793 標量 cast 的邊界事實被過早繫結：u = uint32(i - 1)、i 在 [0, 63] 這種與 VMI 地址無關的標量，會在記憶體合法化裡拋錯；急於 Bind 還會留下一個被低估的區間，等 IntSet 的 cast 分析再拋。修法是先檢查整數 cast 的可表示性再安裝事實，給證明引擎留安全的標量定義、給 IntSet 保守的數值區間，避免直接求 cast 端點。47 項用例透過，代表用例是 Cache（block_size=64、n=65536）。
  • #3796 顯式 SIMD 並行迴圈被當成標量 lane：Engram 反向的初始化用了一個 Parallel 巢狀，其中顯式 SIMD 儲存按整向量塊推進，新的 lane 佈局推斷把它當標量訪問並拒絕（M=2, N=32, L=128）。修法是讓顯式 SIMD 的分類在佈局推斷與並行單元發現之間共享，這類巢狀保留 #3512 之前的通用佈局路徑，自動逐元素單元仍走新的 PTO lane 路徑。37 項用例透過，代表用例是 Engram 反向（512 token、hc 4、hidden 2048）。作者同時寫明一條限制：FP32/128 lane 的顯式加自動混合用例只有生成原始碼級覆蓋，當前 PTOAS/Bisheng 裝置編譯器在該組合上超過 6 KiB 的 VF 棧預算，沒有裝置執行結果。

把這五筆並排看，能讀出昇騰側的工程節奏：#3512 那樣的結構性改動先落地、再按真實運算元逐個暴露邊界，每個補丁都帶硬體實測與用例計數，並且刻意保持彼此獨立。「先合框架、後清迴歸」這套打法在上一輪架構遷移（#1975）裡已經用過一次，現在成了常規流程——代價是主幹在這類視窗裡處於「已知有幾處不工作」的狀態，外部使用者踩到的機率不低。

來源：https://github.com/tile-ai/tilelang/pull/3792、https://github.com/tile-ai/tilelang/pull/3793、https://github.com/tile-ai/tilelang/pull/3794、https://github.com/tile-ai/tilelang/pull/3795、https://github.com/tile-ai/tilelang/pull/3796

2.2 昇騰：Cube 與 Vector 的核心比可配置，不再固定 1:2（#3783）

昇騰的 AI Core 分 Cube 與 Vector 兩類，AscendC 支援 mix(1,2)、mix(1,1)、mix(1,0)、mix(0,1) 等分組，而 TileLang 此前一律按 1:2 啟動，等於把一個硬體自由度的上限鎖死，也讓未來核心比不同的晶片難以適配。這筆改動把核心比變成一個顯式或按裝置推導的引數：T.MixedKernel(..., sids=1) 或 sids=2 按請求的分組大小走並做能力校驗；省略或傳 None 時用當前裝置的 vector_core_num / cube_core_num，取不到裝置資訊就退回 2；T.Kernel 保持原介面與核心型別判定，識別為混合核心時用準備好的向量核數，單核路徑不變，顯式 T.Vector(1/2) 的分組語義保留。選定的組數會一路帶到 sid 域、任務分配、同步與程式碼生成，避免重複查詢裝置與重複歸一化；衝突的宣告與超過裝置能力的請求直接拒絕。

來源：https://github.com/tile-ai/tilelang/pull/3783

2.3 昇騰：AutoSchedule 的週期下界收緊，Z3 求解時間降 10.04%（#3776）

AutoSchedule 用二分搜尋找流水線的啟動間隔（II），原來的下界取所有任務 II 的最大值，偏松，二分要多走幾輪。這筆在資源模型下補了三條必要條件：單任務界（保留 P >= max_i II_i）、資源團界（對兩兩資源衝突的團按殘差排序求和，得到 P >= sum_i II_i）、自依賴界（S_u - S_u >= wait - d*P 推出 P >= ceil(wait/d)，版本相關的距離按流水級數截斷）。用三者最大值初始化二分，其餘邏輯不動。測試環境是昇騰 950PR、CANN 9.3.0 周更版，關快取、一暖三測交替：examples/ascend 全部編譯透過且排程結果不變，Z3 求解階段的累計耗時下降 10.04%。這類「不改結果、只改搜尋起點」的改動風險最低、收益卻是全域性的——編譯期的求解時間是每個使用者都要付的。

來源：https://github.com/tile-ai/tilelang/pull/3776

2.4 昇騰：AI Core 異常轉儲相容 CANN 9.2.0，版本門檻從 9.3 下放（#1956）

這筆合在 tilelang-ascend 倉的 ascendc_pto 分支（11-04 10:18），是本期昇騰倉唯一的合入。AI Core 異常轉儲功能原先宣告依賴 CANN 9.3.0 以上，作者在 9.2.0 上做了可行性驗證並修了三處：acldumpTensorInfo 的必填欄位補齊（addrType=ACL_DUMP_ADDR_RAW、placement、shapeNum、originShapeNum），因為 9.2.0 會強制校驗 addrType，而原實現靠 calloc 置零得到的 ACL_DUMP_ADDR_PTR 被拒絕；跳過空地址與零大小張量，9.2.0 不接受 tensorAddr 為 NULL；註冊介面從已廢棄的 aclrtSetExceptionInfoCallback 換成 aclrtExceptionInfoCallbackRegister。工具側 parse_exception_dump() 相容 9.2.0 輸出的 .npy 與舊版的 .bin 兩種格式，中英文教程的版本要求一併改為 CANN 9.2.0 以上。對還在 9.2 上的使用者，這是從「功能不可用」到「可用」的變化。

來源：https://github.com/tile-ai/tilelang-ascend/pull/1956

2.5 海光：feat/hcu-scaled-cast 分支四筆，E8M0 縮放轉換與 64 bank 的 LDS 讀規劃（11-04）

海光倉的主分支停在 11-02，本視窗的動靜全在功能分支上，四筆提交集中在八小時內（北京時間 16:26 到 20:16），作者來自海光（huangteng@hygon.cn 與 zhouweiwang@hygon.cn）：

  • feat(hcu): add scaled_cast and unify GEMM LDS read planning：新增 E8M0 縮放的 FP8/BF8/FP4 轉換，標量與打包兩種 lowering，新增 src/tl_templates/hcu/scaled_cast.h（223 行）與 lower_scaled_cast.cc；同時把 GEMM 消費端後設資料在矩陣讀選擇與 LDS 策略之間共享，支援單級 LDS 切片，主變換檔案改動 650 行。
  • fix(hcu): align GEMM LDS strategies with 64-bank read phases：按 bank-word 幾何推導 plain-copy 的 swizzle，為矩陣與兩種原生 pack 尺寸共用的窄 b16 tile 提供顯式逆變換；混合或受限相位的非同步消費端優先用相容的公共儲存；選擇原生 wrap 時把完整的 b128 讀相位計入。12 個檔案、新增 1536 行、刪除 112 行，新增 5 個測試檔案。
  • fix(hcu): prefer mixed LDS wraps proven at most two-way：給混合 LDS wrap 加「至多兩路」的證明，只在證明成立時優先使用。
  • fix(hcu): use physical read width for mixed wrap checks：混合 wrap 檢查改用物理讀寬度，去掉原來按 kPack 推斷 8 位元組每 lane 的寫法。

主題是共享記憶體的 bank 衝突與相位對齊——這是海光 DCU 上 GEMM 效能的關鍵路徑，也說明該後端已經從「能編譯」進入「按微架構調優」的階段。四筆都在功能分支上，尚未開 PR，時機上可能與 TileLang 主線做一次同步後再提交。

來源：https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-scaled-cast

2.6 mlir-ascend：BishengIR 工具鏈介面卡遷入 contrib（#209）

合入主分支（11-04 16:35），把 BishengIR 工具鏈的適配程式碼從原有位置搬到 contrib 目錄。這是純結構整理：把「隨上游工具鏈演進而需要頻繁改動的膠水層」與核心 lowering 分開，後續升級 BishengIR 時改動面收斂在 contrib 內。同一分支上還有一筆 refactor(contrib): move BishengIR toolchain adapters into contrib 的主體提交，合併提交是第二筆。這類整理單看沒有功能價值，但它是把「適配層」當成獨立生命週期管理的開始——昇騰工具鏈版本迭代快，膠水程式碼的可替換性比優雅更重要。

來源：https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

2.7 沐曦與摩爾執行緒視窗內無推送（11-04）

tile-ai/tilelang-metax 最後推送停在 10-26 15:08，tile-ai/tilelang-musa 停在 09-30 17:50，兩倉在本視窗都沒有提交，也沒有在審 PR 更新。國產四家後端裡，昇騰與海光是本視窗的兩條活躍線，沐曦與摩爾執行緒處於靜默期。上一期報告裡沐曦同樣無推送，已連續兩個視窗；摩爾執行緒則是自 09-30 以來沒有動靜，值得在後續視窗留意是否轉入內部開發。

來源：https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa

三、生態與採用方

3.1 TileOPs：單日 7 筆合入，GQA varlen 打包、ArgMax/ArgMin 每執行緒八向量、預設快取策略（11-04）

7 筆全部在視窗內合入，節奏上是本季度最密的一天，逐筆看：

#2819 給 SM90 的 GQAVarlenFwdOp 加了打包的 warp 特化程式：當一個請求的 KV 頭組能在一個 tile 內按位置鋪滿時，讓請求按 KV 組填充 tile，於是每個 KV 頭只讀一次 key/value；是否走打包由代價模型判斷，只有「形態保證能省下的量」超過實測的 9% 單頭開銷才用；寬窄請求在同一次啟動裡交替領取，不做主機回讀、不加同步；當裝置上統計出的不可拆項少於 SM 數量時，單 tile 的打包請求把 key 掃描拆成多段、用 fp32 工作區合併。另修了三個主線缺陷：sink 加分數縮放超過約 2.4 時無可見 key 的行返回 NaN、回退路徑在請求首個 key 之前讀到 RoPE 表、以及一個負向對照把呼叫方棧幀留在引用環裡導致跨測試的 CUDA graph 被提前銷燬。測試節點 173 增到 197。

#2828 給按行分塊的 ArgMax/ArgMin 在「一波駐留」的前提下讓每執行緒抱八個 16 位元組向量，並把 64 執行緒塊加入調優候選；駐留判斷由 SM 的執行緒數、常駐 CTA 上限與暫存器共同決定。(2048, 4096) 的 bf16 ArgMin 從 8.704 微秒降到 7.393（對比項 FlagGems 8.543，快 1.155 倍），背靠背場景 7.808 降到 6.000；f16 ArgMax 從 6.816 降到 6.400（對比項 9.344，快 1.46 倍）。12 行基準裡 4 行變快、0 行變慢。

#2820 把「每個元素只讀一次」的運算元改回預設快取策略：逐元素、rope、norm、softmax、歸約類核心不再一律用 evict-last/evict-first，只在真正重複讀的地方保留（二元核心的廣播行、多批多頭共享的 rope cos/sin 行、兩遍 RMSNorm 的兩次讀、融合加法的前兩遍）。同時給 RMSNorm 加了配對核與分階段核兩種實現、把 warp 行 Softmax 的 CTA 縮到剛好填滿 SM 的 warp 數、2D neox RoPE 改用 128 執行緒，並把每架構的常駐 CTA 上限收進 RESIDENT_CTAS。代價是 Softmax 的 (32768, 4096) 從 121.792 退到 131.648 微秒（對比項 quack 126.080），作者在表裡如實標出。

#2822 是 L2 刷寫的只讀改法，與主倉 #3534 同源，168 項測試透過，並額外記錄刷寫方法、只在方法一致時與歷史讀數比較。#2825 用單一重擬合的代價模型給 W4A16 的合法配置排序，split-K 的 CTA 只暫存自己那份量化後設資料，共享記憶體記賬把啟動餘量算進去。#2826 是配套文件，寫清稠密與 tiny-M 的標定契約（測量方式、擬合約束、留出驗證），沒有執行時改動。#2827 把 off-CUDA 套件挪到夜間作業：原來 preflight 裡這兩步在 ubuntu-latest 上要 13 到 18 分鐘、比其他檢查晚十分鐘出結果，而其中 tests/ops 的冒煙跑用的是運算元參考實現，檢查的是裝置無關性與派發、不是核心結果；挪走後 preflight 只保留收集檢查、丟棄測試報告與 cuda_only 必失敗檢查（約 3 分鐘），1856 加 2789 個測試節點轉到夜間，釋出門禁仍要求夜間成功。

來源：https://github.com/tile-ai/TileOPs/pull/2819、https://github.com/tile-ai/TileOPs/pull/2828、https://github.com/tile-ai/TileOPs/pull/2820

3.2 TileOPs 在審：一元逐元素單波次啟動、FP8 分組 GEMM 連續佈局、GEMM 可選 bias（#2829 / #2824 / #2823）

#2829 讓「體內只有一兩條單指令操作」的一元逐元素啟動跑成單波次：讀寫位元組數不超過 L2 的四分之三、結果與運算元同寬、體內最多兩個各降為一條指令的操作，滿足三條才縮成 256 執行緒塊的一波；體內條件由 TIR 變數上的節點計數判斷，帶 cast、除法、數學呼叫、守衛或不能在裸變數上求值的體一律保留預設啟動。目前符合的是 relu、abs、neg、bitwise_not 與 bool 上的 logical_not。順帶修了一個真問題：ReLU 對 NaN 輸入返回 0，現在返回 NaN（16 位用 torch 的位模式、float32 用規範 NaN）。基準裡三行 ReLU 與一行 LogicalNotFwdOp 變快，其餘配置不變；作者也標出 GPU 6 上帶刷寫的 decode 行只是打平。

#2824 給 SM90 的 FP8 分組 GEMM 加上連續佈局（緊湊與對齊兩種打包、段尾或逐行兩種專家後設資料），SM89 仍只支援掩碼佈局，並補齊按專家大小偏斜的實測負載與 DeepGEMM 連續佈局的對比。#2823 給 GEMMFwdOp 加可選 bias[N]，覆蓋稠密、GEMV 與 split-K 三條路徑，附 cuBLASLt BIAS 對比。兩筆都對應 #2130 裡列的缺口，說明 TileOPs 正在按「與 torch 生態逐項對齊」的清單推進，而不是隨機加運算元。

來源：https://github.com/tile-ai/TileOPs/pull/2829、https://github.com/tile-ai/TileOPs/pull/2824、https://github.com/tile-ai/TileOPs/pull/2823

3.3 TileFoundry：FSDP 與 HSDP 的訓練語義進 HIR 求值器（#284）

合於 11-04 16:36。此前的求值器把 reshard 當成全域性張量處理，丟掉了不同參與者的貢獻差異，於是訓練語義（尤其是分片與複製混合的 HSDP）在 HIR 層無法正確求值。這筆實現 #255 的訓練子集：按參與者求區域性值、GPU 座標、呼叫與迴圈、gather/scatter，以及沿軸的 Partial(sum) 歸約，都走共享的投影規則；Reshard 的別名證明改成對歸屬敏感，接上 #170 的結構化生命週期，元組實參從張量葉子校驗。配套加了一個受 modded-nanogpt 啟發的兩塊 GQA 研究：顯式前向反向、動量 SGD、獨立的 PyTorch 參考，覆蓋 FSDP、2×2 與 4×4 的 HSDP、兩種駐留策略，報告集合通訊組、載荷與儲存佔用，並嵌入可複用的 SVG，README 裡的每張表都能從程式碼推出來。

驗證：重構後的核心套件 775 項、後端套件 778 項（13 跳過，兩張 sm90 卡），10 組 HIR 配置在 CPU 上透過；兩種駐留策略在四張 sm90 上與真實的 2×2 PyTorch HSDP 對齊（跑兩次更新，比較區域性梯度、引數與動量）。

來源：https://github.com/tile-ai/TileFoundry/pull/284

3.4 TileFoundry 在審：wgmma 同步計劃（#286），另開一條命名統一的 issue（#287）

#286 處理的是 WGMMA 的非同步性：指令發出之後仍可能繼續讀共享記憶體與暫存器，所以後續訪問必須排在它的效果之後。做法是給每個運算元效果宣告完成條件，把 WGMMA 的 fence、commit、wait 掛到需要它們的 HIR 節點上，lowering 在指令的發射點消費這些步驟；若後續作用域放不下一個集合等待，就在發射作用域裡把操作做完。同時修了暫存器 Reshard 的別名，讓每個元素在持有副本的每個執行緒上都留在同一個暫存器，計劃用這個儲存身份追蹤依賴。作者明確寫了一條風險：矩陣指令的發射器還沒接上，這筆不證明 WGMMA 的執行正確性，也不給效能，TMA 的完成語義另行處理。

#287 是一條文件向 issue，指出程式碼庫裡「一個概念五個名字」（mesh 的一個成員被叫作 placement 133 處、program 110 處、instance 92 處、participant 44 處，另有 wave_units 一類帶 unit 的寫法），而 unit 這一個詞又同時指空型別、硬體計算單元與 mesh 成員；Python 與 CUDA 執行時對 local_layout 的簽名還不一致。這類 issue 通常不會立刻改程式碼，但它是編譯器專案從「能寫」到「能被人長期讀」的必經一步。

來源：https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/issues/287

3.5 採用方與第三方：TileKernels 與 FlashQLA 無推送，三個第三方倉在 24 小時內活躍（11-04）

兩個採用方倉在本視窗都沒有提交：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 上的屏障佔用回退，說明編譯器的改動會以「被動迴歸」的方式作用到它們的形狀上。

第三方側有三條由倉庫發現撈到的活躍項，都是 24 小時內新建或推送：sablin39/tilelang-cuda-skills（149 星，給編寫 TileLang 與用 CUDA 工具鏈除錯提供技能集）、xwhzz/tilelang-hexagon（2 星，把 TileLang 往高通 Hexagon DSP 上帶）、Index4NekoMeteor/Tilelang-AttnRes（1 星，用 TileLang 實現 Kimi K3 的注意力殘差結構）。三條都很小，但方向有意思：一條做工具鏈側的開發體驗，兩條做新目標架構與前沿模型結構的移植，屬於社群自發擴充套件 TileLang 邊界的那類工作。

來源：https://github.com/deepseek-ai/TileKernels、https://github.com/QwenLM/FlashQLA、https://github.com/sablin39/tilelang-cuda-skills

四、社群、教程與活動

4.1 新增 issue 4 條：warp 尺寸 64 的證明、PTO 有符號原子、CuTeDSL 的 NVFP4 量化器（11-04）

主倉三條。#3791 是「[BUG][ice-on-valid-code]」：ThreadSync 只在 warp 尺寸 32 下能證明 warp 一致謂詞，同一份 IR 在 thread_warp_size 為 64 時被拒——這條與本期在審的 #3787 咬得很緊，說明證明器對 warp 尺寸的假設是硬編碼的，而 AMD 與部分國內加速器是 64 寬 warp。#3789 報 PTO 後端為整數 T.atomic_add 發射帶符號的 pto.si32/si64 導致校驗失敗，當天就有 #3792 與 #3794 兩個方向的修法。#3781 報 CuTeDSL 的 NVFP4 block-scaled 量化器無法編譯：打包 FP4 的重解釋與缺失的 ceil/floor。加上 TileFoundry 的 #287（命名統一），本視窗共 4 條新 issue，全部是「內部實現細節洩漏到使用者可見面」的型別，沒有一條是效能抱怨。

來源：https://github.com/tile-ai/tilelang/issues/3791、https://github.com/tile-ai/tilelang/issues/3789、https://github.com/tile-ai/tilelang/issues/3781

4.2 新聞側：四組查詢、HN 與 arXiv 均無視窗內命中（11-04）

Google News 的四組查詢（英文 TileLang GPU kernel、"TileLang" OR "tile-lang"，中文 TileLang 運算元、TileLang 編譯器）裡，最新一條命中停在 352 小時前（10-21 SemiAnalysis 的華為與 DeepSeek 報道），其餘集中在 09-30 那一波「DeepSeek 開源昇騰版 TileLang、DeepGEMM、DeepEP」的報道（856 小時前），全部早於本視窗；Hacker News 以 tilelang 檢索最近 24 小時為 0 條；arXiv 的 all:tilelang 最新命中仍是 07-24 的 TileSight，視窗內沒有新論文。GitHub 倉庫發現撈到的 9 條候選裡，7 條是倉庫自身，1 條是 Swarm World 這類與主題無關的 Show HN，1 條是無關的 2048 類專案。可以確認本視窗新聞側為零，報告主體只能來自提交與 PR——這已經是連續多期的穩定形態。

來源：https://news.google.com/、https://news.ycombinator.com/、https://arxiv.org/

五、趨勢觀察

第一，昇騰 PTO 的清賬模式已經定型，它的收益與代價都值得記錄。#3512 在 10-14 合入，三週內不斷有真實運算元把它踩出邊界，今天的五筆補齊只是最新一批。這種「先落架構、後按運算元清迴歸」的節奏在 #1975 整倉遷移時也出現過，好處是架構迭代快、單筆改動可控，壞處是主幹在相當長的視窗裡處於「已知有幾處不工作」的狀態。五個補丁刻意保持互不依賴，說明維護者有意識地在降低迴滾粒度——這比攢成一個大回滾包要成熟，但也意味著需要更強的釋出前回歸紀律，否則「獨立可合」會變成「獨立可漏」。

第二，效能度量開始被當作產品特性來對待。主倉 #3534 與 TileOPs #2822 在同一天把 L2 刷寫從寫零改成只讀，理由都是「髒行寫回汙染下一輪計時」；TileOPs 更進一步，記錄刷寫方法、只在方法一致時比較歷史讀數。這不是小修小補：一個基準數字如果在不同版本間不可比，那麼所有「快了多少」的結論都要重算。把度量方法本身納入版本管理，是效能工程走向可信的前提。

第三，運算元庫的競爭點從「運算元數量」轉向「選擇器質量」。TileOPs 今天合入的 W4A16 單一重擬合代價模型與稠密 GEMM 的標定契約文件，加上在審的 FP8 分組 GEMM 連續佈局與 GEMM bias，構成一條清晰的線：先把每個形狀的合法配置列舉出來，再用一個統一的代價模型排序，標定過程寫成可復現的契約。這條線做對了，運算元庫的效能就不再依賴個別核心的手工調優；做不對，就會退化成「某些形狀特別快、其他形狀靠運氣」。

第四，架構支援的邊界在兩頭擴。一頭是新卡：B300 的 sm_103a、gr100 的 sm_107a 進白名單，SM120 的 FP8 換成 block-scaled MMA；另一頭是國產加速器：海光在 64 bank 的 LDS 相位上做微架構級調優，昇騰在核心比與異常轉儲這類工程細節上補齊。這兩種擴法的工作量結構完全不同——新卡主要是「讓路徑能走通」，國產卡更多是「讓同一條路徑跑出接近硬體的效率」，後者需要的實測密度高得多，也正是海光那四筆提交裡測試檔案佔比那麼大的原因。

最後一點風險面：一個視窗內 42 筆在審 PR 更新、其中 12 筆帶昇騰/PTO 標籤、7 筆帶 CUDA/SM/CuTeDSL 標籤、4 筆集中在 Profiler，這個結構說明倉庫的變更壓力正從「加功能」轉向「修邊界」。對使用者來說，好訊息是缺陷在被逐個指認；壞訊息是這些缺陷大多落在驗證器、快取鍵、證明器這類「平時看不見」的地方，等它們顯形時往往已經影響到結果的可信度。

附、素材與核查說明

信源核查表

信源 核查結果
tile-ai 組織（29 倉庫） 視窗內推送 8 個倉庫
主倉 tilelang 主分支 3 筆合入；新開 issue 3 條；在審更新 38 筆（昇騰/PTO 12、CUDA 與 CuTeDSL 7、Profiler 4）
TileOPs 主分支 7 筆合入；在審更新 3 筆；無新 issue
TileFoundry 主分支 1 筆合入；在審更新 1 筆；新開 issue 1 條
tilelang-ascend 預設分支 ascendc_pto 1 筆合入；在審清單因匿名介面限額未取到，改以主分支與提交原子源核查
tilelang-mlir-ascend 主分支 2 筆（含 #209 合入）
tilelang-hygon 主分支無新提交（停在 11-02）；feat/hcu-scaled-cast 分支 4 筆
tilelang-metax / tilelang-musa 視窗內無推送
TileOPs-nightly 僅分支推送，主分支無新提交
釋出頁（tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon） 視窗內無新版本，tilelang 最新仍為 09-30 的 v0.1.15
FlashQLA / TileKernels 視窗內無推送
Google News（中英文四組查詢）/ Hacker News / arXiv 無視窗內實質命中，最新命中為 352 小時前

完整信源清單

*[24]** tilelang-ascend 異常轉儲相容 CANN 9.2.0（#1956）— https://github.com/tile-ai/tilelang-ascend/pull/1956

核查說明

本期所有合入、在審、issue 與提交數量均來自 GitHub 檢索與原子源的實測輸出，未做估算；昇騰倉與海光倉的在審 PR 清單因匿名介面額度耗盡未能取到，相關結論只基於主分支提交與分支級原子源，已在信源核查表中註明。tilelang-hygon 的四筆提交位於未開 PR 的功能分支，屬分支級動態，不代表主線已合入。Google News、Hacker News 與 arXiv 三條新聞通道在本視窗均無命中，報告主體因此全部來自程式碼側。