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

信源:GitHub tile-ai 組織逐倉核查(tilelang / TileOPs / TileOPs-nightly / TileFoundry / tilelang-hygon / tilelang-metax / tilelang-ascend / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 文件站與夜間基準倉)+ 合入與在審真值檢索(合入 20 筆、在審更新 45 筆)+ 新增 issue 全量核對(7 條)+ 提交級 diff 核查 + 釋出頁與 PyPI 製品核查 + 外部後端倉核查(Hexagon / CANN 運算元倉 / PPU 包裝器 / 模型片段倉)+ Google News 中英文多組查詢(走代理)+ Hacker News + arXiv。


本期索引

  • 今日重點:TileOPs 0.1.0 正式釋出,206 個運算元、2265 行基準裡 2262 行不慢於最快實現;24 小時合入 20 筆、在審更新 45 筆、新開 issue 7 條;海光 scale 緩衝雙 bank 合入 1389 行;第三方 Hexagon 後端單日 15 筆
  • 一、核心專案進展
    • 1.1 主倉:ROCm 的 exp/sin/cos 不再預設降到快速內建,sin/cos 誤差從 7.8e-04 降到 6.6e-08(#3621)(10-27)
    • 1.2 主倉:自動目標改讀當前 HIP 裝置,多卡機上不再固定取 0 號卡架構(#3594)(10-27)
    • 1.3 主倉:TMA gather/scatter 的常量行索引轉為 int32,字面量寫法不再編譯崩(#3365)(10-27)
    • 1.4 主倉:NVRTC 與 CuTeDSL 按首個輸入張量所在裝置取當前流(#3359)(10-27)
    • 1.5 主倉:自動調優快取鍵納入 compile_flags,同一核心換編譯選項不再複用舊二進位制(#3338)(10-27)
    • 1.6 主倉:ThreadSync 把 address_of 記為寫,不透明呼叫後的柵欄補上(#3346)(10-27)
    • 1.7 主倉:執行時行程數的流水線 epilogue 直接展開,分頁解碼恢復可編譯(#3356)(10-27)
    • 1.8 主倉:切片填充與負索引的靜態邊界兩修,越界寫與靜默負地址關閉(#3352 / #3347)(10-27)
    • 1.9 主倉:T.ws(-1) 區域在所有執行緒上執行,落單寫法不再丟存(#3354)(10-27)
    • 1.10 主倉:profiler 未指定裝置時用當前裝置,快取沖刷不再打到 0 號卡(#3357)(10-27)
    • 1.11 新增缺陷:核心快取命中破壞六個假定冷編譯的測試;CuTeDSL 的 CUDA graph 捕獲為空圖(#3694 / #3693)(10-27)
  • 二、多後端適配(昇騰 / 海光 / ROCm / Hexagon)
    • 2.1 海光:HCU scale 緩衝按雙 bank 分割槽,blockscaled GEMM 在手動 warp 特化下共用一套行對映(#29)(10-27)
    • 2.2 昇騰主倉:PerCoreTask 候選體允許內部 if/else,條件 L1 填充不再被整體拒絕(#3686)(10-27)
    • 2.3 昇騰主倉:CI 把 AscendC 與 PTO 拆成兩個矩陣作業,可分別重跑(#3687)(10-27)
    • 2.4 昇騰倉在審:Cython 包裝器惰性構造分析器,直指昨日記錄的每次呼叫重建問題(#1948)(10-27)
    • 2.5 昇騰倉在審:GLM-5.3 的 kpool 壓縮運算元、PTO 上的張量移位量(#1910 / #1949)(10-27)
    • 2.6 昇騰主倉在審:PTO 的 T.Parallel 向量化、AutoSimVF 通道、批次三維 ND2NZ、寫計數多緩衝、Bisheng 構建身份進快取鍵(#3512 / #3320 / #3689 / #3629 / #3691)(10-27)
    • 2.7 ROCm 在審:pip 版 ROCm 10 的架構探測與映象、FP4 轉換與 E8M0 型別(#3549 / #3665)(10-27)
    • 2.8 昇騰倉日測:2903 項全透過,連續兩日無失敗(#1950)(10-28)
    • 2.9 沐曦:組織檢視有推送時間戳,六個分支均無視窗內新提交(10-27)
    • 2.10 第三方:Hexagon DSP 後端單日 15 筆,HVX 向量化、VTCM 競技場、Crouton 佈局與 DMA 內建齊備(10-27)
  • 三、生態與採用方(TileOPs / TileFoundry)
    • 3.1 TileOPs 0.1.0 釋出:206 個運算元的首個正式版,公開 API 就此定形(10-27)
    • 3.2 TileOPs 0.1.0 基準:H200 上 2265 行裡 2262 行不慢於最快實現,幾何平均 1.37 倍(10-27)
    • 3.3 TileOPs 0.1.0 的已知缺口與下一步:融合運算元、整模型、計算通訊融合(10-27)
    • 3.4 TileOPs 合入四條:where 與 masked_fill 可調優、D128 因果對角按 64 鍵打分、稀疏長 GDN 預填充拆分、變更日誌按釋出日定日期(#2782 / #2783 / #2779 / #2781)(10-27)
    • 3.5 TileFoundry 合入兩條:整視窗按運算元本身下發、動態共享記憶體按核心實際用量定尺寸(#273 / #271)(10-27)
    • 3.6 外部生態:CANN 社群的 TileLang 運算元倉、PPU 版零改動包裝器、DeepSeek V4.1 Flash 片段示例(10-27)
  • 四、社群、教程與活動
    • 4.1 文件站單日六筆自動更新(10-27)
    • 4.2 昇騰倉的文件事實核查四條,單條評論最多

46 輪(#1628 / #1655 / #1656 / #1659)(10-27)

  • 4.3 在審:把 lower trace 抽成獨立的開發工具包(#3583)(10-27)
  • 4.4 新聞側:Google News 與 HN 24 小時內無本次視窗的新報道(10-27)
  • 五、趨勢觀察
  • 附、素材與核查說明

今日重點

TileOPs 0.1.0 在今天視窗內落地,結束了連續多日的「卡在最後一格」狀態。PyPI 上的 0.1.0 製品上傳於 10-27 07:45(北京時間),比上一期日報的截點晚 45 分鐘;GitHub 釋出頁在 13:15 掛出,倉庫首次出現正式 tag。釋出說明的硬指標是三組:206 個運算元全部有 manifest 條目(簽名、測試與基準負載、roofline 公式,CI 裡逐條校驗);H200 上 2265 行夜間基準中 2262 行不慢於最快對照實現,幾何平均 1.37 倍,對照方包括 torch.compile、FlashAttention 3、FlashInfer、vLLM、DeepGEMM 與 FlagGems;206 個運算元裡有 195 筆帶 foundry 標籤的合入是由 TileFoundry 生成核心並透過同一套 CI 進庫的。釋出同時把公開 API 凍結,之後 0.x 只增不改,並新增 API 相容性檢查指令碼,任何按上一版寫法呼叫的 PR 會直接失敗。

24 小時內 tile-ai 組織合入 20 筆(tilelang 13、TileOPs 4、TileFoundry 2、tilelang-hygon 1),在審更新 45 筆,新開 issue 7 條(其中 2 條是昇騰倉的每日測試機器人)。合入重心分兩條線:一條是浮點語義與執行上下文,ROCm 把七個預設數學函式從快速內建改回精確實現(MI300X 上 sin/cos 誤差從 7.8e-04 降到 6.6e-08),自動目標改讀當前裝置而不是固定 0 號卡,NVRTC 與 CuTeDSL 改為按輸入張量所在裝置取流,自動調優的快取鍵補進 compile_flags;另一條是靜態邊界與執行緒域,切片填充按區域末端而不是按長度校驗、超範圍負索引在編譯期拒絕、T.ws(-1) 的區域改為在所有執行緒上執行。海光合入 12 檔案 1389 行,把 scale 緩衝拆成獨立雙 bank 並與 LDS bank 分開。第三方側 Hexagon DSP 後端單日 15 筆,把 HVX 向量化、VTCM 競技場、Crouton 佈局與 DMA 內建一路鋪到 Crouton 上的對映與歸約。


一、核心專案進展

1.1 主倉:ROCm 的 exp/sin/cos 不再預設降到快速內建,sin/cos 誤差從 7.8e-04 降到 6.6e-08(#3621)

合於 10-27 22:39,5 檔案 +107/-17,修 #3618。在 HIP 上,T.exp、T.exp10、T.log、T.log2、T.log10、T.sin、T.cos 七個預設函式無論 tl.enable_fast_math 如何設定都會降到 __expf 一類快速內建,而 CUDA 側按 #875 的約定走精確實現;反過來,顯式的 T.__exp 等快速運算元在 HIP 上根本不下沉,在 tilelang.rocm.language 裡直接拋 AttributeError。本筆讓 HIP 與 CUDA 採用同一策略:七個預設運算元改派 HIPMath,八個帶下劃線的快速運算元新增最低下沉規則指向 HIPFastMath,並統一挪進 FAST_MATH_INTRINSICS 列表由 ROCm 語言層匯出。精度對比取 #3618 的復現:MI300X 上 exp 與 exp10 看相對誤差、sin 與 cos 看絕對誤差(x 到 1e4),四列依次為 T.exp、T.exp10、T.sin、T.cos,主幹為 3.7e-06 / 4.1e-06 / 7.8e-04 / 7.8e-04,本筆為 8.3e-08 / 8.1e-08 / 6.6e-08 / 6.4e-08。

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

1.2 主倉:自動目標改讀當前 HIP 裝置,多卡機上不再固定取 0 號卡架構(#3594)

合於 10-27 22:37,2 檔案 +24/-1。determine_target(“auto”) 在 ROCm 上落到註冊的 HIP 探測器,而該探測器無條件讀 torch.cuda.get_device_properties(0),於是當前 HIP 裝置不是 0 號卡時,自動探測給出的是 0 號卡的架構;CUDA 側早就在用 torch.cuda.current_device(),ROCm 沒有跟上。修法是先取當前裝置索引再查屬性,並用 normalize_rocm_arch 去掉特性字尾。新增測試把 current_device() 打樁返回 1,同時記錄所有被探測過的索引,斷言探測到 gfx942 且只碰過 1 號卡;修復前該斷言得到 probed == [0]。作者註明本機無 ROCm 裝置,測試只用打樁 torch 驗證探測助手,未跑真實 HIP 執行時;與本日另一筆 #3549 改動的是不同函式(#3549 在 tilelang/contrib/rocm.py,本筆在 tilelang/rocm/target.py)。

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

1.3 主倉:TMA gather/scatter 的常量行索引轉為 int32,字面量寫法不再編譯崩(#3365)

合於 10-27 09:15,2 檔案 +44/-2,修 #2601。tma_gather4 與 tma_scatter4 兩個前端助手在收到 Python 整數形式的行索引時直接透傳,落到 TIR 裡沒有確定的整型寬度,於是在 Blackwell 上編譯崩潰;其他行表示式保持原樣。本筆只把 Python 整數轉成 int32 的 TIR 常量,並補一個受 Blackwell 門控的編譯迴歸,用字面量行索引分別編譯 gather 與 scatter,再檢查生成的 CUDA 原始碼裡確實出現對應指令。作者說明該測試只驗證能編譯,不驗證執行時數值正確;且提交者本機的 Windows Python 3.13 環境裡 Home Assistant 的 pytest 外掛依賴 fcntl,pytest 收集階段就失敗,本筆的 TileLang 測試在本地並未真正跑起來。

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

1.4 主倉:NVRTC 與 CuTeDSL 按首個輸入張量所在裝置取當前流(#3359)

合於 10-27 08:11,2 檔案 +5/-4,接 #3301 的 F149 與 #3408。兩個介面卡此前用字串字首判斷目標型別,且預設取呼叫方當前裝置的當前流,於是呼叫方當前 GPU 與張量所在 GPU 不一致時選錯流。本筆改用目標種類判斷,並按首個輸入張量的裝置查當前流,顯式傳入的流控制代碼仍然最高優先,無張量的 NVRTC 呼叫維持原語義。這是繼昨日 out_idx 輸出分配流(#3649)之後,流語義清理的第二筆。作者劃清邊界:本筆只修流選擇,跨裝置的預設流啟動上下文仍是未解項,這類呼叫方必須自己在啟動前把張量所在裝置設為當前裝置,補丁裡沒有逐呼叫的裝置守衛;驗證時用真實 Torch 流與 GPU 啟動覆蓋了同裝置與跨裝置的非預設流,但未建立跨裝置預設流上下文的正確性。

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

1.5 主倉:自動調優快取鍵納入 compile_flags,同一核心換編譯選項不再複用舊二進位制(#3338)

合於 10-27 16:57,3 檔案 +28/-1,接 #3301 的 F131。自動調優器把編譯結果持久化快取,但快取標識裡沒有 compile_flags,於是同一個核心換一組裝置編譯器選項後仍會命中舊二進位制,測出來的效能與真實編譯結果對不上。本筆讓 CompileArgs 真正存下並按原有順序參與雜湊,並把 flag 從 JIT 實現經 AutoTuner.set_compile_args() 一路透傳到候選編譯、逐配置的 pass 合併、分組裝置編譯與快取核心載入。驗證覆蓋三條既有快取鍵測試,以及用原生執行時對 CUDA 做的定點對比:列表與字串兩種 flag 形式都能在逐配置覆蓋中存活;作者指出原先的合併邏輯會把編譯器宏丟掉,正是這次要堵的口子。

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

1.6 主倉:ThreadSync 把 address_of 記為寫,不透明呼叫後的柵欄補上(#3346)

合於 10-27 05:49,2 檔案 +62/-1,接 #3301 的 F305,延續昨日的柵欄語義主線。address_of 本身不攜帶訪問資訊,它只把指標交給下游,但訪問分類器把它的型別硬編碼成讀,於是一塊共享緩衝經不透明呼叫寫入後不會進 writes 集合,後續的全塊讀找不到衝突寫、也不再規劃柵欄,執行緒就這樣搶過寫入讀到寫前內容。實測的 CUDA 輸出很直白:T.address_of(B[0]) 形式在 extern 之後沒有柵欄,而 T.access_ptr(B, “w”) 形式有;後者顯式宣告瞭訪問型別,所以這個洞一直沒被發現。loop_unswitching.cc 早已按寫保守處理 address_of,ThreadSync 是唯一的例外。本筆把 address_of 記為寫,方向是保守的:指標用途未知就可能被寫,而柵欄規劃只需在安全側錯一次;新增兩個用例,斷言該核心下沉後 extern 之後跟著 __syncthreads()。

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

1.7 主倉:執行時行程數的流水線 epilogue 直接展開,分頁解碼恢復可編譯(#3356)

合於 10-27 07:02,1 檔案 +5/-1。T.Pipelined 迴圈的行程數是執行期值時,只要 epilogue 展開兩步以上且帶非同步複製,CUDA codegen 就報 TypeError: Downcast from tirx.Sub to ir.IntImm failed。根因在 EmitImpl:prologue 與 epilogue 的 extent 是常量時才展開,以保證每次迭代的等待計數是常量;而 epilogue 的 extent 是 end - start,即 (n + max_stage) - n,n 為常量時在構造期摺疊,n 為執行期時保持為 Sub,as_const_int 返回空,於是 epilogue 被當成 T.unroll 迴圈發射,其等待寫成 T.ptx_wait_group(1 - k) 這樣讀迴圈變數的形式,而 ptx_wait_group 只收立即數。修法是判斷前先看化簡後的 extent,常量邊界的情形化簡回自身,與原來一致。該缺陷由下游 TileOPs#1852 報出並二分到 afcebed1 之後,破壞物件是分頁解碼——其行程數來自 KV 長度。作者在 sm_89、CUDA 13.0 上雙版編譯對比,新增測試在主幹失敗、在本筆透過。

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

1.8 主倉:切片填充與負索引的靜態邊界兩修,越界寫與靜默負地址關閉(#3352 / #3347)

合於 10-27 05:37 與 05:14,合計 4 檔案 +100/-3。前者針對 Fill 構造器的靜態邊界檢查:原來逐維只校驗 min >= 0 與 extent <= shape,而 extent <= shape 只對從緩衝起點開始的區域成立;一個起點靠裡的切片兩項都透過,卻被下沉成越界寫——T.fill(As[6:12], 7) 在長度 8 的緩衝上編譯透過,發射出 As[threadIdx.x + 6] = 7,寫到 6 到 11。修法是把區域末端也對照目標形狀校驗,且新條件要求起點與長度都是靜態,錨定在零的區域與符號起點區域都不受影響;T.clear 共用同一節點,缺陷相同也一併關閉。後者針對 LegalizeNegativeIndex:原來把負索引加一次緩衝長度就了事,沒有範圍檢查,長度 8 的緩衝上 A[-9] 得到 -1、A[-10] 得到 -2,跳出分配起點且無任何診斷,而 NumPy 與 PyTorch 對同樣的索引拋 IndexError。修法是僅在分析器能證明迴繞後仍為負時拒絕,並報出索引與長度,所有在範圍內的常量與全部動態索引維持原行為。

來源:https://github.com/tile-ai/tilelang/pull/3352、https://github.com/tile-ai/tilelang/pull/3347

1.9 主倉:T.ws(-1) 區域在所有執行緒上執行,落單寫法不再丟存(#3354)

合於 10-27 04:52,2 檔案 +40/-3,接 #2978。WarpSpecialize 的文件把 -1 定義為「跳過 threadIdx.x 繫結」,也就是該區域不受 threadIdx 約束;實現裡卻把 -1 當普通 warp 組索引,構造出 [-warp_group_size, 0) 這個對任何執行緒都不成立的區間,於是區域在零個執行緒上執行,區間內所有存全部被丟棄且無診斷。最小復現很乾淨:256 執行緒裡先全寫 0,再在 T.ws(-1) 裡寫 5,結果 256 個執行緒寫出的仍是 0;同樣的存不加 ws 包裝則 256 個執行緒都寫 5。修法是當 id 裡含 -1 時不發射任何範圍約束,真實組索引仍按自身執行緒約束,並在文件串裡補上 -1 這一例——原文件的示例恰好缺它。作者核查全倉未出現 ws(-1) 用法,因此沒有既有程式碼依賴那個空區域。

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

1.10 主倉:profiler 未指定裝置時用當前裝置,快取沖刷不再打到 0 號卡(#3357)

合於 10-27 07:26,1 檔案 +1/-1。do_bench 在未顯式指定裝置而當前 CUDA 裝置非零時,_cache_device(None) 落到模組級預設的 cuda:0,於是為沖刷快取而做的分配可能落在與實測工作不同的 GPU 上,測出來的數字與真實負載不在同一張卡。修法是預設串改成不帶索引的 “cuda”,由 PyTorch 在分配緩衝時解析當前裝置;顯式指定裝置與 NPU、MPS 的預設值保持不變。上游 CI(run 37915443170)在 Quick Lint、CUDA、ROCm、Metal、Ascend 與 CuTeDSL 示例上透過;作者註明後續的 PR 觸發式執行需維護者批准,尚未執行測試,本機也沒有 CUDA 硬體或計時基準。

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

1.11 新增缺陷:核心快取命中破壞六個假定冷編譯的測試;CuTeDSL 的 CUDA graph 捕獲為空圖(#3694 / #3693)

兩條 issue 均由 lucifer1004 在 10-27 開出。前者指出有六個測試在冷快取下透過、用同一快取目錄再跑一次即失敗,原因是它們假定編譯一定會發生:例如 test_tilelang_kernel_gemm_sm75.py 裡的 sm75 邊界用例在四個引數下報 AttributeError: ‘NoneType’ object …;後者更硬:從 CuTeDSL 後端發出的核心啟動沒有被記進 CUDA graph,torch.cuda.CUDAGraph 捕獲得到空圖,replay() 什麼都不跑,而同一核心用 cuda 後端編譯則捕獲回放正常,報告附了最小復現。兩條都是「快取與執行上下文」這一族缺陷的延伸,與 1.4、1.5 兩筆合入指向同一片區域。

來源:https://github.com/tile-ai/tilelang/issues/3694、https://github.com/tile-ai/tilelang/issues/3693


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

2.1 海光:HCU scale 緩衝按雙 bank 分割槽,blockscaled GEMM 在手動 warp 特化下共用一套行對映(#29)

合於 10-27 08:26,12 檔案 +1389/-314,是本視窗軟體側最大的一筆。改動把 scale 緩衝的行規劃成獨立的雙 bank 分割槽,與 LDS bank 分開,並把同一套行對映同時交給分配、scale 複製與 MMAC 下沉三處使用,避免三處各自推一遍佈局。細節包括:保留物理區間化的生命週期與 m0 對齊分頁;壓縮 scale 複製的下發條數;把空閒 wave 的工作攤到兩個 scale bank 上;該佈局用於手動 warp 特化下的 blockscaled GEMM。這一筆接上昨日合入的原生 B16 儲存與動態 WASP GEMM(#28),海光在兩天裡完成了三筆主幹合入。

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

2.2 昇騰主倉:PerCoreTask 候選體允許內部 if/else,條件 L1 填充不再被整體拒絕(#3686)

合於 10-27 07:24,2 檔案 +85/-24。T.PerCoreTask 裡的 T.Task 只要體內含 if/else 就被拒絕,哪怕分支只控制內部工作——例如一次複製之後按條件補 L1 填充。復現於 d85a70a9,物化階段報「選擇候選的控制流必須位於 task 作用域之外」。本筆刪掉這條一刀切的拒絕,改為對 task 體做統一校驗:同步柵欄包住整個候選,同時每個分支仍逐項檢查管線與核的相容性、巢狀 task、以及不受支援的跨核同步。迴歸測試接受內部 if 與 if/else、保留候選體、並保留原有的拒絕用例,另檢查完成標誌在外層守衛下跟隨整個候選。task、排程單元與延遲三套共 113 例全透過,帶條件 L1 填充的核心也能完成完整昇騰原始碼下沉;裝置執行未測。

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

2.3 昇騰主倉:CI 把 AscendC 與 PTO 拆成兩個矩陣作業,可分別重跑(#3687)

合於 10-27 07:05,7 檔案 +56/-27。此前 AscendC 與 PTO 是同一個 Test for Ascend 作業裡的兩個步驟,任一後端專屬失敗都會讓整條檢查變紅,且重跑必然連帶另一套示例。本筆把示例與單元測試拆成兩個矩陣作業:Test for AscendC 跑 examples/ascend(target 為 ascend,4 併發)與 testing/ascend -m “not pto”(32 併發),Test for PTO 跑 examples/ascend(target 為 pto)與 testing/ascend -m pto(32 併發)。劃分靠既有的 pto 標記,並給只跑 PTO 的測試與雙後端測試裡的 PTO 引數補上標記;兩個標記表示式互斥且合起來覆蓋全套,不重複執行。fail-fast 置為 false,一個作業失敗不取消另一個;兩邊都保留「示例失敗但構建成功時仍跑單測」的行為,各自發布構建與測試摘要。共用測試(含同一前端 IR 分別過兩個程式碼生成器的契約檢查)只在 AscendC 一側跑一次,約定寫進 testing/ascend/README.md。

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

2.4 昇騰倉在審:Cython 包裝器惰性構造分析器,直指昨日記錄的每次呼叫重建問題(#1948)

更新於 10-27 07:42,作者 zwh1025。昨日日報記錄的新缺陷 #1945 指出,昇騰倉的 Cython 包裝器每次呼叫都新建一個 TVM 分析器,實測約 0.28 毫秒的固定開銷;本筆以「惰性構造 TVM 分析器」為題,正對這個開銷。這類小開銷在大模型推理的逐層呼叫裡會累積成可觀測的延遲,屬於昇騰側效能清掃的低垂果實。目前仍在審。

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

2.5 昇騰倉在審:GLM-5.3 的 kpool 壓縮運算元、PTO 上的張量移位量(#1910 / #1949)

兩條均在 10-27 更新。#1910 給昇騰 NPU 加 GLM-5.3 的 kpool 壓縮運算元示例——這是把具體模型的結構件搬進示例倉的做法,與 TileOPs 側「運算元按模型需求長」的思路一致;#1949 讓 T.tile.bitwise_lshift 在 PTO 上支援張量形式的移位量,此前只接受常量。PTO 是昇騰側相對獨立的一條程式碼生成路徑,移位量張量化意味著該路徑在掩碼與位打包類運算元上向通用表達靠了一步。

來源:https://github.com/tile-ai/tilelang-ascend/pull/1910、https://github.com/tile-ai/tilelang-ascend/pull/1949

2.6 昇騰主倉在審:PTO 的 T.Parallel 向量化、AutoSimVF 通道、批次三維 ND2NZ、寫計數多緩衝、Bisheng 構建身份進快取鍵(#3512 / #3320 / #3689 / #3629 / #3691)

五條在審構成昇騰側本視窗的排隊佇列,均於 10-27 有更新。#3512 是 PTO 上把 T.Parallel 向量化的流水線框架與逐元素支援,屬結構性改動;#3320 新增 AutoSimtVF 通道並註冊進昇騰流水線,方向上是在 SIMT 與向量之間自動擇路;#3689 支援批次三維 GM 到 L1 的 ND2NZ 複製,直接對應批次注意力與分塊 GEMM 的搬運需求;#3629 增加寫計數式多緩衝並重排 DeepGEMM 的載入;#3691 把 Bisheng 編譯器構建身份納入二進位制快取鍵,與 1.5 的 compile_flags 進快取鍵是同一類思路——昇騰側同樣存在換編譯器後誤命中舊二進位制的問題。

來源:https://github.com/tile-ai/tilelang/pull/3512、https://github.com/tile-ai/tilelang/pull/3320、https://github.com/tile-ai/tilelang/pull/3689、https://github.com/tile-ai/tilelang/pull/3629、https://github.com/tile-ai/tilelang/pull/3691

2.7 ROCm 在審:pip 版 ROCm 10 的架構探測與映象、FP4 轉換與 E8M0 型別(#3549 / #3665)

#3549 支援 pip 安裝的 ROCm 10 SDK,改動含架構探測與 Dockerfile;這與 1.2 的 auto 目標修正相關但落在不同函式——本筆動的是基於 /opt/rocm 的 get_rocm_arch 探測,1.2 動的是 tilelang/rocm/target.py 裡的探測助手。#3665 修 HIP 模板裡的浮點到 FP4 轉換,並補 E8M0 型別支援,指向分塊縮放與低精度量化在 ROCm 上的可用性。兩條均在審。

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

2.8 昇騰倉日測:2903 項全透過,連續兩日無失敗(#1950)

昇騰倉的每日測試機器人 10-28 06:49 的報告顯示總計 2903 項、透過 2903、透過率 100%,工作流執行號 40590261193;前一日的報告(#1947)同樣是 2903 項全透過。對一條仍在快速演進的 NPU 後端來說,兩日無失敗的日測說明主幹可用性穩定,這也為 2.3 的 CI 拆分提供了基線——拆分後兩個作業各自仍應維持這個數字。

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

2.9 沐曦:組織檢視有推送時間戳,六個分支均無視窗內新提交(10-27)

tilelang-metax 倉的 pushed_at 顯示 10-26 07:08,落在本視窗起點之前;逐分支核查六個分支,視窗內沒有新提交,按無實質變更處理,不列條目級內容。

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

2.10 第三方:Hexagon DSP 後端單日 15 筆,HVX 向量化、VTCM 競技場、Crouton 佈局與 DMA 內建齊備(10-27)

xwhzz/tilelang-hexagon 的 hexagon-backend 分支在 10-27 一天落了 15 筆提交,是本視窗後端側最密集的外部動作。當日主線可以分成三段:底層從 CodeGenTileLangCBase 重建程式碼生成器,把 DMA 做成帶提交組的 IR 內建,把啟動網格與工作執行緒池下沉進 TIR,並在 TIR 裡規劃 VTCM 競技場;中層鋪 HVX 的單條向量化通路(算術對映、向量數學加一條逐元素路徑)、用 HVX 轉置樹做多行歸約、軟體流水化 HVX 對映,並支援片段與 T.reduce_*;上層引入 Crouton 佈局與區域複製、讓 HVX 對映與歸約直接在 Crouton 瓦片上跑、每個會話對 DSP 頻率投票,最後補構建快取與並行裝置測試。提交裡帶一個破壞性標記(feat(hexagon)!),說明 Crouton 佈局與區域複製改動了既有介面。Hexagon 是手機與嵌入式側的 DSP,這條線若走通,TileLang 的硬體覆蓋面就從 GPU、NPU 延伸到 DSP。

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


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

3.1 TileOPs 0.1.0 釋出:206 個運算元的首個正式版,公開 API 就此定形(10-27)

10-27 13:15 掛出 v0.1.0,是倉庫首個 GitHub release;PyPI 上的 0.1.0 製品上傳更早,為 10-27 07:45(北京時間,即 10-26 23:45 UTC),比上一期日報的截點晚 45 分鐘——這正是昨日日報寫「首發仍在最後一格」之後立刻發生的變化。釋出說明定位很明確:一個為 agent 時代重做的運算元庫,每個運算元先寫 YAML 規格(宣告簽名、測試與基準用的負載、roofline 公式),由編碼 agent 開發與調優核心,再由規格派生出的測試與基準決定誰能進庫;195 筆帶 foundry 標籤的合入就是這樣進來的。本版固定公開 API:運算元按 tileops.. 匯入,docs/stability.md 列出公開模組,scripts/ci/check_api_compat.py 讓任何破壞上一版寫法的 PR 直接失敗;後續 0.x 只增不改。另有幾項能力一併進版:torch.compile(fullgraph=True) 對所有 manifest 條目含呼叫期張量輸入且無組合的運算元可用;H200 上驗證了圖測試組內運算元的 CUDA Graph 捕獲;op.request_tune() 後自動調優,編譯產物與調優結果快取在 TILELANG_CACHE_DIR 下跨程序複用;tileops.backend 允許後端分發版透過 register_kernel_builder 把自己裝置上的整運算元接進來;tileops.trace 記錄核心內時間線。釋出致謝指出 @FeathBow 一人貢獻 98 筆合入。

來源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0、https://pypi.org/project/tileops/

3.2 TileOPs 0.1.0 基準:H200 上 2265 行裡 2262 行不慢於最快實現,幾何平均 1.37 倍(10-27)

夜間基準把每個 TileOPs 核心與同一運算元的其他所有實現放在同一張 GPU 上比,單行比值取最快對照方耗時除以 TileOPs 耗時,大於 1 表示更快。全量 2265 行、16 個運算元族,幾何平均 1.37 倍,2262 行達到 0.99 以上。按族看差異明顯:sequence_modeling 4.00 倍、sampling 2.95 倍、convolution 1.86 倍、mamba 1.61 倍、linear_attention 1.54 倍、pool 1.51 倍、quantization 1.45 倍、attention 1.42 倍;體量最大的 elementwise 527 行只有 1.09 倍,gemm 235 行 1.28 倍。對照方包括 torch、torch.compile、FlashAttention 2 與 3、FlashInfer、FlashMLA、vLLM、DeepGEMM、flash-linear-attention、FlagGems、QuACK、Mamba、Marlin 與 cuBLASLt。落後 0.99 的三行分別是 16M float32 元素上的 Where 與 ClampScalar、以及 DeepSeek-V3 預填充輸出投影上的 FP8 GEMM,差距 1.1% 到 1.7%。這個分佈說明收益主要來自長尾運算元族而非主流 GEMM 與逐元素——後者本來就是各路實現卷得最狠的地方。

來源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.3 TileOPs 0.1.0 的已知缺口與下一步:融合運算元、整模型、計算通訊融合(10-27)

釋出說明列了四條已知問題:FusedAddLayerNormFwdOp 在遠離零的行上丟精度、在接近 1e36 的行上溢位(#2773,昨日日報已記錄該缺陷);同一主機上若有不同架構的 GPU,為一個 GPU 編譯的核心可能跑到另一個上(#2302);Bwd 類運算元需直接呼叫,不接入 torch.autograd;單個運算元例項同時只支援一個執行緒與一條 CUDA 流。驗證範圍只覆蓋 H200、Python 3.12 與 3.13;部分核心另有 SM80、SM86、SM89 路徑,這些路徑在對應裝置上跑但不屬於釋出驗證範圍。路線圖三步:先做跨運算元邊界的複雜融合(注意力連同前後投影與歸一化、整層 MoE、GEMM epilogue 內的量化),再做整模型的預填充與解碼並與服務引擎端到端對比,最後做計算與通訊融合(張量並行 GEMM 配 all-reduce、專家並行 MoE 配 all-to-all)。這三步若落地,TileOPs 的對手就從單運算元庫變成推理框架。

來源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0

3.4 TileOPs 合入四條:where 與 masked_fill 可調優、D128 因果對角按 64 鍵打分、稀疏長 GDN 預填充拆分、變更日誌按釋出日定日期(#2782 / #2783 / #2779 / #2781)

四筆在 10-27 合入,恰是釋出前的收尾。#2782 修 request_tune() 在 WhereFwdOp、MaskedFillTensorFwdOp、MaskedFillScalarFwdOp 上拋 ValueError 的問題:自動調優拒絕隨機生成的整型輸入,而這三個核心把 bool 條件或掩碼當 uint8 傳;隨機選擇器只改變每個元素取哪個運算元、不改變候選的工作量,故宣告隨機整型輸入安全;新增三個測試節點,主幹上三例全失敗。#2783 最佳化 D128 稠密程式的因果對角:當因果呼叫只有一段 key 且兩個長度都在 128 行邊界上時,第一個消費 warpgroup 的對角 key 瓦片對全部行都遮蔽了高 64 鍵,於是這一步改按 m64n64 的 QK 配 64 鍵 softmax 跑;該選擇是編譯期標誌(diag_aligned,進核心鍵),因為執行期在 64 與 128 寬 QK 之間分支會讓 ptxas 在某些特化上序列化每條 WGMMA(C7511),改用標誌後所有 manifest 裡的 D128 預填充特化都不再出 C75xx。#2779 把 SM89 上稀疏長序列的 GDN 預填充拆成狀態掃描加獨立輸出程式,適用於稠密、key 主序的 D128 輸入且至少 16384 token、狀態掃描 CTA 不超過 SM 數四分之一的情形;基準裡 (1, 16384, 4, 128) bf16 的耗時從約 0.612 毫秒降到約 0.470 毫秒,與 flash-linear-attention 的 0.441 毫秒相比為 0.939 倍。#2781 把 CHANGELOG 裡 0.1.0 的日期定為它進 PyPI 的那天(2026-10-27),並刪掉貢獻者文件裡的釋出一次配置與 Actions 點選路徑。

來源:https://github.com/tile-ai/TileOPs/pull/2782、https://github.com/tile-ai/TileOPs/pull/2783、https://github.com/tile-ai/TileOPs/pull/2779、https://github.com/tile-ai/TileOPs/pull/2781

3.5 TileFoundry 合入兩條:整視窗按運算元本身下發、動態共享記憶體按核心實際用量定尺寸(#273 / #271)

兩筆均在 10-27 合入。#273 修 sm_89 GEMM 的下沉:schedule finalize 寫出的內容編不過,因為下沉把整個累加器包進一個從零開始、跨度等於自身的視窗 T.ptr_of(acc[0:0+16, 0:0+8]),而 CUDA codegen 拒絕分片暫存器瓦片的指標切片,報「要求普通源佈局」;MMA 又把暫存器片段讀成 lhs_frag_dst[0:0+16, 0:0+32] 這樣的裸 Slice,codegen 同樣無處理器。修法在下沉側:StoragePlan.operand 本來就把「從零開始且跨滿運算元的視窗」當作運算元本身,但物化與指令下發仍構造 Slice;改為只判定一次——對 MLIR 的 TrivialSubViewOpFolder 做同樣的替換。#271 修動態共享記憶體:合成的宿主入口給每個 TIR 核心都用 dynamic_smem=0 啟動,經整型 TensorView 偏移訪問共享記憶體的核心因此沒有分到空間,超過 48 KiB 還需設定 cudaFuncAttributeMaxDynamicSharedMemorySize 而啟動器從不設定;作者沒有采納 #267 裡「按目標容量給滿」的建議,改為按核心實際用量定尺寸,理由是靜態共享記憶體與動態共享記憶體同計一個上限、且每塊申請多少決定一個 SM 能放幾塊——sm_89 上要滿 99 KiB 會讓用 512 位元組的核心也只放得下一塊。CUTLASS 與 TileLang 都按同樣方式定尺寸。

來源:https://github.com/tile-ai/TileFoundry/pull/273、https://github.com/tile-ai/TileFoundry/pull/271

3.6 外部生態:CANN 社群的 TileLang 運算元倉、PPU 版零改動包裝器、DeepSeek V4.1 Flash 片段示例(10-27)

三個外部倉在視窗內都有動作,合起來看是一條「昇騰側運算元資產」的成形過程。Atlas-Ascend-Fork/ops-tilelang 是 CANN 社群面向昇騰的 TileLang 運算元倉,10-20 建倉、10-27 提交「自動發現增量測試並移除運算元對映表」;README 明確只支援 Ascend 950 及以上架構、要求 CANN 9.3.0 以上並用社群 weekly 版,設計原則是介面與排程分離、按業務域聚合、正確性先於效能、benchmark 與測試分離且基線按測試檔案與晶片叢集分片儲存。kevin3113/models-tilelang 用這個倉的七個運算元串出 DeepSeek V4.1 Flash 的「Attention 輸出加 FFN 前處理」片段(mhc_post_forward、帶平方和 GEMM、歸約加 RMSNorm、mix 拆分、sinkhorn、apply_mix、ffn_norm),HIDDEN=5120、MHC=4、預設 decode 形態 SEQ_LEN=1,支援 SuperKernel 與 npugraph_ex 的 ACL graph,B 從 1 到 256 實測 SuperKernel 收益 1.17% 到 2.75%。flytiger-eco/tilelang-wrapper-for-sail 解決另一個方向的遷移:把 CUDA 風格的 TileLang 程式碼零改動跑在 tilelang-sail(PPU 後端分發版)上,做法是純執行時攔截 import 協議與 site 機制,不修改上游原始碼也不重新打包,無 PPU 對應物的 CUDA 特性顯式報錯而非靜默降級。三者分別對應運算元供給、模型落地、後端遷移三個環節。

來源:https://github.com/Atlas-Ascend-Fork/ops-tilelang、https://github.com/kevin3113/models-tilelang、https://github.com/flytiger-eco/tilelang-wrapper-for-sail


四、社群、教程與活動

4.1 文件站單日六筆自動更新(10-27)

tilelang.github.io 在視窗內有六次 Update docs 提交(05:20、06:22、07:19、09:28、22:42、22:43),均由機器人推送,內容來自主倉合入後的文件重生成。釋出當天文件站更新頻次與主幹合入節奏同步,屬正常聯動,不單列內容。

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

4.2 昇騰倉的文件事實核查四條,單條評論最多 46 輪(#1628 / #1655 / #1656 / #1659)

昇騰倉在 10-27 有四條以「事實核查」為題的文件 PR 更新:T.tile.clear(#1628)、T.tile.bitwise_or(#1655)、T.tile.bitwise_not(#1656)、T.tile.bitwise_lshift(#1659),評論數分別為 10、46、18、20,其中三條帶機器人標記。一輪核查能滾到 46 條評論,說明討論的不只是措辭——這類逐 API 對行為的核對,實際是在為昇騰側的語義約定留檔。對一條新後端線來說,文件與實現對齊的成本往往被低估,這組 PR 是可見的投入。

來源:https://github.com/tile-ai/tilelang-ascend/pull/1655、https://github.com/tile-ai/tilelang-ascend/pull/1656

4.3 在審:把 lower trace 抽成獨立的開發工具包(#3583)

10-27 更新的這條在審 PR 把 lower trace 從主倉抽到 tilelang-devkit。下沉過程的追蹤資訊對除錯後端適配很有用,但放在主倉裡會增加編譯路徑與依賴負擔;抽成獨立工具包後,主倉可以保持輕量,需要時再掛上。這類「把除錯能力外接」的改動通常出現在專案開始有多個下游後端的時候——本視窗內 Hexagon、PPU、昇騰三條線同時在動,正是這種時點。

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

4.4 新聞側:Google News 與 HN 24 小時內無本次視窗的新報道(10-27)

中英文四組查詢(TileLang、TileOPs、TileLang 運算元、tile-ai 智算 kernel)返回的條目全部落在 9 月底到 10 月中:一批是 DeepSeek 與華為在 9 月 29 至 30 日開源昇騰版 TileLang 的報道(SCMP、Reuters、Tom’s Hardware、雷峰網、InfoQ 等),最新一條是 10 月 22 日 SemiAnalysis 關於華為與 DeepSeek 共生關係的分析,均在視窗之外。TileOPs 精確查詢零命中。Hacker News 視窗內只有一條與 GameMaker 重實現相關的條目,按噪音剔除。arXiv 上 TileLang 關鍵字最近一篇為 2026-07-24。今日的內容重心全在 GitHub 側,這與釋出日的時間點吻合——正式釋出尚未進入媒體報道週期。

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


五、趨勢觀察

第一,TileOPs 0.1.0 把「運算元庫」這件事的成本結構換了一種演算法。傳統做法是運算元數量隨硬體與模型需求線性增長,人力是瓶頸;這裡的做法是每個運算元先寫一份 YAML 規格,把簽名、負載與 roofline 宣告清楚,再由 agent 生成與調優核心,最後用規格派生的測試與基準做准入。206 個運算元、195 筆 foundry 標籤合入是這個流程的第一次全量交付。它能否持續,取決於兩件事:規格能否覆蓋複雜融合運算元(路線圖第一步正是這個),以及基準能否擋住「調優到基準上但實際不魯棒」的核心——FusedAddLayerNormFwdOp 在遠離零的行上丟精度這條已知問題,恰好是這類風險的一個樣本。

第二,本視窗的合入幾乎全是「靜默錯誤」。ROCm 的數學函式預設走快速內建會讓 sin/cos 差出三個數量級;address_of 被當成讀會讓不透明呼叫後的共享記憶體讀到寫前內容;T.ws(-1) 的區域在零個執行緒上執行,存全部丟棄且不報錯;切片填充按長度而非末端校驗,直接越界寫;負索引迴繞後仍為負,落到分配起點之前。這五條的共同點是編譯透過、執行不報錯、結果錯,且都在最近兩週內被集中發現。對一門要給大模型寫核心的語言,這類缺陷的代價遠高於編譯期報錯,它們指向同一個工程需求:把「不報錯但行為不對」的寫法在編譯期拒掉。本視窗的邊界校驗兩修與 ws(-1) 修正在做的就是這件事。

第三,後端覆蓋正在從「廠商主導」轉向「社群自發」。官方維護的昇騰、海光、沐曦、摩爾執行緒四條線之外,Hexagon DSP 後端在一天裡落 15 筆並把線性代數原語鋪到 Crouton 瓦片上,PPU 版出現零改動包裝器,CANN 社群自建運算元倉並用它串出 DeepSeek V4.1 Flash 的模型片段。這些動作都不在 tile-ai 組織的倉庫裡,卻都依賴 TileLang 的語言層穩定性。這也是 0.1.0 凍結公開 API、並把 API 相容性檢查做成 CI 門禁的意義所在——下游越多,向後相容的價值越高。

第四,快取身份這件事被兩個後端同時盯上。主倉把 compile_flags 納入自動調優快取鍵(#3338),昇騰把 Bisheng 構建身份納入二進位制快取鍵(#3691)。同一個核心在不同編譯選項或不同編譯器下產生的二進位制不同,快取若只按原始碼與形狀做鍵就會誤命中,測出來的效能與真實編譯結果對不上。兩條獨立路線在同一天做同一件事,說明這類問題在實際使用中已經足夠痛。


附、素材與核查說明

信源核查表

信源 核查結果
GitHub tile-ai 組織(29 倉) 視窗內推送 7 倉:tilelang、TileOPs、TileOPs-nightly、TileOPs.github.io、tilelang-hygon、TileFoundry、tilelang.github.io
合入真值(is:merged) 20 筆,逐筆核對合入時間落在視窗內(tilelang 13、TileOPs 4、TileFoundry 2、tilelang-hygon 1)
在審真值(is:open 更新) 45 筆,按更新時間排序全量核對
新增 issue 7 條,逐條讀正文,其中 2 條為昇騰倉每日測試機器人報告
釋出頁與 PyPI TileOPs v0.1.0 於 10-27 13:15 掛出,PyPI 製品上傳 10-27 07:45;tilelang 最新仍為 0.1.15(09-30),tilefoundry 最新 0.0.4rc1(10-25)
tilelang-ascend 倉 視窗內無提交,但 PR 與每日測試 issue 活動密集
tilelang-metax 倉 組織檢視有推送時間戳,六個分支均無視窗內提交,按無實質變更處理
tilelang-musa 倉 上次推送 09-30,視窗內無提交
採用方(TileKernels / FlashQLA) 均無視窗內推送(10-14)
外部後端倉 Hexagon 後端 15 筆提交;CANN 運算元倉、PPU 包裝器、模型片段倉均有視窗內動作,見 2.10 與 3.6
Google News 中英文多組查詢 精確查詢零命中,返回條目全部落在 9 月底至 10 月中,均在視窗外
Hacker News 視窗內一條 GameMaker 重實現條目,與 TileLang 無關,剔除
arXiv 關鍵字檢索無視窗內新論文,最近一篇為 2026-07-24

完整信源清單

tilelang Bisheng 構建身份進快取鍵(#3691)— https://github.com/tile-ai/tilelang/pull/3691