TileLang 每日動態報告(2026-10-23)
調研視窗:2026-10-22 07:00 ~ 2026-10-23 07:00(北京時間,24 小時)。
信源:GitHub tile-ai 組織逐倉核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon 分支 / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 三個文件站)+ 合入與在審真值檢索(is:merged 與 is:open 兩組,按倉庫分組統計)+ 提交級 diff 與釋出頁 tag 核查 + 採用方倉核查(TileKernels / FlashQLA)+ 社群新倉檢索 + Google News 中英文多組查詢(走代理)+ Hacker News + arXiv + 產業媒體報道。
本期索引
- 今日重點:35 筆合入裡主倉把重心放到分析器生命週期,深尋 V3.2 的 top-k 選擇器越界寫被定位到 CUDA_ERROR_ILLEGAL_ADDRESS;昇騰 NZ 佈局打通部分 UB→L1 複製(18 檔案);CANN 社群新開 ops-tilelang 運算元倉
- 一、核心專案進展
- 1.1 主倉:向量化重寫前克隆 analyzer(#3612)(10-22)
- 1.2 主倉:並行迴圈切分後舊 analyzer 繫結失效(#3226)(10-22)
- 1.3 主倉:deepseek_v32 top-k 選擇器候選緩衝越界(#3280)(10-22)
- 1.4 主倉:CUDA 16 位反三角與二進位制數學降級補齊(#3172)(10-22)
- 1.5 主倉:sm8x 稀疏後設資料非雙射形狀改為拒絕(#3125)(10-22)
- 1.6 主倉:eager 前端保留 Python 原始碼位置(#3558)(10-22)
- 1.7 在審:pass 能力檢查、無符號 mulhi 與 assume 謂詞(10-22)
- 二、多後端適配(昇騰 / 海光 / 沐曦 / 摩爾執行緒)
- 2.1 昇騰:NZ 緩衝支援部分 UB→L1 複製(#3530)(10-22)
- 2.2 昇騰 PTO:MAD 方向謂詞指數增長致編譯停滯(#3581)(10-22)
- 2.3 昇騰 PTO:SIMT 向量區域性變數按 lane 分配(#3523)(10-22)
- 2.4 昇騰倉:cann-bench 補 Dilation2D,同日回退 depthwise_conv_2d(#1916 / #1919)(10-22)
- 2.5 昇騰倉在審:sparse_mla 前向與反傳(#1929 / #1928)(10-22)
- 2.6 海光:合併分支移除 enable_register_pipeline(10-22)
- 2.7 沐曦 / 摩爾執行緒 / MLIR-Ascend:視窗內無推送(10-22)
- 2.8 產業側:SCMP 轉載 SemiAnalysis 報告,TileLang 為六倉中唯一可用(10-22)
- 三、生態與採用方(TileOPs / TileFoundry / 運算元庫)
- 3.1 TileOPs:FP8 分頁 warp 特化 kernel 共享記憶體別名致 NaN(#2708)(10-22)
- 3.2 TileOPs:dense prefill 改為持久化遍歷,D64 三消費者(#2706)(10-22)
- 3.3 TileOPs:SM90 dense GQA decode 走 warp 特化(#2694)(10-22)
- 3.4 TileOPs:因果滑窗用一個 KV head 的兩個 key tile 服務整組 query head(#2696)(10-22)
- 3.5 TileOPs:分頁 decode 自動調優此前共用一張 block table(#2703)(10-22)
- 3.6 TileOPs:masked softmax 的 FMA 溢位與 varlen RoPE 奇數半寬(#2699 / #2700)(10-22)
- 3.7 TileOPs:FP8 grouped GEMM 分級暫存 scale 與 gate/up 乒乓 warp group(#2704 / #2693)(10-22)
- 3.8 TileOPs:謂詞廣播與布林運算元改 16 位元組向量讀(#2702 / #2697)(10-22)
- 3.9 TileOPs:SM89 三條路徑同日推進(#2705 / #2691 / #2690)(10-22)
- 3.10 TileOPs:歸約摺疊單次發射與空 DSA 選擇校驗(#2689 / #2698)(10-22)
- 3.11 TileOPs:基準參考改在 float32 上計算(#2692)(10-22)
- 3.12 TileFoundry:安裝 sm_89 架構與 RTX 6000 Ada 裝置(#259)(10-22)
- 3.13 TileFoundry:gmem 計價改按足跡與執行 mesh,並繫結被調方維度(#258)(10-22)
- 3.14 TileFoundry:裝置拒絕啟動 kernel 時必須拋錯(#256)(10-22)
- 3.15 分發版:Sunrise 打出 0.1.15+sunrise.1.2.0 釋出 PR(#9,在審)(10-22)
- 3.16 採用方:TileKernels 與 FlashQLA 視窗內無推送(10-22)
- 四、社群、教程與活動
- 4.1 CANN 社群開設 ops-tilelang 運算元倉(10-22)
- 4.2 tilelang/agent:把運算元最佳化寫成 agent 的 SKILL 工作流(10-22)
- 4.3 社群新倉:Sail 包裝、遠端自動調優 agent 與教程倉(10-22)
- 4.4 文件與版本:文件站更新三次,VERSION 仍 0.1.
15(10-22)
- 五、趨勢觀察
今日重點
本視窗 tile-ai 組織合入 35 筆(TileOPs 20、tilelang 9、TileFoundry 3、tilelang-ascend 2、TileFoundry 文件站 1),另有 71 筆在審 PR 在視窗內更新。主倉九筆的重心從上週的”形狀與幾何校驗”轉向分析器生命週期:向量化重寫前克隆 analyzer(#3612,單檔案 2 行)、並行迴圈切分後舊繫結失效(#3226,6 檔案)、eager 前端保留 Python 原始碼位置(#3558)——三筆都不會讓 kernel 算錯,而是讓編譯期在改寫後讀到過期狀態。數值側最值得記的是 #3280:deepseek_v32 的 top-k 選擇器把 radix 候選暫存在固定 [2, 4096] 共享緩衝,卻沒有約束第一輪選出的閾值桶大小,低尾數相同的 FP32 分數或全等輸入會越界寫,崩在 CUDA_ERROR_ILLEGAL_ADDRESS。昇騰側 #3530 是當日體量最大的一筆(18 檔案 +993/-256),把壓縮 NZ 填充保留到儲存歸一化,為部分 UB→L1 複製顯式給出 C0 組行距;#3581 修掉 PTO 上一個會讓編譯停滯十分鐘以上的謂詞表示式指數增長。生態側兩條新線:CANN 社群 10-20 建了 ops-tilelang 運算元倉、10-22 首次推送修復;Sunrise 分發版打出 0.1.15+sunrise.1.2.0 的釋出 PR,把 TANG 後端 wheel 納入分發。
一、核心專案進展
1.1 主倉:向量化重寫前克隆 analyzer(#3612)
合於 10-22 23:43,只改向量化那一處原始碼,+2/-1。迴圈向量化在改寫 IR 時複用傳入的算術分析器,而分析器快取著改寫前的常量與區間約束;一旦向量化改變了迴圈邊界或索引表示式,後續 pass 拿到的仍是舊結論。修法是在進入向量化重寫前複製一份分析器,改寫只作用於副本。這一筆沒有隨附迴歸測試,價值在於它從另一個入口堵住與 #3226 同源的問題——同一天那一筆處理的是並行迴圈切分後的失效繫結,兩筆合看,分析上下文的生命週期被當成獨立缺陷類在清。
來源:https://github.com/tile-ai/tilelang/pull/3612
1.2 主倉:並行迴圈切分後舊 analyzer 繫結失效(#3226)
6 檔案 +55/-19,涉及 loop_partition、loop_vectorize、lower_tile_op 三處 pass 及其標頭檔案。並行迴圈被切分後迴圈體的迭代變數被重繫結,但 pass 仍按切分前的變數查分析器,取到的是舊變數的區間與整除資訊,簡化與越界判定因此失真。修復把分析上下文按 pass 隔離,並簡化了迴圈分析快照的初始化路徑,配套加了一組最小化的並行迴圈分析迴歸用例。提交者中含 Codex 聯署。
來源:https://github.com/tile-ai/tilelang/pull/3226
1.3 主倉:deepseek_v32 top-k 選擇器候選緩衝越界(#3280)
2 檔案 +91/-10,修 #1351。tl_topk_impl 把 radix 候選裝進固定的 [2, 4096] 共享緩衝,但沒有任何東西約束第一輪 radix 選出的閾值桶不超過 4096 個元素——FP32 分數低位尾數相同、或輸入全等時,原子追加會寫出緩衝末尾並崩在 CUDA_ERROR_ILLEGAL_ADDRESS。改法是在 tl_topk 包裝層申請 (batch, 2, seq_len) 的全域性 workspace,讓兩個 radix 緩衝都由 seq_len 兜住;公開的 tl_topk() 簽名不變。隨附 79 行復現用例,按 #1351 的輸入構造。
來源:https://github.com/tile-ai/tilelang/pull/3280
1.4 主倉:CUDA 16 位反三角與二進位制數學降級補齊(#3172)
3 檔案 +463/-8。asin、acos、asinh、acosh、atanh、log1p、atan2、hypot、copysign、ldexp、nextafter 缺 CUDA 專用下降規則,通用規則只處理 float32/float64,float16 與 bfloat16 呼叫在 codegen 階段無人接管。修復把它們註冊進 CUDAMath 並補 half_t / bfloat16_t 過載:十個運算元按 float32 求值再轉回,nextafter 改為按 16 位編碼做整數比較步進,以便在 fast math 下保住 bfloat16 次正規數。float32/float64 的函式名保持不變,common.h 側新增 182 行模板實現。
來源:https://github.com/tile-ai/tilelang/pull/3172
1.5 主倉:sm8x 稀疏後設資料非雙射形狀改為拒絕(#3125)
單檔案 tilelang/layout/gemm_sp.py +6/-0,修 #2606。make_cutlass_metadata_layout_sm8x 的 interleave 只在 M 為組大小整數倍且 K 為偶數時才是雙射;不滿足時原先靜默生成錯誤後設資料,稀疏 GEMM 結果無提示地錯。現在兩個條件都改為顯式拒絕。原先隨 PR 加的稀疏後設資料形狀測試在評審中被要求移出,未進入主分支。
來源:https://github.com/tile-ai/tilelang/pull/3125
1.6 主倉:eager 前端保留 Python 原始碼位置(#3558)
4 檔案。eager 前端把 Python 的 for / with 語句降成生成式呼叫後,報錯與除錯對映回的是生成程式碼的行號。修復把生成的迴圈呼叫對映回 for 語句、上下文包裝對映回 with 語句,並補上可迭代引數解包的位置;與 #3125 一樣,評審要求移出隨附的原始碼位置迴歸測試。
來源:https://github.com/tile-ai/tilelang/pull/3558
1.7 在審:pass 能力檢查、無符號 mulhi 與 assume 謂詞(10-22)
主倉在審佇列在視窗內更新 50 筆。值得跟的幾條:[Transform] Locate VerifyBufferInit reads outside tile ops(#3620)、[Codegen] Add unsigned mulhi intrinsic(#3617)、[Transform] Keep assumption checks with local dependencies in scope(#3611)與 [Transform] Preserve tl.assume predicates that simplify to true(#3610)、[CPU] Support clz in C and C host code generation(#3609);昇騰側 [BugFix][Ascend] Reject odd FP4 view geometry in GM->L1 copies(#3608)與 [Ascend] Support exact placement of versioned on-chip buffers(#3607)同日更新。另有 #3613 明確標註”請勿評審”的原生 argmin/argmax 歸約草案。
來源:https://github.com/tile-ai/tilelang/pulls
二、多後端適配(昇騰 / 海光 / 沐曦 / 摩爾執行緒)
2.1 昇騰:NZ 緩衝支援部分 UB→L1 複製(#3530)
當日體量最大的一筆:18 檔案 +993/-256。壓縮 NZ 佈局的填充此前在邏輯佈局階段就被抹平,做部分 UB→L1 複製時無法表達源與目的各自的 C0 組行距。修復把緊湊 NZ 填充保留到儲存歸一化階段、再按容量精確物化一次,複製區域下降時顯式給出以 32 位元組為塊的源/目的 C0 組行距,並把行距貫穿到 scatter 模板、Ascend/PTO 的 DMA codegen 與訪問足跡;同時加了幾何、對齊與越界校驗,保留物理檢視所有者,呼叫方遷移到邏輯分配形狀加顯式行距。src/ascend/op/nd2nz.h 新增 83 行。
來源:https://github.com/tile-ai/tilelang/pull/3530
2.2 昇騰 PTO:MAD 方向謂詞指數增長致編譯停滯(#3581)
2 檔案 +48/-5。帶重複可變守衛的 PTO kernel 可以花十分鐘以上做分析,即使裡面沒有矩陣乘——這是 NPU 執行之前的編譯期停頓。根因是 PTOMmadDirectionAnalyzer 在守衛讀取可變記憶體時保守保留進入狀態,若兩條分支都保持 MAD 方向,Merge 收到同樣的謂詞 S,而實現總是構造並化簡 S || S;算術化簡器在復現用例裡消不掉這個重複,連續分支便讓謂詞表示式按倍數翻番。修復在兩側謂詞相同時跳過重建,並補 20 行定向測試。
來源:https://github.com/tile-ai/tilelang/pull/3581
2.3 昇騰 PTO:SIMT 向量區域性變數按 lane 分配(#3523)
4 檔案 +273/-29。SIMT 向量區域性變數被當標量 PTODSL 緩衝分配,沒算 lane 數,卻用 Python 下標訪問;AllocatedBufferValue 不支援逐項賦值,迴圈索引寫入在 PTODSL AST 重寫階段就更早失敗。修復按每個向量槽分配足夠標量元素並把槽下標翻譯成標量偏移,複用既有的連續 load/store 路徑,保持 AscendC 的向量區域性語義,不改迴圈展開、不加新的下降 pass。補 float32x2 / float32x4 兩組 codegen 測試,共 209 行。
來源:https://github.com/tile-ai/tilelang/pull/3523
2.4 昇騰倉:cann-bench 補 Dilation2D,同日回退 depthwise_conv_2d(#1916 / #1919)
#1916 合入 Dilation2D 前向(cann-bench level3 任務):NHWC、fp16/fp32、任意 kernelSize/stride/rate 與 SAME/VALID,越界 tap 的貢獻取 -inf + filter 與 golden 逐位一致,NaN 輸入傳播、+inf filter tap 使越界貢獻變 NaN;實現為純 Vector kernel 加 Python wrapper,按昇騰複製引擎的 32B 行跨距對齊約束派發,grid 按幾何自適應而不用寫死核數。同一天 #1919 把 10-21 剛合入的 depthwise_conv_2d(#1915)整體回退,提交資訊未給原因,示例倉處在”進一個、退一個”的狀態。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1916
2.5 昇騰倉在審:sparse_mla 前向與反傳(#1929 / #1928)
#1929 新增 sparse_mla_fwd(DSA 稀疏 MLA 前向:varlen packed 隨機行 gather、線上 softmax、QK/PV 雙 GEMM 融合),落 examples/dsa_sparse_finetune/,雙 kernel(fast 特化加 portable 通用)配合 host 分派;fast 路徑的適用域寫死為 H=128 且 G=1、dim/tail 為 512/64、topk 為 256 的整數倍、NI 在 1 到 6 之間,域外走通用路徑。#1928 新增 sparse_mla_bwd(dQ GEMM 鏈、P/dP 向量鏈、dKV 隨機行散射),落 examples_experiment/dsa/ 並登記進運算元測試清單,前向結果 k1 已折進本運算元。昇騰倉在審合計 11 筆。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1929
2.6 海光:合併分支移除 enable_register_pipeline(10-22)
gcy_merge_main 於 10-22 11:11 推送一筆 6 檔案 +10/-24 的清理,把 enable_register_pipeline 開關整體刪掉,涉及 HCU 暫存器軟流水的注入與流水規劃兩個 pass、pipeline_utils.h 以及語言層的 loop 擴充套件,示例與效能指令碼同步改呼叫。刪開關通常意味著暫存器流水路徑已經預設生效、雙軌維護不再划算;但 main 分支仍停在 10-20 的 preserve constant K origins in hoisted MLS windows(#25),這批改動還在合併分支上等落地。
來源:https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main
2.7 沐曦 / 摩爾執行緒 / MLIR-Ascend:視窗內無推送(10-22)
tilelang-metax 最後推送停在 10-16,tilelang-musa 停在 09-30,tilelang-mlir-ascend 停在 10-19,三條適配線在視窗內均無提交。昇騰、海光兩條線承擔了本視窗全部後端改動。
來源:https://github.com/tile-ai/tilelang-metax
2.8 產業側:SCMP 轉載 SemiAnalysis 報告,TileLang 為六倉中唯一可用(10-22)
南華早報 10-22 發文《Jensen’s nightmare》,轉述 SemiAnalysis 週四釋出的華為-DeepSeek 報告:DeepSeek 最新軟體釋出加上對昇騰 950PR 的更多實驗,說明與美國的差距在縮小,並判斷 DeepSeek 到 2027 或 2028 年將”主要”依賴華為”已相當成熟”的技術棧。該報告原文即上一期已收錄的那份,其中”六個獲得昇騰支援的倉庫中只有 TileLang 可用”的判斷未變,另提到昇騰 950 後端是併入 TileLang 主倉而非停留在獨立適配層。屬同一信源的二次傳播,不構成新的技術進展,但說明該判斷正在進入主流媒體口徑。
三、生態與採用方(TileOPs / TileFoundry / 運算元庫)
3.1 TileOPs:FP8 分頁 warp 特化 kernel 共享記憶體別名致 NaN(#2708)
2 檔案 +80/-30。GQAPagedFP8WSKernel 在”七個單 token 請求排在 300 token 請求之前、4099 行 FP8 cache、48 行頁”的組合下返回 NaN。根因是共享記憶體規劃器把 warp 角色各自的緩衝讀成不同時段活躍,於是把 split handover 放到了仍被 TMA 填充的 key 或 value 環地址上,handover 的 running maxima 也被壓到環標誌位。修復把這幾塊緩衝在共享記憶體裡顯式隔開,並補 24 行迴歸用例(qwen3-fp8-ragged 口徑)。
來源:https://github.com/tile-ai/TileOPs/pull/2708
3.2 TileOPs:dense prefill 改為持久化遍歷,D64 三消費者(#2706)
6 檔案 +590/-5,新增 GQADensePersistentWSKernel(對外名 gqa_dense_persistent,優先於 gqa_dense):query tile 數超過 SM 數時,一個 CTA 佔一個 SM 交替輪次走完自己的 tile,並在當前 tile 的 epilogue 期間裝載下一個 tile 的 query。D64 上每個 tile 是 192 行 query 攤給三個 64 行消費者 warpgroup,K/V 重讀減少三分之一;D128 保持 128 行。
來源:https://github.com/tile-ai/TileOPs/pull/2706
3.3 TileOPs:SM90 dense GQA decode 走 warp 特化(#2694)
新增 GQADecodeWSKernel:TMA 生產者 warp 填兩槽 K/V 環,一個 WGMMA warpgroup 維持 head 組的線上 softmax;split 數按一次填滿常駐 CTA 槽來定,並保證每個 split 至少四個 KV tile,單 split 時在 CTA 內歸一化並直接寫輸出、不再過 combine kernel;KV head 在 CTA 間變化最快。
來源:https://github.com/tile-ai/TileOPs/pull/2694
3.4 TileOPs:因果滑窗用一個 KV head 的兩個 key tile 服務整組 query head(#2696)
新增 GQADenseSlidingWindowBandKernel:GQADenseFwdOp 原先對每個 query head 都過載同樣兩個 key/value tile,一個 CTA 只做一個(query tile, head);新 kernel 一個 CTA 對(query tile, KV head)把 band 的兩個 key/value tile 只裝一次,依次對該組 query head 做歸約,同時讓下一個 head 的 query tile 在飛。針對的是滑窗行數不超過 128 個 key 的場景。
來源:https://github.com/tile-ai/TileOPs/pull/2696
3.5 TileOPs:分頁 decode 自動調優此前共用一張 block table(#2703)
4 檔案 +24/-13。分頁 decode 的 autotune 供給讓每個請求拿到同一張 block table,等於只測了一個請求的頁、且其頁常駐 L2,候選排名反映的是被服務批次實際付出的流量。實測在 1500 MHz H200 上,MHA decode 調優器在 opt-1.3b-b16 上選了 block_N 128(冷 80.0 us)而放棄 block_N 64(72.5 us),掉到 flashinfer 的 0.994。修復讓調優按請求各自讀不同 cache 頁。
來源:https://github.com/tile-ai/TileOPs/pull/2703
3.6 TileOPs:masked softmax 的 FMA 溢位與 varlen RoPE 奇數半寬(#2699 / #2700)
#2699(4 檔案 +54/-6)修 dense GQA 滑窗前向的 NaN 行:因果窗下的 128x64 調優瓦片讓每個 query tile 從首行視窗開始走,最後幾行會連續遇到兩個全遮蔽 key tile;第一塊之後兩個 running maxima 都在 -1e38 鉗位上,fast math 把 prev * scale - curr * scale 融合成 FMA,舍入誤差(約 2 的 99 次方)溢位成 NaN,修法是在縮放重定標前先減 running maxima。#2700 修 varlen query RoPE 在奇數半旋轉寬度(rotary_dim 10、14 等)下的迴圈邊界:並行度 rows 乘 chunks 不能被旋轉分支的執行緒數整除時 TileLang 丟掉了迴圈上界,warp 特化路徑的後續 query 行算錯,通用路徑(dim 256)直接掛死。
來源:https://github.com/tile-ai/TileOPs/pull/2699
3.7 TileOPs:FP8 grouped GEMM 分級暫存 scale 與 gate/up 乒乓 warp group(#2704 / #2693)
#2704 新增配置鍵 stage_scales:讓生產者用 cp.async 把每一級的 A/B scale 拷進共享記憶體並在落地時到達該級 barrier;stage_scales=0 仍生成舊 kernel。背景是 masked FP8 grouped GEMM 在數學 warpgroup 裡從全域性讀每步 K 的 A scale,環加深會壓縮 L1 並讓這些 load 停住,ds-v3-e8-up-masked 因此只有 DeepGEMM 的 0.979。#2693 給 GemmTemplate 加乒乓形態:兩個數學 warpgroup 各跑完整的 64 行 tile、輪流取 CTA 的 tile,一個組的 epilogue 落在另一個組的 mainloop 下,靠輪轉 barrier 排序。
來源:https://github.com/tile-ai/TileOPs/pull/2704
3.8 TileOPs:謂詞廣播與布林運算元改 16 位元組向量讀(#2702 / #2697)
#2702 讓行廣播謂詞(Eq/Ne/Gt/Lt/Ge/Le 與 LogicalAnd/Or)把每 1024 列整塊的 stride-1 運算元按 16 位元組向量經 load16_evict_last 讀,部分塊與其他 staged body 保留 fragment 複製;linear-bias-float32(1024x4096 對 1x4096)六種比較從 0.99 提到 1.06,邏輯運算元 1.04。#2697 把比較與邏輯的二值 kernel 中全尺寸運算元改為 16 位元組 evict-first 向量讀,float32 行此前只在 torch.compile 的 0.99 到 1.00。
來源:https://github.com/tile-ai/TileOPs/pull/2702
3.9 TileOPs:SM89 三條路徑同日推進(#2705 / #2691 / #2690)
#2705 把 SM89 的 SSD chunk scan 預設流水級數降到兩級,減少 scan kernel 的緩衝需求,六個形狀的 FP16/BF16 實測取三輪、每輪 200 次 CUPTI 取樣的中位數。#2691 在現有規劃器選擇多分割槽時拆分 SM89 的 paged DSA key 區間,複用既有 FP32 部分輸出與 log-sum-exp 的 merge;單分割槽時保持 BF16 輸出與單 kernel,並把編譯目標納入 merge 快取標識、為每個目標建獨立 JIT builder。#2690 在 SM89 上對 query head 組 32/64、head dim 64/128 的 NSA top-k 選擇改用 32 列 tile,保留 64 行 tile 與同一套打分排序實現,八個 FP16/BF16 用例的 GPU busy 中位數降 2.1% 到 15.0%。
來源:https://github.com/tile-ai/TileOPs/pull/2691
3.10 TileOPs:歸約摺疊單次發射與空 DSA 選擇校驗(#2689 / #2698)
#2689 新增 ReduceFoldKernel,把每個保留索引的 lead 段連續字尾在一次發射裡摺疊(此前對軸字首與字尾的歸約要兩次發射),並讓 ReduceEdgeKernel 與 VectorNormEdgeKernel 把行能被 512 執行緒塊在飛持有的呼叫都交給它;背景是 Sum 在 3d-multidim-float16(0.977)與 mlp-intermediate-float32(0.990)上落後 torch-compile。#2698 修分頁 DSA 校驗:query 行沒有任何有效選中 key 時參考實現返回零輸出與負無窮 LSE,而減參考 LSE 會得到 NaN 誤差界、把正確結果判錯;修復給空行精確的零 LSE 界,並補”全部行皆空”時校驗器必須拒絕非零輸出與有限 LSE 的控制用例。
來源:https://github.com/tile-ai/TileOPs/pull/2689
3.11 TileOPs:基準參考改在 float32 上計算(#2692)
ref_program 此前既當正確性參考、又是 torch 基線計時的物件,許多參考直接以 16 位輸入 dtype 調 torch,結果取決於裝置側 torch 每一步怎麼舍入。CUDA 掩蓋了這個差異,CPU torch 與 torch_musa 不掩蓋——GroupNorm 與 InstanceNorm 在 MTT S5000 上失敗,而同一 kernel 對 float64 參考是透過的。修復讓 ref_program 在 float32 或更寬上計算,torch 基線則按輸入 dtype 計時。這是少見的一筆由國產 GPU 後端暴露出、但要先在參考實現層修的問題。
來源:https://github.com/tile-ai/TileOPs/pull/2692
3.12 TileFoundry:安裝 sm_89 架構與 RTX 6000 Ada 裝置(#259)
11 檔案 +392/-22。新增 nvidia_sm89.toml 與 nvidia_rtx_6000_ada.toml 兩個目標描述(分別 114 與 124 行),docs/spec/target.md 同步補 42 行。這一筆把 TileFoundry 的裝置建模從 Hopper 擴到 Ada,與同日 #256 暴露的”按 sm_90 目標編譯的測試在 sm_89 單卡上只錯值”是同一件事的兩面;配套還加了 sm_89 的 e4m3 warp mma 宣告(#261,在審)。
來源:https://github.com/tile-ai/TileFoundry/pull/259
3.13 TileFoundry:gmem 計價改按足跡與執行 mesh,並繫結被調方維度(#258)
當日改動面最廣的一筆:74 檔案 +1128/-553,修 #239、#252、#253。原先 132 個 CTA 只探 40 個時,計價按”每個視窗每個位置移動一次”算,under-fill 根本看不出來;只能由被調方宣告的維度綁不上;analyze 的輸出目標與逐函式記憶體行含義不清。修復把 check 與 analyze 的維度繫結集中到 instantiate,被調方或子模組宣告的維度現可繫結。提交由社群貢獻者發起、維護者接手改寫。
來源:https://github.com/tile-ai/TileFoundry/pull/258
3.14 TileFoundry:裝置拒絕啟動 kernel 時必須拋錯(#256)
7 檔案 +69/-11。此前啟動器發出 kernel<<<grid, block, smem, nullptr>>>(...) 後從不讀啟動狀態,裝置拒絕啟動時呼叫照常返回,輸出保持舊內容,錯誤到不了呼叫方;在 sm_89 單卡上,按捆綁 sm_90 目標編譯的測試只表現為值不對,compute-sanitizer 才顯示 cudaLaunchKernel 返回 cudaErrorNoKernelImageForDevice。修復讓 CUDA 啟動器在讀 cudaGetLastError() 後拋錯,CPU 側 module 同步改,規格文件與測試一併更新。
來源:https://github.com/tile-ai/TileFoundry/pull/256
3.15 分發版:Sunrise 打出 0.1.15+sunrise.1.2.0 釋出 PR(#9,在審)
10-22 16:48 建 candidate/sunrise-0.1.15-1.2.0-20261019 分支並開 PR,把 Sunrise 分發版更新到 0.1.15+sunrise.1.2.0,納入 TileLang 0.1.15 及配套評審過的 TileOPs、TileKernels、TVM、tvm-ffi 變更;源匯入與合併、分發適配分開提交,最終適配配置 TANG 後端的 Sunrise wheel 並修掉切片 GEMM 迴歸(兩個累加器標誌都對公開模板校驗,保留 stride、origin 與 offset 斷言)。自報驗證:本機契約測試 55 透過、TANG 流水線 codegen 測試 54 透過、兩個 lint job 與信任邊界檢查透過;但 S2 wheel 構建出包後,裝後 import 冒煙檢查在 tangGetDeviceCount 處失敗。分支上三筆分別是 0.1.15 分發適配、CI 收窄到 TANG 後端與切片 GEMM 累加器標誌測試。
來源:https://github.com/tile-ai/tilelang-sunrise/pull/9
3.16 採用方:TileKernels 與 FlashQLA 視窗內無推送(10-22)
deepseek-ai/TileKernels 最後推送停在 10-14,QwenLM/FlashQLA 同樣停在 10-14,兩家採用方在本視窗均無提交。
來源:https://github.com/deepseek-ai/TileKernels
四、社群、教程與活動
4.1 CANN 社群開設 ops-tilelang 運算元倉(10-22)
Atlas-Ascend-Fork/ops-tilelang 於 10-20 建倉,定位是 CANN 社群面向昇騰 NPU 的 TileLang 高效能運算元倉:統一管理 Kernel 實現、自動化測試用例、標準工作負載與效能基線,目標是形成可複用、可驗證、可持續演進的運算元資產。倉內已有 src、tests、scripts、docs 四層與 OAT.xml、SECURITY.md 等開放原子社群規範檔案,10-22 推送一筆修復”合法上界種子與用例雜湊組合後溢位導致 CPU 測試失敗”;更早的提交裡有”把 Ascend A5 引用替換為 Ascend 950”與”benchmark 失敗或中斷的會話仍覆寫效能基線”。以單一 DSL 為唯一實現語言的社群級運算元倉此前未見,這條線值得按周跟蹤。
來源:https://github.com/Atlas-Ascend-Fork/ops-tilelang
4.2 tilelang/agent:把運算元最佳化寫成 agent 的 SKILL 工作流(10-22)
tilelang/agent(建倉 08-28,3 星)10-22 推送,README 本身是一段可執行的 prompt 模板:以 @tilelang-operator-dev/SKILL.md 注入領域知識,輸入只有運算元名、GPU/CPU 參考實現、shape 與”tilelang npu 交付、中文交付”約束;註釋寫明 agent 需在 A5 環境執行、載入後按 10 步閉環走,Skill 內含選型決策樹、8-CSV 規則、L0 到 L4 策略與 11 條陷阱。把 4.1 的運算元倉、3.15 的指令級排程與本倉合看,昇騰側正在形成”人寫規則、agent 調優、倉裡沉澱基線”的分工。
來源:https://github.com/tilelang/agent
4.3 社群新倉:Sail 包裝、遠端自動調優 agent 與教程倉(10-22)
flytiger-eco/tilelang-wrapper-for-sail 於 10-21 建倉、10-22 推送(Python,無說明);cc-c122/TileLang-Remote-Autotuning-Agent(Python,10-22 推送)與 easy-tilelang/tilelang-devkit、easy-tilelang/tilelang-tutorials 同日推送。數量多、單倉星標低,屬教程與工具層的自然生長,尚無可以判斷的技術方向。
來源:https://github.com/flytiger-eco/tilelang-wrapper-for-sail
4.4 文件與版本:文件站更新三次,VERSION 仍 0.1.15(10-22)
tilelang.github.io 三次 Update docs,TileFoundry.github.io 一筆 render TileFoundry v0.0.3,TileOPs.github.io 一筆。主倉 VERSION 停在 0.1.15,最近釋出仍是 09-30 的 v0.1.15;倉庫 8742 星、915 fork、開放 issue 389 條。
來源:https://github.com/tile-ai/tilelang.github.io
五、趨勢觀察
一,主倉的缺陷型別換了層。 上一週主倉清的是”形狀與幾何校驗”——cast 舍入註解被靜默丟棄、無 arch 的 CUDA target 靜默退到 sm_50、稀疏後設資料形狀不校驗。這一視窗九筆裡有四筆落在分析器與原始碼位置這類編譯期狀態上:analyzer 在改寫後被複用、切分後變數重繫結、eager 前端丟失 Python 位置。這類缺陷的共性是”生成的 kernel 可能仍然正確,但編譯器對你說的和它實際做的不一致”,而 TileLang 正在接入 agent 自動寫 kernel(4.2),編譯期狀態一致性會直接決定 agent 迭代閉環能不能收斂。
二,昇騰側的進展從運算元數量轉到底層資料通路上。 #3530 一筆 18 檔案,做的是 NZ 壓縮佈局在部分 UB→L1 複製中的行距表達;#3581 與 #3523 分別修編譯期謂詞爆炸與 SIMT 向量區域性變數的分配語義。三筆都不新增運算元,都是讓既有路徑在更多形狀上可編譯、可正確執行。配合 CANN 社群另起 ops-tilelang 運算元倉(4.1),可以讀出分工:運算元資產由社群倉承載,主倉專注後端語義與 codegen。
三,TileOPs 兩天連續 20 筆,工作重心在 SM90 特化與 SM89 補位。 視窗內新增三個 SM90 warp 特化 kernel(dense decode、因果滑窗 band、FP8 varlen 重疊),同時把 SM89 的 SSD chunk scan、paged DSA、NSA top-k 三條小規模路徑一併推進。值得注意的是兩筆 NaN 修復(#2708、#2699)都出現在 warp 特化路徑上:共享記憶體規劃器的活躍期判斷和 fast math 下的 FMA 融合,是這類手工編排 kernel 最容易出錯的兩種地方,而這恰好是 TileLang 承諾要替使用者管住的部分。
四,一個反向指標是基準層。 #2692 把參考實現改到 float32 計算,起因是 MTT S5000 上 GroupNorm 與 InstanceNorm 失敗——問題出在參考實現跟著輸入 dtype 走,而不是 kernel 本身。同日 TileFoundry #256 讓”裝置拒絕啟動”從靜默變成拋錯。兩筆都在同一個方向:把”看起來透過了”變成”確實透過了”,對依賴這些工具做晶片適配的團隊,這類修復的實際價值高於幾條效能數字。
五,分發與產業兩條線在加速。 Sunrise 分發版把 0.1.15 與 TANG 後端的 wheel 打包成釋出 PR,說明 TileLang 已經出現不止一個下游分發形態;產業側 SCMP 轉載 SemiAnalysis 的判斷,把”六倉中只有 TileLang 可用”推到主流媒體。技術節奏上這仍是平靜的一天(主倉無釋出、無 breaking change),但生態層的組織動作明顯多於程式碼層。
附錄:素材與核查說明
信源核查表
| 信源 | 核查結果 |
|---|---|
| tile-ai 組織倉庫 | 29 倉,視窗內推送 10 個;合入 35 筆,在審更新 71 筆 |
| tilelang 主倉 | 9 筆合入,1 筆單行修復,1 筆 18 檔案昇騰改動;VERSION 0.1.15,無新發布 |
| TileOPs | 20 筆合入,全部為 Attention / MoE / Elementwise / Reduction 的 kernel 與測試 |
| TileFoundry | 3 筆合入,另 4 筆在審;0.0.3 於 10-21 釋出的說明已在上一期收錄 |
| tilelang-ascend | 2 筆合入(1 新增示例、1 回退示例),11 筆在審 |
| tilelang-hygon | main 無提交;合併分支 gcy_merge_main 一筆開關清理 |
| tilelang-sunrise | 新建 candidate 分支並開 0.1.15+sunrise.1.2.0 釋出 PR,未合入 |
| tilelang-metax / musa / mlir-ascend | 視窗內無推送,分別停在 10-16 / 09-30 / 10-19 |
| TileRT / tilescale / DeepStack | 視窗內無推送 |
| 採用方 TileKernels / FlashQLA | 均停在 10-14,無推送 |
| Google News(中英文多組,走代理) | 命中 3 條,2 條為同一份 SemiAnalysis 報告的媒體轉載,1 條為 DeepSeek 融資與主題無關已剔除 |
| Hacker News | 無 TileLang 相關討論 |
| arXiv | 以 TileLang 為標題或摘要關鍵詞檢索,視窗內 0 條 |
| 社群新倉 | 5 個新倉或活躍倉,均已核對建倉日期與推送時間 |
| 會員或廠商官網 | 本主題不適用 |
完整信源清單
- [1] 主倉向量化前克隆分析器(#3612)— https://github.com/tile-ai/tilelang/pull/3612
- [2] 主倉並行迴圈切分後分析器繫結(#3226)— https://github.com/tile-ai/tilelang/pull/3226
- [3] 主倉 deepseek_v32 top-k 候選緩衝(#3280)— https://github.com/tile-ai/tilelang/pull/3280
- [4] 主倉 CUDA 16 位數學降級(#3172)— https://github.com/tile-ai/tilelang/pull/3172
- [5] 主倉 sm8x 稀疏後設資料形狀校驗(#3125)— https://github.com/tile-ai/tilelang/pull/3125
- [6] 主倉 eager 前端原始碼位置(#3558)— https://github.com/tile-ai/tilelang/pull/3558
- [7] 主倉在審佇列 — https://github.com/tile-ai/tilelang/pulls
- [8] 昇騰 NZ 緩衝部分 UB→L1 複製(#3530)— https://github.com/tile-ai/tilelang/pull/3530
- [9] 昇騰 PTO MAD 方向謂詞增長(#3581)— https://github.com/tile-ai/tilelang/pull/3581
- [10] 昇騰 PTO SIMT 向量區域性變數(#3523)— https://github.com/tile-ai/tilelang/pull/3523
- [11] 昇騰倉 Dilation2D 示例(#1916)— https://github.com/tile-ai/tilelang-ascend/pull/1916
- [12] 昇騰倉回退 depthwise_conv_2d(#1919)— https://github.com/tile-ai/tilelang-ascend/pull/1919
- [13] 昇騰倉 sparse_mla_fwd(#1929)— https://github.com/tile-ai/tilelang-ascend/pull/1929
- [14] 昇騰倉 sparse_mla_bwd(#1928)— https://github.com/tile-ai/tilelang-ascend/pull/1928
- [15] 海光合並分支提交 — https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main
- [16] 沐曦適配倉 — https://github.com/tile-ai/tilelang-metax
- [17] SCMP 轉載 SemiAnalysis 報告 — https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report
- [18] TileOPs FP8 分頁共享記憶體別名(#2708)— https://github.com/tile-ai/TileOPs/pull/2708
- [19] TileOPs dense prefill 持久化(#2706)— https://github.com/tile-ai/TileOPs/pull/2706
- [20] TileOPs SM90 dense decode(#2694)— https://github.com/tile-ai/TileOPs/pull/2694
- [21] TileOPs 因果滑窗 band kernel(#2696)— https://github.com/tile-ai/TileOPs/pull/2696
- [22] TileOPs 分頁 decode 調優 block table(#2703)— https://github.com/tile-ai/TileOPs/pull/2703
- [23] TileOPs masked softmax NaN(#2699)— https://github.com/tile-ai/TileOPs/pull/2699
- [24] TileOPs varlen RoPE 迴圈邊界(#2700)— https://github.com/tile-ai/TileOPs/pull/2700
- [25] TileOPs FP8 grouped GEMM 分級 s
cale(#2704)— https://github.com/tile-ai/TileOPs/pull/2704
- [26] TileOPs gate/up 乒乓 warp group(#2693)— https://github.com/tile-ai/TileOPs/pull/2693
- [27] TileOPs 謂詞廣播向量讀(#2702)— https://github.com/tile-ai/TileOPs/pull/2702
- [28] TileOPs 布林運算元 evict-first(#2697)— https://github.com/tile-ai/TileOPs/pull/2697
- [29] TileOPs SM89 SSD 流水分級(#2705)— https://github.com/tile-ai/TileOPs/pull/2705
- [30] TileOPs SM89 paged DSA 拆分(#2691)— https://github.com/tile-ai/TileOPs/pull/2691
- [31] TileOPs NSA top-k 列 tile(#2690)— https://github.com/tile-ai/TileOPs/pull/2690
- [32] TileOPs 歸約摺疊(#2689)— https://github.com/tile-ai/TileOPs/pull/2689
- [33] TileOPs 空 DSA 選擇校驗(#2698)— https://github.com/tile-ai/TileOPs/pull/2698
- [34] TileOPs 參考實現改 float32(#2692)— https://github.com/tile-ai/TileOPs/pull/2692
- [35] TileFoundry sm_89 架構與 RTX 6000 Ada(#259)— https://github.com/tile-ai/TileFoundry/pull/259
- [36] TileFoundry gmem 計價與維度繫結(#258)— https://github.com/tile-ai/TileFoundry/pull/258
- [37] TileFoundry 啟動失敗拋錯(#256)— https://github.com/tile-ai/TileFoundry/pull/256
- [38] TileFoundry 0.0.3 釋出說明 — https://github.com/tile-ai/TileFoundry/releases/tag/v0.0.3
- [39] Sunrise 0.1.15+sunrise.1.2.0 釋出 PR — https://github.com/tile-ai/tilelang-sunrise/pull/9
- [40] 採用方 TileKernels — https://github.com/deepseek-ai/TileKernels
- [41] CANN 社群 ops-tilelang 運算元倉 — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [42] tilelang/agent 運算元最佳化 skill — https://github.com/tilelang/agent
- [43] Sail 包裝倉 — https://github.com/flytiger-eco/tilelang-wrapper-for-sail
- [44] 遠端自動調優 agent 倉 — https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent
- [45] 主倉文件站 — https://github.com/tile-ai/tilelang.github.io