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

信源:GitHub tile-ai 組織逐倉核查(Atom 提交流,覆蓋 tilelang / TileOPs / TileFoundry / tilelang-ascend〔ascendc_pto〕/ tilelang-hygon〔main 與 feature 分支〕/ tilelang-metax〔dev〕/ tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack)+ 合入與在審真值檢索(is:merged / is:open 兩組,按倉庫分組統計)+ 釋出頁 tag 核查 + 文件站提交核查 + 採用方倉核查(TileKernels / FlashQLA)+ 社群新倉檢索 + Google News 中英文多組查詢(走代理)+ Hacker News + arXiv + 第三方產業分析報告。


本期索引

  • 今日重點:TileFoundry 0.0.3 正式釋出,把指令級排程交給 agent;TileOPs 一天合入 20 筆,MLA 反向與 SM89 小規模路徑同日補齊;主倉八筆合入集中在”不報錯但結果不可信”的編譯器缺陷;昇騰 DeepGEMM 對齊由在審轉合入(10-21)
  • 一、核心專案進展
    • 1.1 主倉:批式 finalize_reducer 此前只歸約首個輸出塊(#3070)(10-22)
    • 1.2 主倉:BF16→FP32 向量化轉換改走內聯 PTX 打包 lane 轉換(#3595)(10-21)
    • 1.3 主倉:展開 pass 保留空 stepped 迴圈,ROCm 的 Select 改逐 lane 下降(#2887)(10-21)
    • 1.4 主倉:軟體流水保留傳遞性複製依賴,執行緒域參與區域衝突判定(#2822)(10-21)
    • 1.5 主倉:cast 的源/目標/舍入三組合前置校驗,拒絕靜默丟註解(#3350)(10-21)
    • 1.6 主倉:無 arch 的 CUDA target 拒絕而非靜默填 sm_50(#2860)(10-21)
    • 1.7 在審:移位符號性與 TIRx 迴圈作用域修復在 tvm fork 與主倉成對提交(#3597 / tvm #87 / #86)(10-21)
  • 二、多後端適配(昇騰 / 沐曦 / 海光 / 摩爾執行緒等)
    • 2.1 昇騰評測:SemiAnalysis 深度報告設 TileLang Ascend 專章,稱六倉中當時只有 TileLang 可用(10-21)
    • 2.2 主倉昇騰:DeepGEMM 啟發式與 epilogue 對齊由在審轉合入(#3577)(10-21)
    • 2.3 主倉昇騰:warp_reduce_max/min 在 PTO 上的整數符號性修復(#3582)(10-21)
    • 2.4 昇騰倉:亞 32B 緊湊行距源無法表達,keepdim reduce 輸出錯亂(#1795)(10-21)
    • 2.5 昇騰倉:JIT 快取改原子釋出,並在快取被刪除時恢復(#1903)(10-21)
    • 2.6 昇騰倉:cann-bench 補 depthwise_conv_2d 示例(#1915)(10-21)
    • 2.7 昇騰在審:indexer_topk_reducesum、GLM-5.3 kpool compress 與 WY fast 系列運算元遷移(10-21)
    • 2.8 海光:feature 分支推進 FP8 解碼 GEMM,定位到地址鏈佔 kernel 13%(10-21)
    • 2.9 沐曦 / 摩爾執行緒 / MLIR-Ascend:視窗內無推送,三條線停在既有版本(10-21)
  • 三、生態與採用方(TileOPs / TileFoundry / 運算元庫)
    • 3.1 TileFoundry:0.0.3 打 tag 釋出,指令級排程 + 重構分析 + H200 與 Apple AMX 示例(10-21)
    • 3.2 TileOPs:MLA 預填充反向補齊 op 層與 SM90 warp 特化 kernel(#2678 / #2680 / #2688 / #2686)(10-22)
    • 3.3 TileOPs:dense GQA 支援 FP32 attention sink,-inf 關閉、+inf 清零該 head(#2515)(10-21)
    • 3.4 TileOPs:SM89 小規模路徑集中補位,FP8 BMM 預設瓦片降到 64×64(#2681 / #2672 / #2669 / #2654)(10-21)
    • 3.5 TileOPs:NSA top-k 在 SM89 上擴到 2048 壓縮塊並補齊輸出槽(#2674 / #2675 / #2673)(10-21)
    • 3.6 TileOPs:GDN 前向在 walk 填滿 SM 時單 kernel 完成,分量化歸一化按 cluster 切寬行(#2682 / #2679)(10-21)
    • 3.7 TileOPs:GQA 反向在 warp 特化 kernel 裡直接服務分組 KV head(#2685)(10-21)
    • 3.8 TileOPs 在審:摺疊 edge 軸單次發射、NSA top-k 收窄列 tile(#2689 / #2690)(10-22)
    • 3.9 採用方:TileKernels 與 FlashQLA 視窗內無推送(10-21)
  • 四、社群、教程與活動
    • 4.1 社群教程:中文《TVM / TileLang Cookbook》更新,從 jit 講到 cubin(10-21)
    • 4.2 社群新倉:tilelang-wrapper-for-sail 建倉,暫無說明(10-21)
    • 4.3 社群工具:遠端自動調優 agent 倉繼續推送(10-21)
    • 4.4 文件與版本:文件站七次自動更新,主倉 VERSION 仍 0.1.15(10-22)
  • 五、趨勢觀察

今日重點

本視窗合入 32 筆(TileOPs 20、tilelang 8、tilelang-ascend 3、TileOPs 文件站 1),另有一項正式釋出:TileFoundry 於 10-21 10:26 打出 0.0.3。三條線的重心各不相同。主倉八筆合入幾乎全是”不報錯但結果不可信”的缺陷清理——批式歸約只處理首個輸出塊、cast 舍入註解被靜默丟棄、無 arch 的 CUDA target 靜默退到 sm_50、流水規劃丟失傳遞性依賴;TileOPs 一天把 MLA 預填充反向從 op 層做到 SM90 warp 特化 kernel,同時把 SM89 上的 FP8 BMM、W4A16、分頁 DSA、NSA top-k 四條小規模路徑一併推進,在審佇列從 16 筆降到 9 筆;昇騰側昨日仍在審的 DeepGEMM 啟發式與 epilogue 對齊(#3577)今日合入,昇騰倉則修掉一個會造成 keepdim reduce 輸出錯亂的記憶體訪問建模缺陷(#1795)。產業側,SemiAnalysis 的華為-DeepSeek 深度報告專門寫了一節 TileLang Ascend,給出的判斷是六個獲得昇騰支援的倉庫中當時只有 TileLang 可用。


一、核心專案進展

1.1 主倉:批式 finalize_reducer 此前只歸約首個輸出塊(#3070)

合於 10-22 01:16,2 檔案 +29/-5,修 #2623。批式 finalize_reducer 的下降只在緩衝區開頭呼叫一次 run_batch,其餘連續輸出塊根本沒有被歸約:輸出維度大於 batch 時,後面各段拿到的不是歸約結果,而是未初始化或上一輪遺留的值。修復按 ceil(OutputDim / batch) 逐塊發射帶正確偏移的批式 all-reduce——128 個輸出、batch=4 時發射 32 次呼叫,同時把重對映 reducer 緩衝改用扁平檢視、每批以常量線性偏移定址,去掉原先的形狀與步長重建,並保留全部 chunk 呼叫。驗證在 L40S 上完成:新增 sum/max/min 數值迴歸,覆蓋 fp16 與 fp32、batch 4/8/16,23 項 finalize_reducer 相關用例透過,外部復現矩陣 16 例零失配。屬典型的”小尺寸正確、大尺寸靜默錯”的一類缺陷。

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

1.2 主倉:BF16→FP32 向量化轉換改走內聯 PTX 打包 lane 轉換(#3595)

合於 10-21 19:28,4 檔案 +19/-7。
新增 tl::__tl_cvt_bf16x2_to_f32x2,用內聯 PTX 把成對的 BF16 lane 一次轉成 f32x2,替掉逐 lane 的標量提升路徑;2、4、8 三種 lane 寬度均保持支援。同步更新了向量化 cast 與佈局推斷的期望值,並把 BF16 cast 的測試收斂到既有覆蓋範圍,不擴大斷言面。改動不涉及語義,只動生成的指令序列,影響面是所有走向量化 cast 的 BF16 計算路徑。

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

1.3 主倉:展開 pass 保留空 stepped 迴圈,ROCm 的 Select 改逐 lane 下降(#2887)

合於 10-21 19:20。迴圈展開 pass 此前會把 trip count 為 0 的 stepped 迴圈當成空殼消掉,改變後續 pass 看到的程式結構;修復改為在展開過程中保留這類空迴圈。同筆還改了 ROCm 側的向量化 Select:改為逐 lane 下降,避免整組下降在不支援的目標上失敗。兩處改動一在 IR 結構、一在後端下降,合在同一筆裡,說明它們由同一個復現場景引出。

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

1.4 主倉:軟體流水保留傳遞性複製依賴,執行緒域參與區域衝突判定(#2822)

合於 10-21 18:13。流水規劃此前只跟直接依賴,間接(傳遞性)的複製依賴會在規劃階段丟失,dependent copy 的等待隨之被刪;同時物理區域衝突檢查在規劃與下降兩側各有一套口徑。修復用一個後向 worklist 同時跟隨標量與 buffer 複製依賴,讓規劃與下降共享同一份物理區域衝突檢查,保留重放繫結在作用域內,並補齊傳遞生產者、strided alias 與順序 TMA 回落路徑的覆蓋。第二筆提交把執行緒域引入區域衝突判定:訪問先展開到物化執行緒域再比較區域,使互不相交的共享記憶體切片仍可流水,同時保住跨執行緒依賴與 strided 緩衝的物理重疊;同一套域用於流水規劃、依賴校驗與 buffer 版本化。

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

1.5 主倉:cast 的源/目標/舍入三組合前置校驗(#3350)

合於 10-21 13:49。round 只在顯式列出的組合下可下降:PTX cvt.rs 要求 f32 源、32 位 rbits 運算元與特定的打包目標。其餘組合此前在分支分派時被放過,舍入註解就此靜默消失;標量迴圈本來就會拒絕這些組合,向量化迴圈不拒絕,於是同一個 cast 在兩種迴圈形態下行為不同。典型例子是 f32→E8M0:它的 helper 向正無窮舍入、根本不接受舍入引數,而向量化分支從未讀過 round。現在在分派之前統一拒絕非法組合,任何分支都無法再丟註解;配套測試只給真正接受 rbits 的 rs 模式傳該引數。

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

1.6 主倉:無 arch 的 CUDA target 拒絕而非靜默填 sm_50(#2860)

合於 10-21 13:32,1 檔案,修 #1049。在有 CUDA toolkit 但沒有可用 GPU 的機器上,未指定 arch 的 CUDA target 會退化成裸 cuda,TVM 隨後靜默填入 arch=sm_50,編譯繼續推進到一個錯誤的 SM,全程沒有任何診斷——這正是 tilelang.compile(func) 配 target="auto" 的陷阱。修復落在 _cuda_target_from_arch:自動檢測與顯式 target="cuda" 共用的唯一入口,改為丟擲可操作的錯誤,並在錯誤裡給出顯式 arch 的兩種寫法({"kind":"cuda","arch":"sm_90"} 或 TILELANG_DEFAULT_TARGET)。headless 編譯只要給了具體 arch 依然可用,因為 codegen 與 nvcc 只依賴 toolkit,裝置訪問推遲到 kernel 啟動。PR 過程中一度加入三個無 GPU 單測,最終提交把它們撤下,改動只剩入口校驗。

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

1.7 在審:移位符號性與 TIRx 迴圈作用域修復成對提交

主倉 #3597 修混合符號移位時丟失符號性的問題,同期在 tile-ai/tvm fork 提交 #87「TIRx 保留移位左運算元型別」,兩者是同一處邏輯在語言層與 TIRx 層的對應改動。另有 tvm #86 把上游 apache/tvm#19927 的迴圈作用域修復回移。fork 在審 25 筆、主倉在審 163 筆,TIRx 側的先行使 fork 承擔了部分改動的前置驗證,主倉再跟進,這條協作路徑本視窗第二次出現。

來源:https://github.com/tile-ai/tilelang/pull/3597、https://github.com/tile-ai/tvm/pull/87、https://github.com/tile-ai/tvm/pull/86


二、多後端適配(昇騰 / 沐曦 / 海光 / 摩爾執行緒等)

2.1 昇騰評測:SemiAnalysis 深度報告設 TileLang Ascend 專章

SemiAnalysis 於 10-21 釋出《The Huawei-DeepSeek Symbiosis》,副題直接列出「華為 Ascend 950PR 深挖與微基準、DSv4.1 Flash 分析、DeepSeek 的昇騰支援、AICPU、TileLang Ascend」。與 TileLang 直接相關的三點:其一,獲得昇騰支援的六個倉庫(TileLang、DeepGEMM、DeepEP、TileKernels、FlashMLA、DeepSelect)中,文章寫作時只有 TileLang 可用,因為 CANN 9.2.0 尚未釋出,而 DeepSeek 要求 CANN 9.2.0 與 C++20 編譯器,9.2.0-beta.2 不提供後者;其二,昇騰 950 後端已併入 TileLang 主倉、而非只存在於單獨的昇騰適配倉,提供原生編譯、排程、同步與執行時整合,AutoSchedule 依據資料依賴與預估執行時間決定重疊與片上緩衝槽數,在 950PR 上對小尺寸單運算元最多縮短 40%,編譯器還會自動插同步、刪掉多餘等待,並檢查條件分支;其三,與 CANN、Triton-ascend 對比時「某些尺寸可以打平或超過 CANN」,最優實現隨尺寸切換(作者註明這是熱快取下的樸素基準)。文章同時點出生態缺口:多數最佳化面向超節點,x86 主機上 vLLM-Ascend 0.26.0rc1 仍有多處問題。

來源:https://newsletter.semianalysis.com/p/huawei-deepseek

2.2 主倉昇騰:DeepGEMM 啟發式與 epilogue 對齊由在審轉合入(#3577)

合於 10-21 10:20,29 檔案 +1392/-1036。上一期記為在審,本期完成合入:補齊 DeepGEMM-Ascend 的配置選擇與 epilogue,覆蓋密集、批次與分組 GEMM 的融合 bias、Alpha、BF16 隨機舍入與 FP8 量化,並在作者自己的提交裡刪掉冗餘 setup、恢復測試門控。這筆與 2.1 的外部判斷互相印證——昇騰側反覆打磨的正是 GEMM 路徑的啟發式與數值收尾,而這恰是「可用」與「跑得好」之間的差距所在。

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

2.3 主倉昇騰:warp_reduce_max/min 在 PTO 上的整數符號性修復(#3582)

合於 10-21 10:21。PTO 的 pto.redux_max/min 拒絕帶型別的整數:PTOAS 要求無符號性載體加一個顯式 signedness,於是 int32/uint32 的 T.warp_reduce_max/min 直接編譯失敗,報「requires signedness=’signed’ or ‘unsigned’ for integer values」。修復把 warp_reduce_sum/max/min 統一走 contrib.ptodsl 的 tl.redux_* 包裝(現已匯出),由包裝提供載體與符號性,並按 AscendAllReduce 的方式對 value dtype 設門;問題單引用了主倉 #3541 與 PTOAS #1579。驗證分兩層:op × dtype 網格的 lower-only codegen 測試,以及稀疏的雙目標(ascend/pto)執行測試——用 ≥2^31 的 uint32 lane 把符號性寫錯直接暴露出來。

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

2.4 昇騰倉:亞 32B 緊湊行距源無法表達,keepdim reduce 輸出錯亂(#1795)

合於 10-21 17:15,2 檔案 +302/-3,修 #1682,是本視窗技術含量最高的一筆。根因在模板層:copy_ub_to_gm 無條件生成一條多 burst 的 DataCopyPad,而 MTE3 引擎實現的是系統級「UB 窄行按 32B 槽存放」約定——源地址每 burst 前進 ceil(blockLen/32) + srcStride 個 32B 塊。對按 32B 邊界寫行的生產者這是精確匹配;但 T.reduce_max 的 (M,1) fp32 輸出是 V-pipe 緊湊寫入、行距只有 4B,引擎無法表達。實測第二個 burst 讀到的是 element 8 而不是 element 1([127, 1151, 71, 199] 裡 1151 恰是 element 8 的值),寫入側沒問題,錯在 DMA 讀源的地址序列。修復分兩半:兩側皆 packed 且 total ≤ 65535 時摺疊成一條 flat 單 burst(M×L 位元組,利用「單條 burst 長度任意」的契約);不可表達的非 32B 倍數行距改為 trap() 顯式中止,用編譯期 if constexpr 守衛,對 32B 倍數與亞 32B 槽佈局不生成額外程式碼。為什麼不降級:bisheng 的跨管道自動同步只覆蓋直線 intrinsics,用條件分支包裹 DataCopyPad、或兄弟分支裡放 GM 標量寫,都會讓 bisheng 生成壞的同步程式碼並導致裝置死鎖(實測逐行 DMA 被執行時分支包裹即觸發 aicore 異常)。14 個用例覆蓋 flat/multi 路徑、keepdim reduce 復現、UB 迴圈複用、不可表達 pitch 的合法單 burst 形式與顯式中止;測試環境為 Ascend910 A3(dav-c220)/ CANN 9.0.0 本地構建。PR 自述保留兩個邊界:亞 32B 緊湊行距且帶行距目的(如把 (M,1) fp32 寫進寬張量的某一列)仍走槽步進;blockCount > 4095(12 位欄位)仍會靜默截斷(4200 行實測只拷前 104 行),擬另立 issue。

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

2.5 昇騰倉:JIT 快取改原子釋出,並在快取被刪除時恢復(#1903)

合於 10-21 16:20,7 檔案 +1026/-86,關聯 #1836。原實現把編譯產物原地覆蓋到 kernel_lib.so,其他程序若已載入該共享庫,就會讀到被改寫中的檔案。修復先把產物複製到同目錄下的獨立臨時檔案,再用 os.replace() 原子釋出;載入前若快取庫已被刪除,經過錯誤訊息與檔案狀態確認後返回 cache miss,重新編譯;儲存階段目錄消失則跳過儲存、保留已經得到的編譯或調優結果;許可權錯誤、依賴庫缺失、無效 ELF 與序列化內部異常繼續向上傳播,不被吞成 miss。覆蓋面包括 KernelCache、AutoTunerCache 與 AutotuneResult,不改 cache key、運算元計算邏輯與 CI 並行度。41 項 host-only 用例透過(原子釋出 9 項、快取刪除恢復 32 項),PR 自述未做 NPU 側整合驗證。原子釋出部分沿用 #1853 的實現並保留原作者資訊。

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

2.6 昇騰倉:cann-bench 補 depthwise_conv_2d 示例(#1915)

合於 10-21 16:19。examples/cann-bench/ 新增 depthwise 二維卷積基準。附帶的冒煙測試細節值得記一筆:bench_test.sh 判定 PASSED 的標準是輸出裡出現 KERNEL OUTPUT MATCH 或 TEST PASSED!,指令碼原先列印的是「smoke test passed」,退出碼為 0 卻被判失敗;修復按 LSTM/GRU 的冒煙測試對齊,避免新增用例在 CI 裡持續假紅。

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

2.7 昇騰在審:indexer_topk_reducesum、GLM-5.3 kpool compress 與 WY fast 系列遷移

昇騰倉的在審佇列本視窗有 13 筆更新,集中在三組。#1917 實現 indexer_topk_reducesum 運算元,對應稀疏注意力的索引器取 top-k 後求和;#1910 增加 GLM-5.3 的 kpool 壓縮運算元,指向新一代模型的 KV 壓縮路徑;#1827 / #1828 / #1802 / #1803 / #1799 是一組 WY fast 與 chunk 系列運算元遷移(WY fast 本體及其反向、chunk_o 反向、Chunk Scaled Dot KKT、Chunk Delta H),明顯對著線性注意力/門控 delta 網路的訓練鏈。另有 #1785 修多維標量 BufferLoad 在 T.tile 二元運算裡的索引、#1854 校驗自動跨核同步配置、#1731 修 gemm_v0 的 W4A4 int4×int4 結果錯誤(#438)。

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

2.8 海光:feature 分支推進 FP8 解碼 GEMM,定位到地址鏈佔 kernel 13%

tilelang-hygon 的 main 分支視窗內沒有新提交(最新仍是 10-20 合入的 #25),但 feat/hcu-smls-b-vgpr 分支在 10-21 11:04 與 15:27 有三次推送,內容是可複用的效能定位與調參結論。11:04 那筆把解碼 GEMM 的 k-loop 展開,並讓 M>96 的 tile 回到 bn96;作者的關鍵判斷是「每步固定開銷來自地址鏈而不是 barrier」——符號化共享級索引 (k+d) % 3 會下降成 v_mul_lo_u16 / v_sub_u16 / v_readfirstlane_b32 並落在 matrix_load 的地址路徑上,佔掉 kernel 13%,而 T.Pipelined 本來就會發同樣的兩個 barrier。15:27 那筆把 M∈(64,128] 的 wide-N 中段優先按 block 數切分,gate_up 在 M∈(64,96] 取 bm96/bn128/bk128/ns2,並把 builder 移進包內的 decode_cfg;同時把解碼 GEMM 的記錄文件從 3825 行整理成 588 行的「結果 + 最佳化日誌」兩段式,另附已關閉槓桿索引與頻寬模型。

來源:https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-smls-b-vgpr

2.9 沐曦 / 摩爾執行緒 / MLIR-Ascend:視窗內無推送

tilelang-metax(預設分支 dev,最新推送 10-16,釋出仍停在 09-17 的 v0.1.14)、tilelang-musa(最新 09-30,v0.1.15+musa.1)、tilelang-mlir-ascend(最新 10-19,tag 停在 09-09 的 v0.1.2.020.sglang.poc)三倉本視窗均無提交,也無新 tag。

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


三、生態與採用方(TileOPs / TileFoundry / 運算元庫)

3.1 TileFoundry:0.0.3 打 tag 釋出,指令級排程 + 重構分析 + H200 與 Apple AMX 示例

上一期記錄「0.0.3 公告已進 README、尚待打 tag」,本期 10-21 10:26 正式釋出,PyPI 為 tilefoundry==0.0.3,代號「Agent-driven scheduling」。核心變化是排程權歸屬:編譯器列出每個 HIR 操作可由哪些指令實現、各自需要什麼,agent 在候選中選定,編譯器再把選定的排程下降為 TIR。配套 CLI 有 tilefoundry schedule candidates 與 schedule facts 展示候選與約束,tf.schedule 記錄選擇、tilefoundry schedule finalize 完成下降;指令以 AtomSched 操作宣告並帶顯式 TIR 運算元契約,從 CUDA MMA atom、SM90 WGMMA 與非同步複製起步。分析部分重構為記憶體(分段分配、footprint 與複用視窗、迴圈感知的 buffer 生命週期)、成本(按執行作用域計工作量與流量、每條訪問關係宣告讀什麼)、結構(定義事實、估值上下文與作用域快照統一走 visit context,投影複用使每個語料用例只分析一次)。新教程沿一個 8192×5120 乘 5120×17408 的 GEMM 一路走到 AtomSched 在 H200 上的最優 kernel。

效能面給出三組真實數字:Qwen3.8-27B 預填充在單張 H200 上,1024 / 4096 / 8192 token 分別 80.29 / 309.61 / 631.19 ms,SGLang 為 89.06 / 317.11 / 638.25 ms;短序列 128 / 256 / 512 / 1024 token 分別是 SGLang 的 1.88× / 1.41× / 1.10× / 1.03×;Apple M2 Pro 上的 Qwen 預填充示例把大投影做成 FP16 AMX 外積。兩組 Qwen3.8-27B 執行在每個長度上都與 transformers 產生相同首 token。已知問題兩項:wheel 尚未打包 TileFoundry 的 CUDA 執行時標頭檔案,從安裝包出發 tf.compile 編不出 CUDA kernel(需原始碼 checkout,下個版本修),以及效能時間線對 insert_slice / tf.zeros 跨 mesh 位置的定價有誤且不顯示 mesh lane 欠填(#239)。

來源:https://github.com/tile-ai/TileFoundry/releases/tag/v0.0.3

3.2 TileOPs:MLA 預填充反向補齊 op 層與 SM90 warp 特化 kernel(#2678 / #2680 / #2688 / #2686)

四筆連成一條線。#2678(10-21 11:34)加上 MLADenseBwdOp 與 MLAVarlenBwdOp 兩個 op 層,對齊 FA3 的 _flash_attn_backward 介面;#2680(10-21 20:58)補 SM90 kernel,針對 D_qk=192、D_v=128 給出 MLADenseBwdWSKernel 與 MLAVarlenBwdWSKernel 兩套 warp 特化實現;#2688(10-21 23:28)改 MLAVarlenPrefillWSFwdKernel 的常駐 CTA 取數順序——此前先給每個 head 領一個 query tile,首個 tile 走完之後只能按 head 逐個排隊,現在改成一次只走同一個 head 的 query tile;#2686(10-22 03:08)做收尾清理,把單讀者常量與 helper 移到使用處,不改行為。文件站 #78 同步把兩個新 op 寫進 Attention 頁,起因是 check_api_pages.py 報主倉匯出 207 個 op 中有兩頁缺失。

來源:https://github.com/tile-ai/TileOPs/pull/2678、https://github.com/tile-ai/TileOPs/pull/2680、https://github.com/tile-ai/TileOPs/pull/2688

3.3 TileOPs:dense GQA 支援 FP32 attention sink(#2515)

合於 10-21 22:42。為 dense GQA 的預填充與解碼增加可選的逐 head FP32 attention sink,覆蓋 SM90 上的 FP16/BF16/FP8、RoPE、softcap 與滑動視窗;其他架構對 sink 呼叫直接拒絕。語義上每個 sink 是分母裡一個不縮放、不加 softcap 的 logit:-inf 關閉該 sink、+inf 直接把該 head 輸出清零,split decode 在大 logit 下也要保住 sink 權重。順帶修了一處基準失真——分塊預填充的基準此前被當成單 query 解碼在測;融合 RoPE 的尾部用例把 KV 長度調到 271/1025,以便覆蓋拆分執行。

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

3.4 TileOPs:SM89 小規模路徑集中補位(#2681 / #2672 / #2669 / #2654)

四筆合入都在補 L40S/4090 級的短板。#2681 把 FP8 BMM 在 SM89 上的預設瓦片改成 64×64——條件是結果 grid 每 SM 至多一塊,大 grid 仍保留 128×128;作者給出 PyTorch 2.13.0 + CUDA 13.0 + TileLang 0.1.12 的單卡實測,1×16×64×128 從 5.161 降到 3.799 us、16×128×128×1024 從 12.943 降到 7.414 us,4096³ 反而從 512.573 變為 583.680 us(該行瓦片未變,屬噪聲)。#2672 把 64 列 W4A16 預設的 M 上界從 16 擴到 64,M=17 / 32 / 64 的重放時間分別下降 43.35% / 47.20% / 40.00%,同時保留顯式配置覆蓋與數值容差。#2669 用 MMA 實現 SM89 的分頁 DSA,保留 656 位元組 FP8 latent cache 佈局、BF16 舍入、自然對數 LSE 與 dead position 行為:batch 4 / 128 head / 256 位置 0.269 ms 對 PyTorch FP32 參考 0.727 ms,48×64 head×2048 位置 2.543 對 2.535 ms,作者明確寫「大用例基本打平,不構成普遍優勢」。#2654 針對小 token 數只計算並搬運覆蓋的 8-token 累加器字首,在 132 SM 的 H200、N 可被 128 整除、133–225 輸出 tile、K∈{7168, 7680, 8192} 的條件下把 40-token 解碼開啟到 M33–40。

來源:https://github.com/tile-ai/TileOPs/pull/2681、https://github.com/tile-ai/TileOPs/pull/2672、https://github.com/tile-ai/TileOPs/pull/2669、https://github.com/tile-ai/TileOPs/pull/2654

3.5 TileOPs:NSA top-k 在 SM89 上擴到 2048 壓縮塊並補齊輸出槽(#2674 / #2675 / #2673)

三筆圍繞同一處 SM89 路徑。#2674 讓 group48 複用既有 tiled top-k 程式:一個 token 一個 tile(48 個 query head 行)、列向 GEMM warp 佈局;#2675 把 SM89 上的 top-k 選擇從原來的上限擴到 1024–2048 個壓縮塊,覆蓋 head dim 64 / 128 / 256;#2673 修候選不足時的輸出填充——請求的每個 top-k 槽都要填 -1,此前 32 槽的排序帶會讓 32–63 號槽留著舊值。另有 #2684 給打包請求的 NSA 排序容量加上界,按 min(total_chunks, ceil(total_tokens / block_size)) 再對齊到打分 tile。

來源:https://github.com/tile-ai/TileOPs/pull/2674、https://github.com/tile-ai/TileOPs/pull/2675、https://github.com/tile-ai/TileOPs/pull/2673

3.6 TileOPs:GDN 單 kernel 預填充與分量化歸一化切分(#2682 / #2679)

#2682 直指線性注意力的 GDN 前向:其預填充在夜間跑批的 79 行上慢於 FlashInfer(0.57×–0.98×),改法是在「walk 能填滿 SM」的情形下把預填充收進單個 kernel。#2679 處理 FusedAddRMSNormPerTokenQuant:單元素約 25 條 SASS 指令花在逐元素數學上,改為把元素數學打包,並把寬行跨 cluster 切分。兩筆都是把「單運算元內序列開銷」換成「更寬的並行與更少的指令」,與同日 SM89 那組小規模路徑調優同一思路。

來源:https://github.com/tile-ai/TileOPs/pull/2682、https://github.com/tile-ai/TileOPs/pull/2679

3.7 TileOPs:GQA 反向在 warp 特化 kernel 裡直接服務分組 KV head(#2685)

合於 10-21 22:58。SM90、head dim 128 時,GQABwdOp 此前把分組 KV head 交給流水化 kernel,按 query head 逐個 CTA 走;現在改由 warp 特化的 GQA 反向 kernel 直接吃分組 KV head,減少重複搬運與排程開銷。同日 #2687 修 PrivateUse1 重新命名迴歸測試裡的裝置後端自動載入問題,屬於測試基礎設施的清理。

來源:https://github.com/tile-ai/TileOPs/pull/2685、https://github.com/tile-ai/TileOPs/pull/2687

3.8 TileOPs 在審:摺疊 edge 軸單次發射、NSA top-k 收窄列 tile(#2689 / #2690)

視窗末尾兩筆仍在審。#2689 針對 reduction:把邊緣軸摺疊進一次發射,並讓 fold 的尾部裝載先發;#2690 在 SM89 上給 NSA top-k 用更窄的列 tile。合入 20 筆之後,TileOPs 的在審佇列從上一期的 16 筆降到 9 筆,是本視窗組織內唯一在審數量明顯下降的倉庫。

來源:https://github.com/tile-ai/TileOPs/pull/2689、https://github.com/tile-ai/TileOPs/pull/2690

3.9 採用方:TileKernels 與 FlashQLA 視窗內無推送

deepseek-ai/TileKernels(1991★)與 QwenLM/FlashQLA(732★)本視窗均無提交,最新推送都停在 10-14;FlashQLA 的版本仍為 09-30 的 v0.1.3。採用方側的平靜已經持續一週,主線的動能仍在 tile-ai 組織內部。

來源:https://github.com/deepseek-ai/TileKernels、https://github.com/QwenLM/FlashQLA


四、社群、教程與活動

4.1 社群教程:中文《TVM / TileLang Cookbook》更新

RuneFang/tvm_tilelang_cookbook(34★,7-21 建倉)10-21 有推送。自我定位寫得很清楚:面向「能寫點 Python、聽過 CUDA 但沒讀過編譯器程式碼」的讀者,手把手從 @tilelang.jit 一路挖到 .cubin。這類把 DSL 到編譯產物全鏈路講透的中文材料仍然稀缺,對昇騰、海光等新後端的入門者尤其有用。

來源:https://github.com/RuneFang/tvm_tilelang_cookbook

4.2 社群新倉:tilelang-wrapper-for-sail 建倉,暫無說明

flytiger-eco/tilelang-wrapper-for-sail 於 10-21 09:04 建倉,無描述、無 README,倉名指向為 SAIL 提供 TileLang 包裝。這類「廠商內部執行時 + TileLang 外殼」的小倉通常是新硬體接入的前哨,但內容尚未公開,本期只作記錄。

來源:https://github.com/flytiger-eco/tilelang-wrapper-for-sail

4.3 社群工具:遠端自動調優 agent 倉繼續推送

cc-c122/TileLang-Remote-Autotuning-Agent(1★,7-01 建倉)10-21 22:57 有推送。與 TileFoundry 把排程交給 agent 的方向一致,社群側也在嘗試把自動調優做成可遠端驅動的 agent,而非本地指令碼遍歷。

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

4.4 文件與版本:文件站七次自動更新,主倉 VERSION 仍 0.1.15

tilelang.github.io 在視窗內自動更新七次(10-21 10:18、11:26、11:33、13:54、18:26、19:23 與 10-22 01:23),均隨主倉合入觸發。主倉 VERSION 仍為 0.1.15(該版本 09-30 釋出),README 的最新里程碑仍停在 09-30 的 Ascend 950 後端;TileOPs 無新 tag;TileFoundry 今日補上 v0.0.3 tag,是視窗內唯一的新版本。TileOPs-nightly 的 pushed_at 雖有變動,但其預設分支自 08-30 起無新提交,變動來自非預設分支,不計入開發活動。

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


五、趨勢觀察

編譯器缺陷的性質從「算錯數」擴到「選錯目標」。 本期主倉八筆合入裡,四筆改變的是「編譯器以為自己在做什麼」:無 arch 的 CUDA target 靜默退到 sm_50(#2860)、cast 的舍入註解按分支靜默丟棄(#3350)、批式歸約只處理首個輸出塊(#3070)、流水規劃丟失傳遞性依賴(#2822)。它們的共同點是不報錯,只在某些尺寸或某些迴圈形態下給出不可信結果——比崩潰更難定位,也說明專案把靜默正確性放在了功能擴張之前。

昇騰線的瓶頸已移到記憶體模型與執行時邊界。 亞 32B 緊湊行距無法用 DMA 表達、只能顯式中止(#1795);快取併發改為原子釋出並補刪除恢復(#1903);PTO 的整數符號性要靠包裝層補齊(#3582)。這三筆都不是「加運算元」,而是「讓已有運算元在邊界條件下不出錯」,配合 DeepGEMM 啟發式對齊合入,昇騰側正從能力覆蓋轉向工程可靠性。

分析機構開始把 TileLang 當作昇騰可用性的判據。 SemiAnalysis 用「六個倉庫中只有 TileLang 當時可用」來定義 DeepSeek 昇騰支援的完成度,這比任何單篇技術部落格都更能說明 TileLang 在國產算力敘事裡的位置:從實驗性 DSL 變成預設承載層。同一篇報告也給出反向訊號——最佳化集中在超節點、x86 主機支援仍不完整,通用性落後於效能。

注意力棧的重心從推理前向移到訓練反向。 TileOPs 一天內補齊 MLA 預填充反向的 op 層、SM90 warp 特化 kernel 與取數順序最佳化,再加上 GQA 反向的分組 KV head 處理,與昇騰側上月以來的反向運算元補位(indexer_bwd、SiTU-GLU 反向、WY fast 反向)完全同向。訓練側的 TileLang 覆蓋度正在被系統性補齊。

SM89 從「能跑」走向「調優」。 FP8 BMM 瓦片、W4A16 的 M 上界、分頁 DSA、NSA top-k 四條線同日動刀,且都給出單卡實測與「不構成普遍優勢」這類限定語。消費級與單卡推理場景成了效能目標之一,同時作者對基準口徑的自我約束明顯變嚴。

排程與正確性的分工在 TileFoundry 與主倉之間變得清晰。 主倉繼續收緊語義與正確性,TileFoundry 0.0.3 則把「選哪條指令」交給 agent、把「保證語義」留給編譯器,並用 H200 與 Apple AMX 兩個示例給出可復現數字。這條分工如果穩定下來,TileLang 生態會出現「語言層守正確、agent 層搏效能」的雙層結構。


附錄:素材與核查說明

信源核查表

信源 核查結果
tile-ai/tilelang(main) 視窗內 8 次提交、8 筆合入;VERSION 仍 0.1.15;在審 163 筆,其中 21 筆視窗內更新
tile-ai/TileOPs(main) 視窗內 20 次提交、20 筆合入;在審 9 筆,3 筆視窗內更新
tile-ai/tilelang-ascend(ascendc_pto) 視窗內 3 次提交、3 筆合入;在審 235 筆,13 筆視窗內更新
tile-ai/tilelang-hygon(main / feat-hcu-smls-b-vgpr) main 無提交;feature 分支 3 次推送(10-21 11:04、15:27);在審 3 筆
tile-ai/TileFoundry(main) 無新提交;v0.0.3 於 10-21 10:26 打 tag 釋出;在審 1 筆
tile-ai/tilelang-metax(dev)/ musa / mlir-ascend 三倉視窗內均無提交(最新分別 10-16、09-30、10-19),無新 tag
tile-ai/TileRT / tilescale / DeepStack 三倉視窗內均無提交(最新分別 09-24、08-25、10-05)
deepseek-ai/TileKernels / QwenLM/FlashQLA 均無提交,最新推送均停在 10-14
tilelang.github.io / TileOPs.github.io 分別 7 次 / 1 次自動更新,均隨主倉合入觸發
TileOPs-nightly 預設分支自 08-30 無新提交,pushed_at 變動來自非預設分支
Google News(中英多組)/ Hacker News / arXiv 視窗內三篇產業報道圍繞華為-DeepSeek 協同,僅 SemiAnalysis 一文含 TileLang 專章;arXiv 最新相關論文仍為 07-24

補充說明:海光線的推送不在預設分支,判定方法是先取 feat/ 分支的提交流,再用倉庫事件介面核對 PushEvent 的 ref 與 before→head;TileOPs-nightly 的 pushed_at 變動同樣按此法排除。在審數量取各倉 is:open 的檢索總數,視窗內更新數取 updated:>=視窗起點 的檢索總數。本視窗起點(10-21 07:00)之前半日內落地的提交由上期報告覆蓋,本期不重複計入。

完整信源清單

/TileOPs/pull/2680>

ek-ai/TileKernels>