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

信源:GitHub tile-ai 組織逐倉核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 文件站 / 語言服務倉 / 夜間基準倉)+ 合入與在審真值檢索(合入 30 筆、在審更新 49 筆)+ 新增 issue 全量核對(11 條)+ 提交級 diff 與釋出頁核查 + PyPI 製品狀態核查 + 採用方倉核查(TileKernels / FlashQLA)+ 社群新倉檢索 + Google News 中英文多組查詢(走代理)+ Hacker News + arXiv + 產業媒體報道。


本期索引

  • 今日重點:24 小時合入 30 筆、在審更新 49 筆、新開 issue 11 條;主線圍繞「流與屏障語義」三處靜默錯;TileOPs 首發仍卡在 PyPI 最後一格;海光兩筆大合入
  • 一、核心專案進展
    • 1.1 主倉:out_idx 輸出改在呼叫方流上分配,H100 復現的五次三次損壞歸零(#3649)(10-26)
    • 1.2 主倉:NVRTC 裝置端編譯接進 TVM-FFI 執行路徑,新增編譯鍵項(#3285)(10-26)
    • 1.3 主倉:批次 TMA 載入的 cache hint 少一個模板引數,DeepSelect 的 TMA 路徑編譯不過(#3683)(10-26)
    • 1.4 主倉:歸約的 FP8/FP4 非有限常量寫出非法字面量,亞位元組累加器改為早拒(#3189)(10-26)
    • 1.5 主倉:SM120 分塊縮放 GEMM 顯式展開 K 原子,FullCol 配置提升 51%(#3286)(10-26)
    • 1.6 主倉:JIT 輸出形狀裡的標量引數在分配前解析(#3602)(10-26)
    • 1.7 主倉:sm8x 稀疏後設資料的運算元 dtype 不再無兜底落入(#3335)(10-26)
    • 1.8 主倉:DeepSelect Top-K 示例合入,129 組配置裡相對 torch.topk 最高 8.00 倍(#3304)(10-26)
    • 1.9 在審:TIRx 重構規模再擴到 1194 檔案,錨定未打補丁的上游版本(#3666)(10-26)
    • 1.10 新增缺陷:ThreadSync 三類錯位柵欄,配三筆在審修復(#3680 / #3682 / #3681)(10-26)
    • 1.11 新增缺陷:sm_100 以上歸約 batch 與佈局衝突,已有一筆在審修復(#3667 / #3684)(10-26)
    • 1.12 在審:轉置目標形狀校驗、區域化 kill 重排(#3677 / #3670)(10-26)
    • 1.13 已關閉:模糊測試系列的六條編譯期崩潰(10-26)
  • 二、多後端適配(昇騰 / 海光 / 沐曦)
    • 2.1 昇騰主倉:帶版本片上緩衝的精確定址介面(#3607)(10-26)
    • 2.2 昇騰在審:手工排程的 FlashMLA,BF16 預填充與 FP8 解碼(#3678)(10-26)
    • 2.3 昇騰在審:片上記憶體溢位提前拒絕、E4M3 的 UB 到 L1 複製、檢視佈局與指標單位(#3655 / #3639 / #3623)(10-24~26)
    • 2.4 昇騰在審:批次排程可選、跨核同步要求批模式、L0C 雙複製直轉(#3675 / #3614 / #3468)(10-26)
    • 2.5 昇騰倉:安裝器的 SHMEM 段雙修,空殼 wheel 不再假成功(#1943 / #1946)(10-26)
    • 2.6 昇騰倉:W4A4 的 int4 矩陣乘結果錯誤,根因是半位元組打包被當成 1 位元組(#1731)(10-26)
    • 2.7 昇騰倉:超過兩維的 GM 到 UB 複製按 DMA 行數摺疊前導維(#1805)(10-26)
    • 2.8 昇騰新缺陷:Cython 包裝器每次呼叫都造一個分析器,約 0.28 毫秒(#1945)(10-26)
    • 2.9 海光:暫存器流水線與 warp 發散合入,28 檔案(#18)(10-26)
    • 2.10 海光:原生 B16 儲存與動態 WASP GEMM 合入,scale 緩衝雙 bank 在審(#28 / #29)(10-26)
    • 2.11 沐曦:組織檢視有推送記錄,六個分支均無視窗內新提交(10-26)
  • 三、生態與採用方(TileOPs / TileFoundry)
    • 3.1 TileOPs:首發仍在最後一格,製品仍是 2025-10-09 的開發版(#2738)(10-26)
    • 3.2 TileOPs:D128 稠密前向拆 key,prefill-rope 從 31.04 微秒降到 17.09(#2777)(10-26)
    • 3.3 TileOPs:D128 持久化前向按執行緒重寫並寫出 LSE(#2774)(10-26)
    • 3.4 TileOPs:融合加法的 LayerNorm 行跨 CTA 拆分,精度缺陷另開新單(#2768 / #2773)(10-26)
    • 3.5 TileOPs:逐元素與歸約四條,clamp 的 bf16 行提速 5.9%(#2766 / #2767 / #2769 / #2770)(10-26)
    • 3.6 TileOPs:呼叫張量的對齊判斷上移到運算元層,三筆收尾(#2772 / #2778 / #2780)(10-26)
    • 3.7 TileOPs:偏移儲存三類修復,SM90 分組 GEMM 與輸出 ABI(#2761 / #2765 / #2762)(10-26)
    • 3.8 TileOPs 在審:FP8 查詢的分頁預填充改按 e4m3 乘 value(#2775)(10-26)
    • 3.9 TileOPs 在審:稀疏長序列的 GD

N 預填充拆分掃描(#2779)(10-26)

  • 3.10 TileFoundry:動態共享記憶體按核心實際檢視定尺寸(#271)(10-26)
  • 3.11 TileFoundry:整視窗按運算元本身下發,sm_89 的 GEMM 恢復可編譯(#273)(10-26)
  • 3.12 TileFoundry 新缺陷:分析側把共享記憶體峰值算重、且不傳到 TIR(#272)(10-26)
  • 四、社群、教程與活動
    • 4.1 語言服務:實驗性 CUDA 除錯上線,接 cuda-gdb 與 Nsight(10-26)
    • 4.2 文件站:單日八筆更新(10-26)
    • 4.3 貢獻流程治理四連:PR 描述規範、純文件跳過後端測試、Agent 標題確認、評審機器人降噪(#3671 / #3672 / #3673 / #3674)(10-26)
    • 4.4 社群新倉:RISC-V 社群樞紐倉、CANN 運算元倉、Hexagon 後端倉、除錯工具倉(10-26)
    • 4.5 產業側:昇騰與 DeepSeek 生態差距收窄的報道(10-26)
  • 五、趨勢觀察
  • 附、素材與核查說明

今日重點

24 小時內 tile-ai 組織合入 30 筆(tilelang 14、TileOPs 14、tilelang-hygon 2),在審更新 49 筆,新開 issue 11 條、關閉 12 條。昨日 72 小時視窗裡 78 筆的量在單日回落到 30 筆,節奏正常,但內容重心明顯從「效能清掃」轉向「流與屏障語義」:主倉合入的輸出分配修復(#3649)把 H100 上五次跑三次輸出損壞的問題降到零,原因是三個介面卡接受了呼叫方的 stream 卻在當前流上分配 out_idx 輸出,快取分配器按分配流排序複用,於是別的流上排隊的活覆蓋了核心結果;同一天浮出的三筆 ThreadSync 缺陷(#3680)指出柵欄被放進發散分支、分支裡的柵欄錯誤地清掉外層掛起依賴、以及按執行緒計數迴圈的柵欄位置,配 #3682 與 #3681 兩筆在審修復,另有 tvm 分叉倉一筆把布林位運算譯成邏輯運算,讓證明器不再在 10000 步預算內放棄(#88)。釋出側沒有進展:TileOPs 的首發追蹤單(#2738)只剩 TestPyPI 試跑與首個 tag,PyPI 上最新仍是 2025-10-09 的 0.0.1.dev2,倉庫無 tag 無 release。後端側海光一天合入兩筆大會:暫存器流水線加 warp 發散(#18,28 檔案 +5312/-64)和原生 B16 儲存加動態 WASP GEMM(#28);昇騰主線合入了帶版本片上緩衝的精確定址介面(#3607,14 檔案 +982/-26)。在審最重的仍是 #3666,規模從昨日的一千出頭擴到 1194 檔案 +106648/-104497。


一、核心專案進展

1.1 主倉:out_idx 輸出改在呼叫方流上分配,H100 復現的五次三次損壞歸零(#3649)

合於 10-26 22:36,5 檔案 +183/-31,修 #3648。cython、nvrtc、cutedsl 三個介面卡都接受 stream= 並在該流上啟動,但 out_idx 輸出是用 torch.empty 在當前 torch 流上分配的;快取分配器只按「分配流」排序複用,於是一塊仍有排隊工作的視訊記憶體被當成空閒塊返回,那批工作隨後覆蓋核心結果。修法是新增 output_allocation_stream(stream, device):呼叫方流與當前流不同時返回撥用方流,否則返回空;控制代碼為零時對映到預設流,因為舊版 PyTorch 的 ExternalStream(0) 會給出池化流而不是預設流。H100 80GB 上的復現從 4/5、5/5、5/5 損壞變為三個介面卡各 0/5,新增測試檔案 10 項透過。作者明確否決了加 wait_stream 的寫法:那會讓每次呼叫都把啟動流排在當前流的全部工作之後。

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

1.2 主倉:NVRTC 裝置端編譯接進 TVM-FFI 執行路徑,新增編譯鍵項(#3285)

合於 10-27 00:02,12 檔案 +133/-133。新增 tl.cuda_compiler="nvrtc",與既有 TVM-FFI 宿主執行時配合:NVRTC 把裝置端 CUDA 源編成 cubin,共享 CUDA 執行時模組匯入這些位元組,宿主側繫結、標量轉換、條件、迴圈與啟動準備仍由 TVM-FFI 生成,預設值仍是 nvcc,舊的 execution_backend="nvrtc" 介面卡保持原實現。二進位制快取鍵里加進了編譯器名與 NVCC/NVRTC 版本,與目標細節和有效編譯選項一起參與雜湊;架構 token 如 90a 與 100f 原樣保留,不相容目標在編譯前報診斷;每個成功建立的 NVRTC program 都在 finally 裡銷燬。該 PR 的 CI 在 H20 上跑出 5506 透過、491 跳過,示例 76 透過,ROCm 7.2 通道 3220 透過,昇騰通道 1797 項透過。

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

1.3 主倉:批次 TMA 載入的 cache hint 少一個模板引數,DeepSelect 的 TMA 路徑編譯不過(#3683)

合於 10-27 00:11,2 檔案 +20/-11。eviction_policy="evict_first" 或 "evict_last" 的批次 TMA 載入在 main 上編譯失敗:程式碼生成發出一個 CacheHintSm90 模板實參,而裸指標版的 tl::tma_load 過載只接受屏障型別。這條缺陷直接壓住了 DeepSelect 的 TMA 路徑,在 GB200 加 CUDA 13.1 上用該示例的 ring 與 replay 用例復現。修法是從 #3296 裡抽出快取提示這一塊,補上缺失的模板引數並把選定的策略傳給 PTX 指令;共享記憶體對齊那一部分仍留在 #3296。修復後同一臺機器上 DeepSelect 全檔案 144 項透過、TMA 複製測試 22 項透過,CUDA 12.9 的 SM90 編譯檢查也確認三種策略的 PTX 都帶上了應有的運算元(#3296 的其餘改動尚未合入)。

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

1.4 主倉:歸約的 FP8/FP4 非有限常量寫出非法字面量,亞位元組累加器改為早拒(#3189)

合於 10-26 15:44,7 檔案 +204/-4,修 #2998。T.reduce_max 與 T.reduce_min 在輸出為 FP8 E4M3 時發出非法的 inff 字面量;現在非有限的 E4M3、E8M0、FP4 常量改走既有的轉換構造器並傳入合法的 CUDART_INF_F 與 CUDART_NAN_F 表示式,飽和與 NaN 行為保持不變。同一筆改動讓普通與延遲歸約提前拒絕亞位元組數值累加器(FP4、int4 之類),給出一句明確診斷;FP4 輸入配 FP16 或 FP32 累加仍受支援,布林歸約豁免,E8M0 歸約支援仍不在本次範圍內。B300 加 CUDA 13.1 上 28 項聚焦測試全過,完整歸約測試檔案 150 項透過,另有三項批處理校驗失敗是 main 上已有的。

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

1.5 主倉:SM120 分塊縮放 GEMM 顯式展開 K 原子,FullCol 配置提升 51%(#3286)

合於 10-26 16:36,4 檔案 +308/-2。這是 #3257 引入 T.gemm_blockscaled 與 T.mma_gemm_blockscaled 後的收尾:A 運算元取自暫存器分片時,K 原子迴圈原先用 T.serial,每次片訪問都是執行期下標,nvcc 在不自行展開的流水化與 warp 特化核心裡會把分片留在區域性記憶體,也就是把 MMA 路徑溢位掉;改為 T.unroll(..., explicit=True) 後每個訪問都是編譯期常量。第二處是 SFA 與 SFB 都在共享記憶體時,原來每對 MMA 要讀三次共享,現在每 lane 每 K 原子只裝一次緊湊選擇包,行主序與分塊 K 主序兩種佈局都覆蓋。RTX PRO 6000 Blackwell 上 143 項透過、4 項跳過,輸出與 float32 參考逐位一致;暫存器 A 基準裡 FullRow 提升 2.7%、Square 11.1%、FullCol 51%(348.5 到 526.7 TFLOPS),ptxas 報告的溢寫從 664 位元組與 484 位元組降到 320 與 160。作者同時標註:block_K = 128 配行主序共享 scale 的路徑在 main 上本就不過校驗,與本次改動無關,留待另查。

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

1.6 主倉:JIT 輸出形狀裡的標量引數在分配前解析(#3602)

合於 10-26 10:06,7 檔案 +242/-3,修 #3406。Cython 與 PTO 包裝器以前在分配輸出時,遇到用顯式標量 PrimFunc 參數列達的輸出維度就解析不出來;現在介面卡傳出「輸出維度到標量引數位置」的對映,包裝器直接用實參值,且不追加重複的啟動引數。迴歸覆蓋兩個後端、零維與非零維、標量與多維輸出、輸出順序、裝置放置與啟動引數個數。同一族缺陷的後續兩筆仍在審:#3636 處理 (n, 2*n) 這類複合維度表示式,#3638 處理同名但彼此獨立的動態符號被按名字合併的問題。

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

1.7 主倉:sm8x 稀疏後設資料的運算元 dtype 不再無兜底落入(#3335)

合於 10-26 16:54,1 檔案 +12/-5。稀疏後設資料的佈局建構函式按架構分派,把 MMA 運算元 dtype 與後設資料緩衝寬度配對;sm90 分支會拒絕不支援的 dtype,sm8x 分支卻寫成兩條白名單檢查且沒有兜底,於是其他拼寫都會落到下面的放置邏輯,而那段邏輯只從緩衝寬度推 group 與 interweave。後果是 float8_e4m3fn(torch 的正規拼寫)被按 16 位後設資料接受,並套用了 fp16 的 group 32 與 interweave 4 約定;任何未識別的拼寫包括打字錯誤也在任意寬度被接受。現在不支援的 dtype 直接報錯。

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

1.8 主倉:DeepSelect Top-K 示例合入,129 組配置裡相對 torch.topk 最高 8.00 倍(#3304)

合於 10-26 16:53,4 檔案 +1265/-0。示例把 DeepSeek 的 DeepSelect 用 TileLang 重寫,不含內聯 PTX;作者實測 129 組配置、10 組不可用,在 RTX 5090(170 個 SM)加 CUDA 12.9 與 PyTorch 2.7 上跑重複輸入與清快取兩種方法。重複輸入口徑下 BF16、批 6、N 為 16384、K 為 512 的一組時延 5.059 微秒對 torch.topk 的 40.448,加速 8.00 倍;N 為 1048576 的一組 15.503 對 100.254,6.47 倍。FP32 與 K 為 4096 的組合在該封裝裡不可用,K 為 1 的原生路徑有對齊地址錯誤,作者在表裡用橫線標註「不可用」而非零。該示例依賴的快取提示修復正是同日合入的 #3683。

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

1.9 在審:TIRx 重構規模再擴到 1194 檔案,錨定未打補丁的上游版本(#3666)

10-25 開出、10-26 三度更新,當前 1194 檔案 +106648/-104497。方案是把 TileLang 從內建 TVM 遷到獨立釋出的 TIRx,且對固定的上游基線打零補丁:刪除 3rdparty/tvm 及其子模組項,3rdparty/ 只剩 CUTLASS、HIP 標頭檔案與忽略檔案;未修改的 apache/tvm-tirx 原始碼釘在 9af5544d,依賴 apache-tvm-tirx==0.1.0a2 與 apache-tvm-ffi==0.1.15.post0,原先帶 +tilelang.1 字尾的整合補丁不再需要。這是當前佇列裡體量最大的一筆,也決定後續 TVM 側修復(例如同日 tvm 分叉倉的 Z3 布林位運算)如何進入主線。

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

1.10 新增缺陷:ThreadSync 三類錯位柵欄,配三筆在審修復(#3680 / #3682 / #3681)

#3680 於 10-26 19:56 開出,指出 tl.ThreadSync 在三處放錯柵欄:把針對逐執行緒取值的守衛當成塊內一致,於是柵欄留在分支裡,等的是跳過分支的執行緒;if 某一分支裡的柵欄清掉了外層序列的掛起依賴,塊跳過該分支時 if 前的寫與 if 後的讀就衝突;迴圈的行程數是逐執行緒的(如按執行緒號取模)時,柵欄被留在迴圈體內,部分同步的下降直接以「無法證明執行緒矩形」失敗。#3682(3 檔案 +39/-7 級別改動,12 項新測試)用一遍前置分析收集可能逐執行緒取值的緩衝,並把同步摘要按分支與迴圈的可達性收縮;作者宣告該改動只報告競爭、不消除競爭,也不做分支拆分。#3681 修的是另一條同源崩潰:sm_100 起 float32 迴圈被向量化到 256 位,廣播到 8 個 lane 時生成的表示式取了臨時物件地址,NVRTC 以「必須是左值」拒絕編譯,改為用 __float_as_uint 加移位直接拼位。配套的 tvm 分叉 PR #88 把布林位運算譯成邏輯運算,使證明器在 10000 步預算內能得出同一結論。

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

1.11 新增缺陷:sm_100 以上歸約 batch 與佈局衝突,已有一筆在審修復(#3667 / #3684)

#3667 於 10-26 08:31 開出:T.reduce_* 在 sm_100 與 sm_120 上指定 batch 等於 N 時無法下降,因為片段複製在 sm_100 以上走 256 位載入,每執行緒持有的輸出數變少,一個在 sm_90 上合法的 batch 觸發了硬校驗。#3684 於 10-27 05:44 開出並附修法:ReduceOp 改為取 N 的不超過 batch 的最大真因數並給警告,finalize_reducer 保持嚴格因為其測試釘住了行為;新增的 sm_100a 與 sm_120a 下降測試在 main 上失敗、在本分支透過,實測機為 RTX 5070 Ti。同一批歸約問題的另一筆 #3341(批次全歸約的執行緒範圍解析)仍在審。

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

1.12 在審:轉置目標形狀校驗、區域化 kill 重排(#3677 / #3670)

#3677 把 tl.transpose 的目標形狀校驗放進 TransposeNode 構造,在 NormalizeToAccessRegion 之後執行,因此繞過 Python 包裝器直接構造也能覆蓋;相等性用 arith::Analyzer::CanProveEqual 判定而不是比較 Python 維度列表,(3, n+1) 與獨立構造的 (n+1, 3) 仍被接受,常量不匹配則以軸號與兩側長度報錯。此前不匹配的目標會下降成越界下標,產生錯誤結果或越界寫。CPU 目標 13 項測試透過。#3670 換掉共享記憶體活躍性分析裡殘留的 scope_level 機制:該值只在含 SeqStmt 的 For 體內遞增、While 從不遞增,於是兩個方向都會錯——While 情形下 kill 點沒越過回邊,打包器把互斥區域重疊,sm90 上實測結果每輪翻倍;兄弟迴圈情形下又會讓 kill 越過整段區域邊界。新規則按真實的外層 For 與 While 鏈與每段迴圈的連續區間計算,越界與提前停止兩個方向都被夾住。

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

1.13 已關閉:模糊測試系列的六條編譯期崩潰(10-26)

10-26 一次性關閉六條來自模糊測試的編譯期崩潰或錯誤程式碼類缺陷:int16 與 uint16 目標上的 T.atomic_add 因 nvcc 無 16 位過載直接中止(#2572);非布林的緩衝元素直接當 if 或 assert 條件時以內部 is_bool() 斷言中止(#3009);浮點運算元餵給只接受整數的 div 與 mod 族時同樣內部斷言(#3008);int2、uint2 這類不在 1、4、8、16、32、64 之列的整型寬度讓 CUDA 程式碼生成中止(#2996);reduce_max 與 reduce_min 寫進 FP8 或 FP4 分片時崩潰(#2998);以及整緩衝目的地的原子歸約被靜默丟掉除 0 號執行緒之外的全部執行緒(#2959)。最後一條另有文件側修正合入(#3340),把示例裡的目的地從整緩衝改成具名元素,避免示例本身繼續教出錯誤寫法。

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


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

2.1 昇騰主倉:帶版本片上緩衝的精確定址介面(#3607)

合於 10-26 10:04,14 檔案 +982/-26。新增 T.annotate_buffer_addresses,允許在排程後的昇騰流水線裡,對 UB、L1、L0A、L0B、L0C 與 SS 各池內的分配做精確位元組定址;放入定址的緩衝必須先用 T.annotate_buffer_versions 宣告固定份數。校驗發生在版本展開與佈局補齊之後,覆蓋缺少或不支援的份數、重疊、未對齊、動態尺寸與容量溢位;分配器還要為 SIMT 的 UB 上下文預留與 AIV 私有地址空間讓位,L0 的 scale factor 控制代碼跟隨其資料分配。部分對映只釘住給定區間,其餘緩衝在對齊後的空閒空間裡安置且不允許複用已佔用地址;沒有定址標註的池保持原行為。子核投影時只保留活躍儲存副本上的標註。

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

2.2 昇騰在審:手工排程的 FlashMLA,BF16 預填充與 FP8 解碼(#3678)

10-26 19:31 開出,覆蓋 Ascend 950 上的 BF16 預填充與緊湊 FP8 解碼,含可選的 FP8 與 FP4 額外 KV cache、attention sink、非法下標與逐查詢長度。注意力迴圈、線上 softmax、反量化與排程都用 TileLang 表達,只有小段 intrinsic 標頭檔案包裹單條硬體指令。手工實現顯式控制 AIC 與 AIV 流水與同步;持久化版本用 T.Stage 加編譯器管理的緩衝版本與同步,QK 與 PV 使用同一 L0B 儲存環的不同佈局檢視。

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

2.3 昇騰在審:片上記憶體溢位提前拒絕、E4M3 的 UB 到 L1 複製、檢視佈局與指標單位(#3655 / #3639 / #3623)

#3655 讓 MergeUBAllocations 在下降期就按目標物理容量拒絕溢位:兩個獨立迴圈各自合法地選兩份 32 KB 的 L0A 分片,合併後 64 KB 正好佔滿,此前會發出越過物理緩衝末尾的指標偏移且不帶任何尺寸後設資料,現在改為給出一句短的診斷。#3639 開啟 E4M3 的 UB 到 L1 複製:ND 輸入經既有的 scatter 與後複製流水自動重排,或直接拷入預排的 NZ 輸入,複用既有的幾何檢查並給足跡與延遲分析補上 FP8 元素尺寸處理。#3623 讓 T.view 與 T.reshape 造出的緩衝各自保持佈局,避免經共享資料變數解析佈局時選錯 Cube 對映或 DMA 地址,並讓裸指標偏移以自己的元素單位重對映。

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

2.4 昇騰在審:批次排程可選、跨核同步要求批模式、L0C 雙複製直轉(#3675 / #3614 / #3468)

#3675 給昇騰加可選的批排程(schedmode 1):正常模式下排程器只要所需核資源部分可用就啟動運算元,多個流的全核同步運算元可能交錯,批模式要等全部資源就緒再整體啟動,全核同步場景必須用它;#3614 描述的是同一機制在 mode 0 下的症狀與修法。#3468 處理 A5 上 FixPipe 不能同時開雙目標模式與 quant_pre 的限制,讓 T.dual_copy 直接做 L0C 的 fp32 到 f16/bf16 轉換,省掉一次 Vector 轉換與一塊 fp32 的 UB 緩衝,也省掉 L0C 到 UB 的 fp32 頻寬。

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

2.5 昇騰倉:安裝器的 SHMEM 段雙修,空殼 wheel 不再假成功(#1943 / #1946)

兩筆在審 PR 指向同一個缺陷(#1728):install_ascend.sh --enable-shmem 在 build.sh 失敗後不檢查返回碼,繼續 bdist_wheel,package_data 的通配匹配不到檔案也不報錯,於是打出一個只有 __init__.py 的空殼包並 Successfully installed shmem-1.0.0,二十多行後才在 source ../../../install/set_env.sh 處報「設定 shmem 環境失敗」,把排障方向引向環境配置。#1946 只改兩處:加返回碼檢查立即退出,把 set_env.sh 改成基於指令碼目錄的絕對路徑,並用逐字提取 shmem 段的受控 harness 驗證(修復前 10 項中 6 失敗 4 透過,修復後 10 項全過)。#1943 走另一條路:在解除安裝現有包之前依次檢查原生構建、環境指令碼、wheel 構建與 wheel 存在性,新增 10 項 CPU 控制流測試,並說明它修的是安裝器續行與診斷,不含上游 SHMEM 編譯失敗本身。

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

2.6 昇騰倉:W4A4 的 int4 矩陣乘結果錯誤,根因是半位元組打包被當成 1 位元組(#1731)

#1731 在 10-26 更新,定位與復現都相當完整。整個 C++ 模板棧用 sizeof(AscendC::int4b_t) == 1 推導分形幾何,但該型別是半位元組打包,32 位元組的 C0 塊正確元素數是 64 而非 32,於是三層搬運全部出現二倍偏差:L1 到 L0A/L0B 的 LoadData 引數按錯誤常數計算,資料在片上堆疊到錯誤位置;B 運算元的轉置載入被 load_cbuf_to_cb_s4 硬編碼關閉並靜默丟棄標誌,B 永遠拿不到 mad_s4 期望的 nZ 佈局;UB 與 GM 之間的位元組長度按每元素 1 位元組算,每行多拷一倍並混入相鄰行。計算指令本身是對的,於是症狀是「量級正常、數值全錯」(原報告 99.6% 不匹配)。修復全部落在 src/tl_templates/ascend/common.h 與兩處變換,新增 W4A4 示例與 10 項 int4 迴歸,1024 立方端到端從 99.6% 不匹配變為零不匹配,既有 gemm_v0 套件 48 項、l1_to_l0 與 gm_to_ub 263 項、逐元素 492 項均透過。

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

2.7 昇騰倉:超過兩維的 GM 到 UB 複製按 DMA 行數摺疊前導維(#1805)

#1805 修 #1231。copy_gm_to_ub 與 copy_ub_to_gm 的 DMA 模型是二維矩形,而 AscendCopy::Lower 從 region 的活躍維裡只取最後兩個當行與列,超過兩個活躍維時前導維被靜默丟棄:3D 塊複製 A[t:t+16, :, :](region 為 16、4、8,共 512 元素)被降級成 4 與 8,DMA 只搬 32 個元素,UB 分片其餘 15 個 token 未初始化,後續計算讀到垃圾值。修法是把前導維摺疊進 DMA 的行數,兩個方向同改。同日昇騰倉另有一批文件事實核查 PR 更新(sin 與 cos、axpy、clamp、bitwise 系列、clear、fill),其中 clamp 那筆順帶修了 PTO 上 T.tile.clamp 下界無效的問題。

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

2.8 昇騰新缺陷:Cython 包裝器每次呼叫都造一個分析器,約 0.28 毫秒(#1945)

10-26 15:52 開出,報告者給出逐層定位與因果驗證。CythonKernelWrapper.forward() 每次呼叫都無條件構造一個 TVM Analyzer(),本機中位數約 0.28 毫秒,而它只在動態 shape 分支裡被用到一次;靜態 shape 核心的 dynamic_symbolic_map 為空字典,simplify() 從不執行。因果驗證的做法是把模組裡的 Analyzer 全域性換成「一調 simplify 就拋錯」的樁,所有呼叫照常成功,全呼叫耗時從 442 至 467 微秒降到 104 至 107 微秒,且與核心規模無關(1 個 block 與 32 個 block 同量級)。實測環境為 Ascend 910B3 加 CANN 9.2.0。這對小運算元密集呼叫的推理鏈路是固定開銷,值得關注修法走向。

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

2.9 海光:暫存器流水線與 warp 發散合入,28 檔案(#18)

合於 10-26 13:33,28 檔案 +5312/-64,10-08 開出、跨兩週多收斂。改動給 tilelang-hygon 加暫存器流水線,可與既有的共享記憶體流水線協同,對應的提交與屏障指令按流水級插入;warp 數超過 4 時啟用 warp 發散。作者給出的結論是 example_gemm 達到與 rocblas 相當的效能。同一倉庫當天還需注意 #18 的合併順序:它先於 #28 合入,屬於把主幹改動併入分支後的大合併,因此行數遠大於 #28。

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

2.10 海光:原生 B16 儲存與動態 WASP GEMM 合入,scale 緩衝雙 bank 在審(#28 / #29)

#28 合於 10-26 16:23,6 檔案 +340/-230,三件事:保持 MLS 的 Let 引數發射順序、把 GEMM 的 warp 劃分傳播到矩陣儲存、拒絕打包的 B16 到 B32 輸出別名;同時讓 WASP 維度可動態,按軸推斷對齊,修 pong 的 K 尾過濾,並把 WASP 端到端迴歸與基礎 MLS 測試分開。#29 於 10-26 17:13 開出並在審:給 scale 緩衝的行規劃一個與 LDS bank 分開的雙 bank 分割槽,讓分配、scale 複製與 MMAC 下降共用同一套行對映,保留物理區間生命週期與 m0 對齊分頁,把空閒 wave 的工作攤到兩個 bank 上,並用於手工 warp 特化下的分塊縮放 GEMM。

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

2.11 沐曦:組織檢視有推送記錄,六個分支均無視窗內新提交(10-26)

組織頁顯示 tilelang-metax 在視窗內有推送時間戳(北京時間 10-26 15:08),但逐分支核對後沒有找到對應提交:預設分支 main 最新仍是 2026-09-29 的版本號提升,dev 分支最新 2026-10-16 的一次合併,0.1.14 分支 2026-09-17,0.1.12 與 race 更早;事件流裡也沒有視窗內的推送事件。據此按「無實質變更」處理,不計入本期條目,也不猜測其性質。該倉最近一次正式釋出仍是 10-16 的 0.1.15。

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


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

3.1 TileOPs:首發仍在最後一格,製品仍是 2025-10-09 的開發版(#2738)

首發追蹤單 10-26 更新,八項清理已完成六項:命名與清單引數、打包與倉庫檔案、Python 3.13、公共 API 與擴充套件點、運算元行為約定、輸入校驗生成、文件站。剩下兩格:夜間基準(修兩條壞掉的基線 #2763 已合入;把新行拉到所有對照器之上,累計掃描改成單次啟動 #2766 已合入,GQADenseWithLSEFwdOp 的 D128 程式 #2774 已合入,但 MHA 行仍只有 FA3 的 0.996,調優繼續)與釋出本身(第三方宣告、PyPI 安裝文件、開發者依賴組、CHANGELOG 與釋出步驟 #2753,以及需要維護者執行的 TestPyPI 試跑與首個 tag)。核對 PyPI 與倉庫:tileops 最新仍是 2025-10-09 的 0.0.1.dev2,共三個版本,倉庫沒有任何 tag 與 release。首發處在「文件就緒、製品未上傳」的最後一格,已持續多日。

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

3.2 TileOPs:D128 稠密前向拆 key,prefill-rope 從 31.04 微秒降到 17.09(#2777)

合於 10-26 22:45,2 檔案 +527/-124。D128 稠密呼叫若 128 行的瓦片讓 SM 閒著,就把每個瓦片的 key 拆成 min(SM 數除以瓦片數, 最短瓦片的 key 步數) 份,且只在拆分能從最長瓦片勻掉兩步 key 時才拆;每份寫出帶本行縮放最大值與和的未歸一 fp32 輸出,合併核心按最大最大值加權,因此抵消項與大 logit 都不丟精度。瓦片最後一步 key 中,若某 warpgroup 最後需要的 key 落在低 64 個 key 內,softmax 只對那 64 個做掩碼與指數,warpgroup 0 的因果對角步必然命中。H200 1500 MHz 上 prefill-rope 的 f16 與 bf16 從 31.04 與 31.01 降到 17.09 與 17.09 微秒,對 FA3 是 2.324 與 2.328 倍;qwen3-8b 分塊預填充從 31.30 降到 17.79。測試節點淨增 14。

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

3.3 TileOPs:D128 持久化前向按執行緒重寫並寫出 LSE(#2774)

合於 10-26 18:22,2 檔案 +656/-48。根因是 GQADenseWithLSEFwdOp 以前跑在非持久化的 warp 特化核心上,因為持久化核心拒絕 return_lse,夜間基準實測只有 FA3 的 0.73 至 0.78。本次用 TileLang 的逐執行緒原語(TMA 對映、mbarrier、WGMMA 描述符、warp shuffle、stmatrix)重寫 D128 持久化稠密前向,直接寫出 log-sum-exp,同一份實現同時服務 GQADenseFwdOp 與 GQADenseWithLSEFwdOp;當全部 key 與 value 超過 L2 的一半時,一個頭的查詢塊相鄰排布,使併發瓦片共享 L2 裡的 K 與 V;每個瓦片的首次 QK 在上一個瓦片的尾聲裡啟動。公共介面與清單未變。合併後 MHA 行仍為 FA3 的 0.996,追蹤單標註調優繼續,這與 #2777 屬於兩條獨立的最佳化線。

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

3.4 TileOPs:融合加法的 LayerNorm 行跨 CTA 拆分,精度缺陷另開新單(#2768 / #2773)

#2768 合於 10-26 15:45,5 檔案 +525/-24。起點是一次夜間基準的回退(llama-70b 解碼行從 0.0037 升到 0.0044 毫秒),查下來不是程式碼變化而是調優重選:調優鍵改動後 FusedAddLayerNormKernel 重新調優,調優器從 L2 背靠背計時候選,選中了 1024 執行緒,冷啟動比 256 執行緒慢 19%。新寫的 FusedAddLayerNormHeldKernel 把一行(SM90 上是一個 cluster CTA 的份額)留在暫存器裡,在最忙 SM 的成本更低時把一行拆到最多 8 個 CTA,塊大小取整個 warp 數中能整除份額訪問量的那個;派發按「排隊行數」與「暫存器超過 255」兩條互相拒絕來選,調優只提供受調構建份額能填滿的寬度,使解碼行不可能再被重選。llama-70b 解碼 2.368 微秒對 flash-attn 的 4.800。同一筆改動暴露的精度問題另開 #2773:行遠離零時 fp32 求和的均值與方差失去精度(常數 10000.1 的行歸一化後約 0.7 而非偏置),靠近 1e36 時平方和溢位導致全 NaN,作者的建議是改用 Welford 或縮放兩遍法。

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

3.5 TileOPs:逐元素與歸約四條,clamp 的 bf16 行提速 5.9%(#2766 / #2767 / #2769 / #2770)

四條同日合入,都屬「先找根因再改」。clamp(#2767)改為在自身精度內比較,根因是舊實現把每個元素加寬到 fp32、每個界都做一次 NaN 判定與選擇,約 5 條指令一個元素;現在每對元素一條 HMNMX2.NAN,16384 見方 bf16 從 0.25704 毫秒降到 0.25024,對 torch-compile 1.0587 倍,並在全部 16 位模式、NaN、正負無窮與帶符號零的邊界上驗證與 torch 的 fp32 比較逐位一致。累計掃描(#2766)把 dtype=float32 的 cumsum 與 cumprod 從「先加寬成一份再掃」改成一次啟動直接寫目標精度,hidden-state-f32 工作量下搬運從 112 MiB 降到 48 MiB,cumprod 47.36 降到 14.62 微秒、cumsum 47.86 降到 14.62。GELU 乘(#2769)改為經共享 gelu() 在 fp32 求值並只舍入一次,舊的 erf 路徑帶 NaN 選擇分支、三次 16 位舍入與更長的依賴鏈;小門控網格的塊數從 56 降到 14。二元運算(#2770)把二位元組運算元的每執行緒元素數從 8 提到 16,保持兩條 16 位元組載入在飛,34 行預設配置隨之變化,MinimumFwdOp 的 llama-7b 行從 8.42 降到 8.00 微秒。

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

3.6 TileOPs:呼叫張量的對齊判斷上移到運算元層,三筆收尾(#2772 / #2778 / #2780)

#2772 合於 10-26 15:28,55 檔案 +373/-224:執行器在實現執行前把儲存起點不在 16 位元組邊界的呼叫張量拷到對齊儲存,成功後把被寫的那幾個拷回,並返回撥用方的張量;若被複製的張量與其他實參重疊(且至少一方是被寫的),在執行前直接拒絕,因為複製會掩蓋核心自身的檢查;無地址的儲存(假張量或拒絕讀取的包裝)不複製;核心側的對齊與暫存邏輯(含 MoE 分組 GEMM、因果卷積狀態、融合 QK-norm RoPE 的 qkv)全部刪除。代價是熱路徑上對齊呼叫的宿主時間漲 1 至 2 個百分點,MoEPostPermuteFwdOp 帶 out 時 4.4 個百分點。#2778 把這項判斷折進生成的觀察器並讓計劃快取決定複製,熱路徑只讀該決定,執行器原來單獨的地址遍歷被刪。#2780 於 10-26 開出:目前只有四個核心檢查 out 與輸入重疊,其餘帶 out 的運算元會靜默競爭,且每次檢查都重複讀 data_ptr(MoEPostPermuteFwdOp 一次呼叫多 12 次),方向是在清單裡逐運算元宣告是否允許重疊,在運算元層統一檢查。

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

3.7 TileOPs:偏移儲存三類修復,SM90 分組 GEMM 與輸出 ABI(#2761 / #2765 / #2762)

#2761 讓分組 GEMM 的 MMA 實現接受起點帶偏移的連續張量,並對齊向量訪問,把結果寫回偏移輸出時保持標識與未觸及的專家行,複製前檢查重疊;原實現遇到這類檢視會以 CUDA 地址未對齊失敗,新增 16 個迴歸覆蓋 FP16 與 BF16 各自的 a、b、bias、out 偏移,配緊湊與掩碼兩種佈局。#2765 把同一能力補到 SM90 的 MoE 分組 GEMM:儲存起點不在 16 位元組邊界時拷 a、b 並暫存 out,vector_aligned 現在寫明「被寫的張量要暫存並在啟動後拷回」這條規則,BF16 與 FP16 的 MMA 核心改用它,FP8 核心從 SM90 FP8 核心繼承。#2762 修輸出 ABI:當 TVM FFI 介面卡宣告由被調方分配輸出時,交給 TileLang 課執行體自己分配並返回,其他介面卡保持原有分配路徑;報告的迴歸是舊實現在新構建上 27 個 Mamba 用例全敗,候選實現在公開呼叫路徑上透過。

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

3.8 TileOPs 在審:FP8 查詢的分頁預填充改按 e4m3 乘 value(#2775)

10-26 17:20 開出,10-27 07:00 仍更新。GQAPagedFP8WSKernel 把 FP8 查詢的 value 乘積改在 e4m3 裡做:負責轉換的 warp 把每個 value 瓦片轉置成 K 主序的 e4m3 運算元(ldmatrix 轉置、按字做一次位元組置換、stmatrix),在每個 16 內排好 key 的順序,使每執行緒的分數正是它供給的機率,再用暫存器內打包把機率舍入到 e4m3。softmax 按 FA3 的 FP8 形式:2 的 8 次方偏移讓小機率不進 e4m3 的次正規數,行最大值隨每個更大的分數移動,key 瓦片從最後一個往回掃。H200 上 llama-8b-fp8-prefill 從 369.3 降到 339.8 微秒(相對 FA3 從 1.000 到 1.087),qwen3-fp8-ragged-odd 從 130.9 降到 123.9。作者如實標註精度容差從 2 的負 4 次方放寬到 2 的負 3 次方,原因是 e4m3 機率本身有半個 2 的負 4 次方量級的舍入,FA3 在同類輸入上同樣超出。

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

3.9 TileOPs 在審:稀疏長序列的 GDN 預填充拆分掃描(#2779)

10-26 開出。在 SM89 上對稀疏的長序列 GDN 預填充網格複用既有狀態掃描與獨立輸出程式:適用條件是稠密、key 主序的 D128 輸入、至少 16384 個 token、且狀態掃描 CTA 不超過 SM 數的四分之一,其他呼叫保持原實現。分解方式寫程序序快取標識,並補了尾塊與初始狀態的常駐迴歸。實測(CUDA 13、PyTorch 2.13、TileLang 0.1.15)32768 長度 bf16 為 1.0384 毫秒對 FLA 的 1.0656,65536 為 2.1178 對 2.1786,16384 則仍是 FLA 更快(0.4695 對 0.4409)。作者明確說明收益只在這組四個 value head 上成立,八個與十六個 value head 的既有模型形狀不選這條路徑,候選也不普遍快於 FLA。

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

3.10 TileFoundry:動態共享記憶體按核心實際檢視定尺寸(#271)

10-26 開出並在審,12 檔案 +340/-32。合成的宿主入口以前給每個 TIR 核心都傳 dynamic_smem=0,用整數 TensorView 偏移訪問共享記憶體的核心因此在沒有動態共享記憶體的情況下執行,超過 48 KiB 時也從不提升 cudaFuncAttributeMaxDynamicSharedMemorySize。作者選擇按核心實際達到的最遠位元組數來定尺寸,而不是按目標容量,理由有兩條:CUDA 把靜態與動態共享記憶體算在同一個每塊限額裡,而該執行時自帶若干靜態共享記憶體(ProgramMetaData 或共享的分配張量),按滿容量申請會直接起不來;每塊申請多少還決定一個 SM 上能放幾塊,sm_89 上申請 99 KiB 時哪怕只用 512 位元組也只能放一塊。程式碼生成同時改為在首個語句前宣告一次 smem_base,因為在 loop_outputs 裡它原先落在第一個迴圈的塊內,第二個迴圈看不見,nvcc 報未定義識別符號。H200 與 sm_89 上兩項 GPU 測試透過,未打補丁時前者非法訪存、後者編譯失敗。

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

3.11 TileFoundry:整視窗按運算元本身下發,sm_89 的 GEMM 恢復可編譯(#273)

10-26 開出並在審,5 檔案 +54/-24。schedule finalize 給 sm_89 GEMM 寫出的程式碼編不過:下降把整個累加器包進一個從零開始、覆蓋其全長的視窗 T.ptr_of(acc[0:0+16, 0:0+8]),而 CUDA 程式碼生成對分片暫存器瓦片的指標切片報「需要普通源佈局」;再往下 MMA 把暫存器分片讀成 lhs_frag_dst[0:0+16, 0:0+32] 這類裸切片,同樣沒有處理分支。這些視窗其實都是整個運算元。修法放進下降:ConvertHIRToTIR 一次判定視窗是否覆蓋其運算元(起始為靜態零、尺寸等於運算元形狀),三個構造視窗的地方共用該判定,最終寫出 T.ptr_of(acc),MMA 直接讀分片。49 個有 TIR 孿生的排程夾具裡 29 個的 finalize 輸出按此改變,sm_89 GEMM 與 mesh 檢視恢復可編譯,兩個 wgmma 輸出越過了指標切片拒絕、停在下一處檢視,而那一處佈局在 main 上同樣被拒。

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

3.12 TileFoundry 新缺陷:分析側把共享記憶體峰值算重、且不傳到 TIR(#272)

10-26 開出,作者是在評審 #271 時發現的,兩條缺口都指向同一種後果:分析能透過的程式在啟動時被拒,或最終 TIR 無法自述需要多少共享記憶體。第一條,記憶體分析把 smem 與 facts.memory.shared.per_cta 比較(sm_89 上 101376 位元組),但 CUDA 把這個限額同時作用於靜態與動態共享記憶體,而執行時自身持有 12 位元組的 __shared__ ProgramMetaData,這項佔用只寫在 runtime.cuh 裡,沒有告知分析,於是恰好差在 12 位元組以內的放置能透過分析而在啟動時被拒。第二條,分析把每 CTA 的共享記憶體高水位記在 RegionMemoryMetadata.peaks,但 ConvertHIRToTIR 不把它寫進產出的 TIR,PrimFunction.metadata 既不列印也不解析,於是 finalize 之後必須重算一遍(#271 正是靠掃描所有整數地址的共享記憶體檢視重算),而作者寫的宿主入口申請過少也不會在執行前被發現。

來源:https://github.com/tile-ai/TileFoundry/issues/272


四、社群、教程與活動

4.1 語言服務:實驗性 CUDA 除錯上線,接 cuda-gdb 與 Nsight(10-26)

tilelang-lsp 在 10-26 17:11 推入首個除錯提交(78a5761,標題為加入實驗性 TileLang CUDA 除錯)。內容四塊:新增 tlsp.debug 啟動器,做 Python 行號對映、加 -G 與 -lineinfo 編譯、隔離快取並抓取生成源;把 Nsight 的 cuda-gdb 接進 VS Code,支援 Python 斷點、啟動前檢查與生成 CUDA 源的瀏覽;補單元、啟動配置與可選 GPU 驗收測試(覆蓋待定斷點、裝置端單步、執行緒檢視與數值);中英文文件與 CI 同步更新。這是該倉自 8 月以來第一次實質提交,此前最後一次更新是 8-04 的 CI 測試。

來源:https://github.com/tile-ai/tilelang-lsp/commit/78a5761404

4.2 文件站:單日八筆更新(10-26)

tilelang.github.io 在視窗內推送八次(北京時間 07:32、10:16、10:21、16:51、17:25、22:42、次日 00:08、00:16),提交標題統一是 Update docs,屬於隨主倉合入自動同步的文件構建。同日還有第三方本地化倉(tilelang-docs-l10n)一次推送與中文 RISC-V 社群樞紐倉的首次內容提交,說明文件同步鏈路正常運轉。

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

4.3 貢獻流程治理四連:PR 描述規範、純文件跳過後端測試、Agent 標題確認、評審機器人降噪(#3671 / #3672 / #3673 / #3674)

四條同日合入,全部圍繞「評審成本」。 #3671 在 CONTRIBUTING.md 裡寫明 PR 描述應有問題、方案與取捨,證據要貼在它支援的論斷旁邊,缺陷修復寫前後行為或迴歸測試的意圖與結果,效能改動要給可比測量與條件。#3672 讓純文件 PR 跳過 CUDA、ROCm、Metal、Ascend 與 CuTeDSL 的構建與測試,做法是用既有 lint 作業的檢出看合併基差異,改動路徑全是 Markdown、reStructuredText 或 docs/_static/ 下時跳過,白名單刻意收窄,可執行示例、構建檔案與工作流不在此列。#3673 給提交技能加確認步:最終 diff、驗證與 PR 描述就緒後,先把擬用標題原文給貢獻者看,由其確認或替換。#3674 關掉評審機器人的自動摘要與文件字串覆蓋率檢查,並加共享評審指示,要求把發現錨定在當前 diff 與實際呼叫路徑上、區分迴歸與既有侷限、先比較已有實現再提修法。

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

4.4 社群新倉:RISC-V 社群樞紐倉、CANN 運算元倉、Hexagon 後端倉、除錯工具倉(10-26)

視窗內出現若干個值得登記的外部推進。TongmingLAIC/tilelang-riscv 於 10-23 建立、10-26 推送,定位是 TileLang 在 RISC-V 上的社群樞紐,聚合編譯器後端、AI 運算元、開發工具與文件,目前仍是空骨架,但方向與國內 RISC-V AI 棧的訴求一致。Atlas-Ascend-Fork/ops-tilelang 是 CANN 社群面向昇騰 NPU 的 TileLang 運算元倉,強調統一管理核心實現、自動化用例、標準負載與效能基線,10-26 有推送。xwhzz/tilelang-hexagon 做高通 Hexagon 後端,倉庫 32 MB,10-26 有推送。superAngGao/tilelang-debugger 是 H200 上按源定位的除錯工作區。另有 RuneFang/tvm_tilelang_cookbook(37 星)與 localizethedocs/tilelang-docs-l10n 的常規更新。

來源:https://github.com/TongmingLAIC/tilelang-riscv

4.5 產業側:昇騰與 DeepSeek 生態差距收窄的報道(10-26)

10-26 有媒體轉述 SemiAnalysis 10-22 的報告,稱 DeepSeek 與昇騰的深度配合正在收窄與英偉達在晶片硬體與軟體生態上的差距,並預計 DeepSeek 到 2027 或 2028 年將主要依賴昇騰技術棧;報道特別提到,這份判斷出現在 DeepSeek 釋出面向昇騰生態的軟體工具數週之後,其中包含六個與英偉達平臺一一對應的軟體模組。該訊息本身不涉及 TileLang 的功能變更,本期僅作生態背景登記:這六個模組裡的高階語言編譯工具即 tilelang-ascend,昇騰倉當日的安裝器與 int4 矩陣乘修復正落在同一條鏈路上。

來源:https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report


五、趨勢觀察

一、正確性缺陷的重心從「效能回退」移到「編譯期不報、只在特定架構或儲存域上錯」。當日主線三筆都屬此類:輸出在錯誤的流上分配(#3649,只在非預設流且當前流有排隊工作時損壞)、TMA 快取提示少一個模板實參(#3683,只在帶 eviction_policy 的批次載入上失敗)、sm8x 稀疏後設資料沒有兜底(#3335,只有 torch 正規拼寫才被認出,且套錯 fp16 的佈局約定)。這類缺陷的共同點是有硬體或引數前提,常規 CI 覆蓋面之外,因此越來越多由模糊測試系列驅動——本章當日關閉的六條裡有五條來自該系列。

二、同步與記憶體規劃成為主戰場。當日浮出的 #3680 三條 ThreadSync 缺陷、配 #3682 與 #3681 兩筆修復、tvm 分叉倉 #88 的證明器修正、主倉 #3670 的區域化 kill 重排、TileFoundry #272 的共享記憶體峰值不落 TIR,指向同一件事:編譯器的同步與記憶體判斷正在從「啟發式」轉向「可證明」。這也解釋了 #3666 為何重要——它把內建 TVM 換成獨立打包的 TIRx,此後這類證明器修正才能以正常版本節奏進入主線,而不是停在分叉倉裡。

三、後端適配進入「精確控制」階段。昇騰的 T.annotate_buffer_addresses 讓 UB、L1、L0A、L0B、L0C、SS 各池可以按位元組定址並做重疊、對齊與容量校驗,海光的暫存器流水線、warp 發散與 scale 緩衝雙 bank 分割槽,都是把「能跑」推向「能按硬體結構擺」。這些能力集中在最近兩週,說明國產後端已過可用性門檻,開始處理效能上限問題。

四、釋出工程仍是短板。TileOPs 的首發在文件、清單、打包、API 穩定性、輸入校驗全部就緒之後,卡在 TestPyPI 試跑與首個 tag 上已多日,PyPI 製品停留在 2025-10-09。對一個 196 星、當日合入 14 筆的運算元庫來說,可安裝性比再最佳化 1 個百分點更能決定外部採用速度;追蹤單裡唯一未完成的技術項是 MHA 行 0.996,管理層不應讓它繼續阻塞 tag。

五、外部生態在同步擴張。RISC-V 社群樞紐倉、CANN 運算元倉、Hexagon 後端倉同日出現或更新,說明 TileLang 正在從「tile-ai 組織的專案」變成多平臺各自建倉的生態。對國內讀者,RISC-V 方向尤其值得跟蹤:它與此前 AI 眼鏡、端側推理的硬體路線直接相關,但目前仍只有空骨架,實質性進展需要看後續是否有可執行的運算元與後端落地。


附、素材與核查說明

信源核查表

信源 核查結果
GitHub tile-ai 組織(29 倉) 視窗內推送 7 倉:tilelang、TileOPs、TileOPs-nightly、tilelang-hygon、tilelang-metax、tilelang-lsp、tilelang.github.io
合入真值(is:merged) 30 筆,逐筆核對合入時間落在視窗內
在審真值(is:open 更新) 49 筆,按更新時間排序全量核對
新增 issue 11 條,逐條讀正文
關閉 issue 12 條,其中 6 條為模糊測試系列
PyPI(tileops / tilelang) tileops 最新 0.0.1.dev2(2025-10-09);tilelang 最新 0.1.15(2026-09-30)
釋出頁 tag 核查 TileOPs 無 tag 無 release;tilelang 最新 v0.1.15(2026-09-30)
tilelang-ascend 倉 視窗內無推送(上次 10-23),但 PR 與 issue 活動密集
tilelang-metax 倉 組織檢視有推送時間戳,六個分支均無視窗內提交,按無實質變更處理
採用方(TileKernels / FlashQLA) 均無視窗內推送(10-14)
Google News 中英文多組查詢 24 小時內無 TileLang 元件精確命中;一條產業側報道見 4.5
Hacker News 視窗內一條前端條目,與 TileLang 無關,剔除
arXiv 關鍵字檢索無視窗內新論文,最近一篇為 2026-07-24
會員與採用方倉社群檢索 六個外部倉視窗內有推送,見 4.4

完整信源清單

/tilelang/pull/3468>

#2775)— https://github.com/tile-ai/TileOPs/pull/2775