調研視窗:2026-10-28 07:00 ~ 2026-10-29 07:00(北京時間,24 小時)。

信源:GitHub tile-ai 組織逐倉核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文件站)+ 合入與在審真值檢索(合入 24 筆、在審更新 48 筆)+ 新增 issue 全量核對(6 條)+ 提交級 diff 與檔案清單核查 + 釋出頁與 tag 核查 + 第三方後端倉核查 + Google News 中英文六組查詢(走代理)+ Hacker News + arXiv。


本期索引

  • 今日重點:昇騰 910 A2/A3 的 NPUIR 後端併入 mlir-ascend 主幹,504 項自測全過;TileOPs 選擇器改按忙時校準,manifest 遺憾從 5.79% 降到 0.48%;24 小時合入 24 筆、在審更新 48 筆、新開 issue 6 條
  • 一、核心專案進展
    • 1.1 主倉:標量讀的索引係數不再混進 lane 數,行寬 32768 以上的編譯崩潰關閉(#3453)(10-28)
    • 1.2 主倉:批次 all-reduce 按標量分支解析執行緒區間,batch>1 時 2048 個輸出全錯的問題修復(#3341)(10-28)
    • 1.3 主倉:SIMT 原子區域的形狀校驗補上,越界區域在編譯期拒絕(#3465)(10-28)
    • 1.4 主倉:稀疏 MMA 的 K 尾塊不再被靜默放過(#3367)(10-28)
    • 1.5 主倉:atomic_addx4 的目標 dtype 收窄到 float*(#3369)(10-28)
    • 1.6 主倉:非 restrict 引數的 handle 字尾不再被裁掉(#3366)(10-28)
    • 1.7 主倉:TMA 描述符只記錄已校驗的名字(#3484)(10-28)
    • 1.8 主倉:profiler 輸出索引統一校驗,標量與列表形式同一套邊界規則(#3486)(10-28)
    • 1.9 主倉:T.Parallel 的 coalesced_width 統一轉 IntImm(#3374)(10-28)
    • 1.10 主倉:測試基建把 cache 敏感用例改走 without_kernel_cache(#3695 / #3696)(10-28)
    • 1.11 新增缺陷三條:SM120 上 13 個示例失敗、IfStmtBinding 指數級編譯、without_kernel_cache 不強制重編(#3701 / #3700 / #3699)(10-28)
  • 二、多後端適配(昇騰 / 海光 / ROCm)
    • 2.1 昇騰:910 A2/A3 的 NPUIR 後端併入 mlir-ascend 主幹,504 項自測全透過、72 個示例僅 1 例環境性失敗(#207)(10-28)
    • 2.2 昇騰主倉:Bisheng 構建身份進二進位制快取鍵(#3691)(10-28)
    • 2.3 昇騰主倉:手寫流水線保留配置暫存器相位(#3624)(10-28)
    • 2.4 昇騰倉在審:W4A4 int4×int4 GEMM 結果錯誤的根因與修法(#1731)(10-28)
    • 2.5 昇騰倉在審:a2a3 排程器遷移,100 檔案 18453 行(#1951)(10-28)
    • 2.6 昇騰倉在審:Carver 按所選 SDK 讀 A3 資源預算(#1955)(10-28)
    • 2.7 昇騰倉在審:AI Core Exception Dump 的依賴降到 CANN 9.2.0(#1956)(10-28)
    • 2.8 昇騰倉在審:文件改用 T.Tensor,棄用 T.Buffer DSL 介面(#1957)(10-28)
    • 2.9 海光:HCU 普通複製的 LDS swizzle 特性分支,18 檔案含 5 個新測試(#30)(10-28)
    • 2.10 ROCm:pip 版 SDK 的 ROCm 10 支援合入,MI355X 不再被誤判成 gfx900(#3549)(10-28)
  • 三、生態與採用方(TileOPs / TileFoundry)
    • 3.1 TileOPs:選擇器改按忙時校準,manifest 遺憾 5.79% 降到 0.48%(#2784)(10-28)
    • 3.2 TileOPs:非 NT 佈局改走轉置持久核心,2112×7168×4096 TN 從 195.3 降到 175.3 微秒(#2789)(10-28)
    • 3.3 TileOPs:cumsum 長行改兩段式種子掃描,(8,65536) 從 12.75 降到 5.18 微秒(#2790)(10-28)
    • 3.4 TileOPs:extrema 歸約改用 PTX 的 max.NaN,NaN 傳播不變而回歸消失(#2792)(10-28)
    • 3.5 TileOPs:並行 cumsum 的有效行 tile、GDN 窄預填充分塊兩筆(#2788 / #2786)(10-28)
    • 3.6 TileOPs 新開議題:背靠背呼叫下大 GEMM 行落後最快基線(#2791)(10-28)
    • 3.7 TileFoundry:ring 階段繫結與編譯流程統一,70 檔案(#277)(10-28)
    • 3.8 TileOPs 在審:釋出映象把 TileOPs 0.1.0 與 TileFoundry 0.0.3 打在一起(#2793)(10-28)
  • 四、社群、教程與活動
    • 4.1 文件站兩筆自動更新(10-28)
    • 4.2 昇騰倉文件事實核查與安裝修復批次在審,七條同批推進(10-28)
    • 4.3 昇騰倉新增運算元示例在審:GroupNorm 與 fla

_chunk 兩個運算元(#1607 / #1953 / #1952)(10-28)

  • 4.4 第三方:社群自動調優代理重新整理中文說明(#74),Hexagon 後端倉庫主分支視窗內無新提交(10-28)
  • 4.5 新聞側:Google News 六組中英文查詢、Hacker News 與 arXiv 均無視窗內實質命中(10-28)
  • 五、趨勢觀察
  • 附、素材與核查說明

今日重點

今天視窗的重心在昇騰側。tilelang-mlir-ascend 把 Ascend 910(A2/A3)的 NPUIR 後端整條併入主幹,一筆 55 檔案、21490 行新增:Python 側是 tilelang/ascend910 包(backend / adapter / codegen / compiler / execution_backend / jit / pipeline / runtime / target 加 language 下的 kernel、copy、gemm、npuir 方言面),C++ 側是 src/ascend910/target 下的開發者與專家兩套發射器、op 層與 LowerNpuirBlock、NpuLoopVectorize、RewriteFDivToMul 三個 pass,執行時補了 npu_launcher.h 與 npu_utils,安裝走 install_npuir.sh,target 別名接受 ascend910 / ascend-910 / a2 / a3 / npuir。作者給出的驗收是倉庫自帶的兩套用例:pytest testing/npuir 504 項全透過、examples 72 例中 71 例透過且 0 超時 0 報錯,唯一失敗的 deepseek_v32/fp8_lighting_indexer.py 在舊後端上同機同樣失敗,屬環境與容差問題。第二筆提交是逐條對齊舊後端的行為差異,重點在 bf16 合法化(HIVM 不接受 bf16 運算元,改為 fp32 計算後顯式轉換)與動態 extent 的複製下沉。

TileOPs 今天合入七筆,兩條主線都在效能側。選擇器不再貼著 manifest 擬合,改為在忙時(device_busy_ms)上按 160 個稠密 shape、60 個 m 9–32 的 shape、264 個 m≤8 的 shape 重新標定,用 softmin 期望對數遺憾做座標下降;端到端忙時遺憾在留出的 manifest 上從 5.79% 降到 0.48%,m≤8 的留出集從 14.18% 降到 1.79%。GEMM 側新增 coop2_swap:把非 NT 佈局換成轉置的持久核心,兩個消費者 warpgroup 各持 128 行,尾寫經 stmatrix.trans 與 TMA 轉置落回,並支援寬窄 tile 混排(例如 2112 = 13×160 + 32 拆成 160 寬 tile 加 128 窄 tile),2112×7168×4096 TN 從 195.3 微秒降到 175.3 微秒,同機上另一塊 H200 讀到 174.3 降到 158.5 微秒,均壓過 QuACK 的 187.9 與 167.0 微秒。

歸約側則同日先壞後修。上午 #2787 用 int32 鍵實現 extrema 歸約的 NaN 傳播、#2788 改並行 cumsum 的行 tile,兩筆都在 07 時前後落地,隨後 #2790 與 #2792 分別把兩個迴歸收掉:cumsum 改成按種子分段、兩次發射((8,65536) float32 從 12.75 微秒降到 5.18 微秒,同時修掉主幹的 65535 tile 以上發射失敗),extrema 改用 PTX 的 max.NaN / min.NaN 並在非整 16 位元組向量的行上走單獨核心(amax (2048,4096) bf16 從 8.352 微秒回到 6.976 微秒,快於 FlagGems)。這類”同一天內引入再修掉”的節奏說明 TileOPs 的夜間基準已經把迴歸攔在了當天。


一、核心專案進展

1.1 主倉:標量讀的索引係數不再混進 lane 數,行寬 32768 以上的編譯崩潰關閉(#3453)

合於 10-28 11:06,單檔案 +8/-5。一個用向量化方式讀整行、同時又按下標讀該行單個元素的 kernel,在行寬對 65536 取模後不小於 32768 時編譯報 InternalError: Can’t fetch the lanes of a scalable vector。根因在 VectorTypeAccessChecker::OnArrayAccess:它把標量讀記為 access_dtype.with_lanes(me->coeff),而 me->coeff 是讀下標的模係數,對 x[row, 0] 就是行寬;DataType 的 lane 數只有 16 位並按 int16 讀回,32768 到 65535 因此被當成可伸縮向量,lanes() 在 get_preferred_dtype 裡直接拋錯,更寬的倍數則靜默迴繞。修法是把係數單獨存進 int64_t,相容性判斷改用它取模,標量讀的 dtype 只保留元素型別。這個缺陷的原始觸發場景是 TileOPs 的 GroupNorm:Stable Diffusion 的 40960 與 163840 兩個組寬都會踩到。驗證用兩次 wheel 構建對比:主幹在 32768 與 40960 兩個寬度上編譯失敗,本筆在所有測試寬度上都能構建。

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

1.2 主倉:批次 all-reduce 按標量分支解析執行緒區間,batch>1 時 2048 個輸出全錯的問題修復(#3341)

合於 10-28 05:46,2 檔案 +170/-8。標量 all-reduce 分支在發射呼叫前會先把參與執行緒的區間解析出來(reducing_threads 大於 32 且目標支援命名柵欄時走 ResolveAllReduceThreadRange),批次分支卻直接取塊邊界當偏移與柵欄寬度。當歸約執行緒是 [32, 96) 這種不從 0 開始的區域性區間時,批次呼叫從 0 號執行緒索引、按整塊等柵欄,非參與執行緒未初始化的暫存器被折進歸約:batch=1 正確,batch=2 在同一 kernel 上 2048 個輸出元素全部錯,最大絕對誤差 1083,且沒有任何診斷。修法是讓批次分支與標量分支同一套解析,把解析出的偏移與範圍傳給 MakeBatchAllReduce,工作區 stride 仍保持塊大小。新增測試在同一 kernel 上跑 batch=1 與 batch=2 並對 torch 參考比對。

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

1.3 主倉:SIMT 原子區域的形狀校驗補上,越界區域在編譯期拒絕(#3465)

合於 10-28 08:14。SIMT 原子操作此前不校驗傳入區域的形狀,形狀不匹配的寫法一路走到 codegen 才表現出錯誤結果或崩。本筆在 CUDA 與 ROCm 兩條路徑上補齊區域形狀校驗,把不合法區域提前到編譯期拒絕。

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

1.4 主倉:稀疏 MMA 的 K 尾塊不再被靜默放過(#3367)

合於 10-28 16:01。稀疏 MMA 在 K 方向上有未處理的尾塊時,此前不報錯而是沿用整塊的搬運與計算引數,數值靜默偏差。本筆顯式拒絕這類 K 尾塊,把問題挪到編譯期。

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

1.5 主倉:atomic_addx4 的目標 dtype 收窄到 float*(#3369)

合於 10-28 16:08,3 檔案 +27/-9。CUDA 的 atomic_addx4 與 atomic_addx4_ret 通用實現此前接受任意目標指標,int32 或 int64 目標會落到 float 語義上算錯。本筆把兩個通用過載限制為 float* 目標,fp16 與 bf16 的既有特化不動;Python 側 API 保持型別中立,CPU 後端的整數目標不受影響。CUDA 側新增整數目標編譯失敗的迴歸用例,CPU C 後端新增兩種整數寬度的累加正確性用例。

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

1.6 主倉:非 restrict 引數的 handle 字尾不再被裁掉(#3366)

合於 10-28 13:19。引數去重時按去掉 _handle 字尾的名字比較,導致同一個核心裡 A 與 A_handle 兩個不同引數被當成同一個,程式碼生成只發一份。修法改為按精確名加變數身份去重,並新增 CUDA 側 A 與 A_handle 並存的迴歸用例。

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

1.7 主倉:TMA 描述符只記錄已校驗的名字(#3484)

合於 10-28 17:12。TMA 描述符的登記此前在校驗之前完成,校驗失敗的描述符名留在表裡,後續同名的合法描述符會被誤判為重複。本筆把登記挪到校驗之後。

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

1.8 主倉:profiler 輸出索引統一校驗,標量與列表形式同一套邊界規則(#3486)

合於 10-28 19:24,2 檔案 +15/-9,修 #2884。profiler 與 adapter 兩側此前對 result_idx 的處理不一致:標量形式做過邊界檢查、列表形式沒有,負索引也未歸一。本筆讓 BaseKernelAdapter 與 Profiler 共用一套規則,接受 None、整數或整數列表,按 -len(params) 到 len(params)-1 判定,合法負索引歸一、越界與非整數元素拋 ValueError。新增 28 項 CPU 側迴歸覆蓋標量與列表兩種路徑。

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

1.9 主倉:T.Parallel 的 coalesced_width 統一轉 IntImm(#3374)

合於 10-28 16:27。coalesced_width 傳 Python 整數時,部分路徑直接以原生整數進入 pass,型別斷言在後續變換裡失敗。本筆在語言層統一轉成 IntImm,消除寫法差異帶來的失敗。

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

1.10 主倉:測試基建把 cache 敏感用例改走 without_kernel_cache(#3695 / #3696)

合於 10-28 01:59 與 09:19。承接昨日記錄的 issue #3694(核心快取命中破壞假定冷編譯的測試),兩筆把檢查編譯產物、原始碼與診斷資訊的用例改走 without_kernel_cache() 或顯式 .compile(…),讓斷言看到的是本次構建的結果而不是程序內快取裡的舊 kernel。

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

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

1.11 新增缺陷三條:SM120 上 13 個示例失敗、IfStmtBinding 指數級編譯、without_kernel_cache 不強制重編(#3701 / #3700 / #3699)

三條新開 issue 指向測試與編譯時間。其一,#3701 在 RTX PRO 6000 Blackwell(sm_120,動態共享記憶體上限 99 KiB)上跑完整 examples 套件,CUDA 與 CuTeDSL 兩個後端都是 13 失敗 199 透過 27 跳過,失敗集逐字相同:7 例是示例在非 –tune 路徑上硬編碼 Hopper 尺寸的 tile(如 attention_sink 要 163840 位元組、deepseek_mla 解碼要 231424 位元組),4 例是參照實現 flash-attn 2.5.8 沒有 sm_12x 核心,另兩例分別是 ReduceOp 下沉與 ThreadSync 無法證明執行緒範圍。其二,#3700 報告在迴圈裡重複寫 x = x + x 並加上 i < N 守衛時,tl.LayoutInference 從 0.087 秒漲到 118.5 秒,關掉 replayable bind 內聯即退回 0.100 秒。其三,#3699 指出 without_kernel_cache() 只翻轉 CacheState(只被磁碟快取讀取),而 @tilelang.jit 包裝器在 JITImpl.compile() 之前還會查程序內的 _call_form_cache 與 _kernel_cache,因此該塊內的一次普通呼叫並不保證重新編譯。

來源:https://github.com/tile-ai/tilelang/issues/3701

來源:https://github.com/tile-ai/tilelang/issues/3700

來源:https://github.com/tile-ai/tilelang/issues/3699

二、多後端適配(昇騰 / 海光 / ROCm)

2.1 昇騰:910 A2/A3 的 NPUIR 後端併入 mlir-ascend 主幹,504 項自測全透過、72 個示例僅 1 例環境性失敗(#207)

合於 10-28 08:04,55 檔案 +21490/-1,作者 CeleNewYear,基於 v0.1.15 的 a35f8ddf,兩個提交分別負責後端落地與行為對齊。後端麵包含 tilelang/ascend910 下的 backend、adapter、codegen、compiler、execution_backend、jit、pipeline、runtime、target,以及 language 下的 kernel、copy、gemm、npuir 方言面;C++ 側是 src/ascend910/target 的開發者與專家兩套發射器、src/ascend910/op 運算元層與 src/ascend910/transform 的 LowerNpuirBlock、NpuLoopVectorize、RewriteFDivToMul;執行時要 npu_launcher.h、npu_utils.cpp、npu_utils.py 生成 torch 相容的啟動包裝並按 *_infer_workspace_shape_function 估每塊 scratch;構建側接 CMakeLists 並給 install_npuir.sh 的 –bishengir-path 選項;target 別名接受 ascend910、ascend-910、a2、a3、npuir,模式開關 TILELANG_ASCEND_MODE 取 expert 或 developer,預設 expert。驗收資料:pytest testing/npuir 504 透過 0 失敗;examples/run_all.py –sequential 72 例中 71 透過 1 失敗,0 超時 0 報錯,唯一失敗的 deepseek_v32/fp8_lighting_indexer.py 在舊後端同機同樣失敗。第二筆提交對齊的差異包括 bf16 合法化(HIVM 拒絕 bf16 運算元,改為 fp32 計算後顯式轉換)與動態 extent 區域的複製下沉。

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

2.2 昇騰主倉:Bisheng 構建身份進二進位制快取鍵(#3691)

合於 10-28 10:26。昇騰側的二進位制快取此前只按核心與 target 做鍵,換 Bisheng 工具鏈版本後仍會命中舊二進位制。本筆把構建身份併入快取鍵,工具鏈變化即失效重編。

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

2.3 昇騰主倉:手寫流水線保留配置暫存器相位(#3624)

合於 10-28 09:57。手動流水線在改寫階段丟失配置暫存器的相位資訊,導致同一份配置在正確的拍上不生效。本筆在改寫中保留相位,恢復與自動流水線一致的時序語義。

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

2.4 昇騰倉在審:W4A4 int4×int4 GEMM 結果錯誤的根因與修法(#1731)

在審,7 檔案 +634/-24,修 #438。根因是整條 C++ 模板棧都用 sizeof(AscendC::int4b_t) == 1 推分形幾何,而 int4b_t 是半位元組打包型別:32 位元組的 C0 塊正確元素數是 64 而不是 32,於是所有複製 stride、repeat 次數與位元組計數全部 2 倍偏差。三層疊加:L1 到 L0A/L0B 的 LoadData 引數由該常數推出,資料堆到錯誤位置;B 運算元的 zN 到 nZ 轉置被 s4 專用指令 load_cbuf_to_cb_s4 靜默關掉,必須改走 LoadDataWithTranspose 到 load_cbuf_to_cb_transpose_s4,而其引數語義在 CANN 標頭檔案裡沒有文件,作者只能用探針 kernel 在裝置上逆向(正確引數 repeatTimes = N/64、srcStride = roundUp16(K)/64、dstGap = 3,且 s4 版本按 2048 位元組塊而不是文件寫的 512 位元組 fractal 工作);UB 到 GM 的位元組長度按每元素 1 位元組算,每行多搬一倍並混入相鄰行。mad_s4 指令本身正確,症狀是 99.6% 不匹配、輸出 -501/627 而正確答案 641/1785。

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

2.5 昇騰倉在審:a2a3 排程器遷移,100 檔案 18453 行(#1951)

在審,100 檔案 +18453/-666,作者 erhsh。把 A2/A3 的排程器按新架構遷移,是視窗內最大的一筆在審改動。PR 描述表格未填,具體邊界需按檔案清單核對。

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

2.6 昇騰倉在審:Carver 按所選 SDK 讀 A3 資源預算(#1955)

在審,3 檔案 +203/-27。在 Ascend910_9392 上 Carver 回落到 Ascend910A,把 256×256 的 FP32 累加器放到 256 KiB 的 L0C 預算上,而所選 SDK 只給 128 KiB。修法是識別 9392 與 9382 兩個具體 SKU、用既有的 ASCEND_HOME 選擇讀各自平臺的 INI(另留跨編譯用的顯式檔案選項),把 SDK 容量裁剪到當前後端 arena 同時保留原始容量、配置身份、BF16 支援、核數與 L2 屬性。六個定向用例覆蓋裝置、target 與顯式晶片探測、超出與合法的 L0 tile、以及 9382 關閉 BF16 的獨立配置;作者給出的迴歸是原生昇騰 CI 主套 1352 透過 3 跳過、示例套 153 透過。

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

2.7 昇騰倉在審:AI Core Exception Dump 的依賴降到 CANN 9.2.0(#1956)

在審,5 檔案 +49/-32。該特性此前宣告依賴 CANN 9.3.0,作者在 9.2.0 上驗證可行並修兩處相容:acldumpTensorInfo 補齊 addrType=ACL_DUMP_ADDR_RAW、placement、shapeNum、originShapeNum 四個必填欄位(9.2.0 的 acldumpSaveExceptionInfo 強校驗 addrType,原實現被 calloc 置零成 ACL_DUMP_ADDR_PTR 而被拒),並跳過 NULL 地址或零大小的 tensor;註冊介面從已棄用的 aclrtSetExceptionInfoCallback 換成 aclrtExceptionInfoCallbackRegister。解析側相容 msaicerr.py 的兩種產物(9.2.0 出 .npy,舊版出 .bin),文件的版本要求同步改為 CANN 不低於 9.2.0。

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

2.8 昇騰倉在審:文件改用 T.Tensor,棄用 T.Buffer DSL 介面(#1957)

在審,單檔案 +124/-122。程式設計指南 3.1 節的 kernel 引數定義與上游 API 對齊:文字與示例統一用 T.Tensor,並補一句說明 T.Buffer(…) 與 T.Buffer[…] 的棄用只針對 TileLang DSL 介面,底層 TVM 的 Buffer 型別與概念並未棄用。依據是 ascendc_pto 分支 tilelang/language/proxy.py 裡的 @deprecated 裝飾器。

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

2.9 海光:HCU 普通複製的 LDS swizzle 特性分支,18 檔案含 5 個新測試(#30)

在審,18 檔案 +2000/-84,分支 Feat/hcu plain copy lds swizzle 於 10-28 13:59 推送。改動集中在新增的 gemm_lds_strategy_utils(+148/-0 與 +29/-0)、annotate_mls_gemm_dep.cc(+332/-45)與 ds_read_format.cc(+61/-22),copy 與 merge_shared_memory_allocations 有配套小改;測試側一次新增 test_hcu_plain_copy_lds_swizzle(396 行)、_formula(237)、_formula_b8(220)、fp8(200)與 _b8_write_phases(80)。方向是普通複製路徑的 LDS 寫相位與 swizzle 策略,配 b8/fp8 兩類資料寬度。

來源:https://github.com/tile-ai/tilelang-hygon/pull/30

2.10 ROCm:pip 版 SDK 的 ROCm 10 支援合入,MI355X 不再被誤判成 gfx900(#3549)

合於 10-28 07:38,2 檔案 +106/-19。ROCm 10.x 的 rocm/pytorch 映象把 ROCm 以 Python wheel 形式釋出(rocm-sdk 與 _rocm_sdk_devel),既沒有 /opt/rocm 目錄樹也沒有 rocm-* 的 apt 包。get_rocm_arch() 此前硬編碼 /opt/rocm 字首,路徑不存在時回落到 gfx900,於是 MI355X(gfx950)上 rocwmma 的靜態斷言直接失敗;設 ROCM_PATH 也無效,因為 TVM 的 tvm_callback_rocm_get_arch 用預設字首呼叫該函式。修法是按呼叫方字首、find_rocm_path()、torch 的 gcnArchName、gfx900 兜底的順序解析,每次 rocminfo 探測加 10 秒超時;另新增 docker/Dockerfile.rocm10,基於 rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12,去掉 rocm-* 的 apt 安裝、把 ROCM_PATH 指向 rocm-sdk path –root、保留基礎映象的 torch 構建並加 import 自檢。作者在 MI355X 上實測架構探測從 gfx900 變為 gfx950,testing/python/amd 套件按正確架構構建並透過。

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

三、生態與採用方(TileOPs / TileFoundry)

3.1 TileOPs:選擇器改按忙時校準,manifest 遺憾 5.79% 降到 0.48%(#2784)

合於 10-28 10:51,4 檔案 +259/-203。稠密選擇器的打分改為一個 tensor_core_tflops 乘結構因子,此前結構差異只體現在一個 max() 裡,記憶體受限的 shape 會把 basic、simple 與 coop2 打成分;launch_us 項刪除(在忙時上擬合為 0)。m≤8 的 NT 帶用流式模型替換手測閾值:每 CTA 速率取 min(ns·rows·bk·2/latency, cta_gbps),波次按 SMEM 受限的駐留算,再加 DRAM 下限 n·k·2/dram_tbps 與 split-K 的歸約項。TN/TT 的持久核心路徑打通(A/B 的 SMEM tile 跟隨儲存佈局,transpose_A/B),176/192 寬度、窄 tile、ping-pong 與 split-K 仍限 NT。標定在忙時上做,三組網格 160/60/264 個 shape、7 輪交錯 CUPTI 取中位,softmin 期望對數遺憾按座標下降、tau 0.1→0.05→0.02 退火。端到端結果:留出的 manifest 5.79% 到 0.48%,標定網格 9.28% 到 2.65%,m 9–32 3.13% 到 2.09%,m≤8 11.20% 到 1.17%,m≤8 留出集 14.18% 到 1.79%。

來源:https://github.com/tile-ai/TileOPs/pull/2784

3.2 TileOPs:非 NT 佈局改走轉置持久核心,2112×7168×4096 TN 從 195.3 降到 175.3 微秒(#2789)

合於 10-28 21:06,4 檔案 +620/-15,解決 #2785。新增 coop2_swap:對 TN/TT/NN 計算 Cᵀ = op(B)ᵀ·op(A)ᵀ,按 256×block_n 的 tile 走持久核心,兩個消費者 warpgroup 各持 128 行,尾寫經 stmatrix.trans 與 TMA 轉置寫回。混排 tile 寬度(narrow_n 與 row_mix)讓 M 不被 padding 浪費,例如 2112 = 13×160 + 32 拆成寬 160 加窄 128,寬的先跑、窄的收尾,拆分點沿用 coop2 窄 tile 的 _busiest_cta 指標。選擇器在 NT 之外列舉 coop2_swap 的 block_n 128 與 160、結構因子 1.2 並按忙時重擬合。H200 忙時:2112×7168×4096 TN 從 195.3 降到 175.3 微秒(QuACK 187.9),24576×1536×4096 TN 從 421.9 降到 416.4 微秒;另一塊 H200 上分別為 174.3 降到 158.5(QuACK 167.0)與 384–402 微秒區間。行組規則把固定 4 改成 16,在 24 個 TN/TT/NN 標定 shape 上與 group 4 相差 0.989–1.022 倍;另有 9 個 TN shape 上混排比均勻 coop2_swap 快 3.6% 到 11%。

來源:https://github.com/tile-ai/TileOPs/pull/2789

3.3 TileOPs:cumsum 長行改兩段式種子掃描,(8,65536) 從 12.75 降到 5.18 微秒(#2790)

合於 10-28 20:45,4 檔案 +255/-246。少量長行的 cumsum 改用種子分段掃描,發射次數從三次降到兩次,既消掉 #2788 引入的迴歸,也修掉其佈局失敗。規則是行不短於 64 KiB 且行數不多時走分段掃描,cumprod 不變;另修掉主幹上行數超過 65535 個 tile(如 (1, 2**24) float32)時的發射失敗。H200 @1500 MHz 忙時: (8,65536) float32 12.75 到 5.18 微秒,(16,65536) 19.14 到 7.01 微秒,(1,65536) bf16 9.03 到 3.45 微秒,(3,8193) fp16 7.52 到 3.29 微秒。384 個 shape 的掃描裡沒有一項比 #2788 之前更慢,CumsumFwdOp 與 CumprodFwdOp 的 manifest 行均快於最快對照實現;新增 few-long-seq 負載(8×65536 float32)5.20 微秒對 FlagGems 6.30 微秒。

來源:https://github.com/tile-ai/TileOPs/pull/2790

3.4 TileOPs:extrema 歸約改用 PTX 的 max.NaN,NaN 傳播不變而回歸消失(#2792)

合於 10-28 23:31,4 檔案 +346/-119。amax/amin 改為在 fp32 上用 PTX 的 max.NaN 與 min.NaN,不再走 #2787 的 int32 鍵方案,NaN 傳播語義不變而當天早些時候的迴歸被收掉;不是整 16 位元組向量的行另走一個 extrema 行核心,因為 T.reduce_max 在 fp32 上會丟掉 NaN。H200 @1500 MHz:amax (2048,4096) bf16 在 #2787 之前是 7.008 微秒、主幹 8.352、本筆 6.976;fp16 7.008 到 7.775 到 6.944;amax (64,32768) bf16 3.552 到 4.544 到 3.520;amin 同形 3.552 到 4.224 到 3.520。

來源:https://github.com/tile-ai/TileOPs/pull/2792

3.5 TileOPs:並行 cumsum 的有效行 tile、GDN 窄預填充分塊兩筆(#2788 / #2786)

#2788 於 10-28 07:13 合入,讓並行 cumsum 只處理有效的行 tile,是全視窗內最早的一筆,隨後被 #2790 以兩段式掃描替換掉問題路徑。#2786 於 07:09 合入,2 檔案 +30/-2,在 SM89 上把 32 列的 GDN 預填充計劃改用既有的帶融合輸出的狀態掃描,把整塊分段與尾塊分開處理,保留準備階段、輸出與最終狀態三個介面;選用條件為稠密輸入、batch 1、16 個 value head、狀態寬 128、key major、序列長大於 16384 且不超過 65536。H=4 時中位裝置執行時間相對 TileOPs 基線降 13.65% 到 18.31%,裝置活動跨度降 13.48% 到 17.07%;與 FLA 交替對比的四輪裡 H=16、長 65536、BF16 的跨度差異在 -2.57% 到 +1.77% 之間,作者自述該行落在測量波動內。

來源:https://github.com/tile-ai/TileOPs/pull/2788

來源:https://github.com/tile-ai/TileOPs/pull/2786

3.6 TileOPs 新開議題:背靠背呼叫下大 GEMM 行落後最快基線(#2791)

新開,承接 #2133 與 #2789。bench_gemm 單次讀一個呼叫、中間有 L2 沖刷與同步,呼叫從功耗上限以下起步,能耗代價被掩蓋;改成訓練裡的背靠背連續呼叫後,十條 manifest 大行裡有六條落後最快基線。典型一行:prefill-o-proj 4096×7168×16384 NT 本實現 1412.7 微秒,QuACK 1371.1、cuBLAS 1360.1、DeepGEMM 1359.2,比值 0.962。方法為單卡 H200、時鐘解鎖、預熱 1 秒後計時 1.5 秒、四輪交替取中位。

來源:https://github.com/tile-ai/TileOPs/issues/2791

3.7 TileFoundry:ring 階段繫結與編譯流程統一,70 檔案(#277)

合於 10-28 16:55,70 檔案 +1567/-1298,取代 #276。動機有兩處:未繫結的 ring 檢視與階段選擇會帶著沒有 SSA 名字的檢視走到 CUDA 運算元上;三個編譯入口各自拼裝不同的 pass 列表。改法是把 ring 階段儲存為一組已繫結的 TensorView、每個消費作用域只繫結一次,讓 default_pipeline 在 finalize、compile 與 scaffold 三處共用,核心宿主入口 pass 經 Target.resolve 解析並由 CUDA 提供 @bind 實現;Module.entry 成為宿主起點,分析的共享記憶體峰值帶進宿主啟動;呼叫鏈上的所有 CPU 宿主都參與連結,根宿主排在最前。契約變化:link_program 去掉 entries,載入器繫結第一個被連結的宿主,build 要求存在 CPU 宿主根。遺留風險寫在 PR 裡:#278 跟蹤 Broadcast 佈局歸一,#255 跟蹤 SymbolRef 與函式引用。

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

3.8 TileOPs 在審:釋出映象把 TileOPs 0.1.0 與 TileFoundry 0.0.3 打在一起(#2793)

在審,9 檔案 +120/-44。.github/runner/Dockerfile 新增 release 目標,在開發映象之上裝 tileops 與 tilefoundry 的 PyPI 釋出版,映象名以 TILEOPS_VERSION 與 TILEFOUNDRY_VERSION 命名,推為 ghcr.io/tile-ai/tileops-release:tileops0.1.0-tilefoundry0.0.3-cuda13.2-torch2.13-tilelang0.1.12;配套 scripts/verify_release_image.sh 拉映象後在 GPU 上跑棧與基線檢查、打包測試與指定的 –bench 檔案,用例按映象內 tileops 版本對應的 git tag 取(wheel 裡不帶測試與基準)。同時移除 sgl-kernel 基線:沒有任何映象再裝它,bench_moe_permute_align 一直跳過它。

來源:https://github.com/tile-ai/TileOPs/pull/2793

四、社群、教程與活動

4.1 文件站兩筆自動更新(10-28)

tilelang.github.io 在 10-28 08:20 與 13:31 各有一筆由 github-actions 提交的 Update docs,內容為文件構建產物隨主幹同步,無獨立內容變更。

來源:https://github.com/tile-ai/tilelang.github.io

4.2 昇騰倉文件事實核查與安裝修復批次在審,七條同批推進(10-28)

tilelang-ascend 的文件事實核查流水線在視窗內更新七條:T.tile.bitwise_lshift、bitwise_or、bitwise_and、fill、clear 五個 API 的核查,clamp 在 PTO 上生效並核對文件(#1620),以及 axpy 的 docstring 與測試加強(#1600)。另有 #1943 把 SHMEM 安裝改成構建失敗即停止,與 #1942 的快照分析器校驗同屬機器人提交。

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

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

4.3 昇騰倉新增運算元示例在審:GroupNorm 與 fla_chunk 兩個運算元(#1607 / #1953 / #1952)

三條示例類 PR 在視窗內更新:#1607 給 GroupNorm 加 20 個 cann-bench 用例,#1953 加 fla_chunk_o 運算元與測試,#1952 加 fla_chunk_intra_token_parallel 運算元與測試。這也是 1.1 那個標量讀崩潰缺陷的原始場景來源(GroupNorm 讀每行首元素)。

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

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

4.4 第三方:社群自動調優代理重新整理中文說明(#74),Hexagon 後端倉庫主分支視窗內無新提交(10-28)

社群倉庫 cc-c122/TileLang-Remote-Autotuning-Agent 於 10-28 09:21 合入 #74,按當前 V2 工作流重新整理中文 README,倉庫此前在 10-24 還有一次基線取消狀態與 mcprofiler 歸因的合併。另一個第三方後端倉庫 xwhzz/tilelang-hexagon 的組織檢視顯示 10-29 06:56 有推送時間戳,但主分支最新提交停在 08-24(CI 依賴升級),視窗內主分支無程式碼提交,昨天的 15 筆屬特性分支活動,本條不作為新增內容計入。

來源:https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent/pull/74

來源:https://github.com/xwhzz/tilelang-hexagon

4.5 新聞側:Google News 六組中英文查詢、Hacker News 與 arXiv 均無視窗內實質命中(10-28)

六組查詢包含 TileLang 精確詞加兩日、中文運算元詞加七日、英文 GPU kernel DSL 加三日、組織名加七日、中文廠商詞(昇騰 / 沐曦 / 海光)加七日以及運算元庫詞加兩日。返回結果裡與 TileLang 相關但落在視窗內的只有一篇轉載稿(昇騰 950PR 矩陣算力對比,屬二手市場評論,無原始資料,不收錄),其餘為 DeepSeek 融資與技術類泛報道,標題語義不含 TileLang 相關內容。英文查詢全部空返回。Hacker News 當日條目為一門名為 Vu 的新語言釋出,與 TileLang 無關。arXiv 無視窗內新增的 tilelang 或 tiled DSL 論文。新聞側本日與 GitHub 側的訊號強度差距繼續擴大,日報主體仍以倉庫為準。

來源:https://news.google.com/

五、趨勢觀察

昇騰側的推進方式這兩天已經很清楚:主倉按小時合入小修(快取鍵、暫存器相位),專用倉 mlir-ascend 用一次兩萬行的合併把 910 的 NPUIR 路徑從特性分支搬進主幹,tilelang-ascend 則在 PTO 與 AscendC 雙後端的既定節奏上做文件與運算元補齊。三處節奏不同但指向同一件事:910 的編譯器路徑正在從”能跑”轉向”可驗收”,504 項自測與 71/72 示例就是這次給出的驗收口徑。值得注意的是在審佇列裡 a2a3 的排程器遷移佔到了 100 檔案,說明下一階段的變動量在排程層而不在運算元層。

效能側的看點在測量口徑本身。TileOPs 這兩筆把選擇器和 GEMM 的結論都挪到了忙時(device_busy_ms),並在同一份 PR 裡給出單次呼叫與背靠背呼叫的差別:單次讀法下本實現領先 QuACK,背靠背、功率封頂的訓練場景裡十條大行有六條落後。這是一個比”某行快了多少微秒”更重要的訊號,因為訓練實際就是背靠背。#2791 把這件事單獨立項,說明團隊已經把口徑問題當成產品問題處理,而不是當作基準偏差粉飾過去。

工程衛生上今天有三條值得記:一是同一批內的自傷與自修(#2787/#2788 上午合入、#2790/#2792 下午收回歸),說明夜間基準的兜底在起作用,但也說明歸約這套程式碼的語義邊界仍然脆;二是 without_kernel_cache() 的語義漏洞(#3699)被從遷移過程中揪出來,程序內兩層快取不吃 CacheState,測試自以為在冷編譯其實拿的是舊 kernel,這類缺陷會讓”我們的測試全過”失去意義;三是 SM120 上 13 個示例失敗(#3701)暴露出示例硬編碼 Hopper 尺寸的問題,Blackwell 裝機量上來後這類”示例不是產品”的債務會被反覆計入。

附、素材與核查說明

信源核查表

信源 核查結果
GitHub tile-ai 組織逐倉(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文件站) 視窗內有推送:tilelang、TileOPs、TileFoundry、tilelang-mlir-ascend、tilelang-hygon、tilelang.github.io;tilelang-metax、tilelang-musa、TileRT、tilescale、DeepStack、tilelang-puzzles 無視窗內主分支提交
合入真值檢索(is:pr is:merged) 24 筆:tilelang 15、TileOPs 7、TileFoundry 1、tilelang-mlir-ascend 1
在審真值檢索(is:pr is:open updated) 48 筆,含昇騰倉文件核查批次與 a2a3 遷移
新增 issue 核對 6 條,其中 1 條為昇騰倉每日測試機器人
主分支提交級核查 tilelang 15 筆、TileOPs 7 筆、TileFoundry 1 筆、文件站 2 筆;合入數與主分支提交數一致
釋出頁與 tag tilelang 最新仍為 v0.1.15(09-30 釋出),視窗內無新版本;TileOPs 最新為 v0.1.0
第三方後端倉 Hexagon 後端主分支最新提交停在 08-24,視窗內無新提交;自動調優代理合入中文說明重新整理;沐曦倉六個分支無視窗內提交
Google News(中英六組) 僅一篇二手轉載稿涉及,無原始資料,不收錄
Hacker News 當日無 TileLang 相關條目
arXiv 無視窗內新增 tilelang 或 tiled DSL 論文

完整信源清單

thub.com/tile-ai/tilelang-hygon/pull/30>