TileLang 每日動態報告(2026-10-30)
調研視窗:2026-10-29 07:00 ~ 2026-10-30 07:00(北京時間,24 小時)。
信源:GitHub tile-ai 組織逐倉核查(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文件站 / 夜間映象倉)+ 合入與在審真值檢索(合入 12 筆、在審更新 54 筆)+ 新增 issue 全量核對(4 條)+ 主分支逐筆提交核查(tilelang 5、tilelang-ascend 5、TileOPs 1、TileFoundry 1、文件站 5)+ 釋出頁與 tag 核查 + 第三方後端與生態倉核查(10 個倉庫)+ Google News 中英文六組查詢(走代理)+ Hacker News + arXiv。
本期索引
- 今日重點:昇騰倉五筆合入補齊演算法側運算元(GLM-5.3 kpool 壓縮、Chunk Scaled Dot KKT、W4A4 GEMM 修復、depthwise 卷積示例、效能剖析 skill);主倉把全 warp 歸約改走 redux.sync、把 autotuner 的獨立配置改成併發編譯;24 小時合入 12 筆、在審更新 54 筆、新開 issue 4 條
- 一、核心專案進展
- 1.1 主倉:全 warp 歸約改走 redux.sync,四個 32 lane 的 FP32 取大從 20 條 shuffle 變 4 條硬體歸約(#3692)(10-29)
- 1.2 主倉:autotuner 的獨立 JIT 配置改為併發編譯,鎖只保留在同鍵候選(#3476)(10-29)
- 1.3 主倉:昇騰緩衝區檢視的佈局與指標單位不再跨別名傳播(#3623)(10-29)
- 1.4 主倉:GM 到 L1 的複製與越界填充合併成一個排程任務,去掉多餘 MTE2 柵欄(#3543)(10-29)
- 1.5 主倉:CI 把 ROCm 與 Metal 任務鎖回相容的 PyTorch wheel(#3706)(10-29)
- 1.6 新開缺陷三條:帶填充的共享緩衝走批次複製報形狀不符、負縮放迭代圖誤編譯反向索引、檢視位寬在 int32 溢位(#3704 / #3705 / #3707)(10-29)
- 1.7 三條修復在審:位寬改 int64、CPU 標量 GEMM 只清目標區域、關快取時繞過兩層程序內快取(#3708 / #3709 / #3710)(10-29)
- 二、多後端適配(昇騰 / 海光 / 摩爾執行緒 / 沐曦)
- 2.1 昇騰:GLM-5.3 稀疏注意力 indexer 的 kpool 壓縮運算元合入,單 kernel 純 Vector 路線(#1910)(10-29)
- 2.2 昇騰:分塊門控線性注意力的 Chunk Scaled Dot KKT 前向運算元合入,靜態四塊展開、零 GM 工作區(#1803)(10-29)
- 2.3 昇騰:W4A4 的 int4 乘 int4 GEMM 結果全錯問題修復合入,根因是型別尺寸謊報與專用轉置指令被靜默關閉(#1731)(10-29)
- 2.4 昇騰:depthwise_conv_2d 示例重提,改用清單化測試通道以避開併發批次(#1920)(10-29)
- 2.5 昇騰:新增運算元效能剖析 skill,用 msprof op 輸出定最佳化方向(#1931)(10-29)
- 2.6 昇騰在審:NPUIR 語言層三筆檢視與廣播修正(#1962 / #1960 / #1961)(10-29)
- 2.7 昇騰在審:PTO 後端按位異或的無符號位元組工作區檢視修正(#1959)(10-29)
- 2.8 海光:HCU 兩筆在審有更新、社群出現 DCU 融合運算元倉;摩爾執行緒與沐曦視窗內無推送(10-29)
- 三、生態與採用方(TileOPs / TileFoundry / TileKernels / FlashQLA / 第三方)
- 3.1 TileOPs:釋出映象合入,運算元庫與 TileFoundry 首次打進同一映象(#2793)(10-29)
- 3.2 TileOPs 在審:GEMM 按功耗封頂重算柵格分組與常駐塊數,q-b 的 DRAM 讀從 168 MB 降到 94 MB(#2794)(10-29)
- 3.3 TileFoundry:AMX 執行時運算元改成 39 個靜態 TIR 契約,屬破壞性變更(#279)(10-29)
- 3.4 採用方:TileKernels 與 FlashQLA 主分支視窗內無推送(10-29)
- 3.5 昇騰側出現獨立的社群運算元倉 ops-tilelang,視窗內有推送(10-29)
- 3.6 PTO-ISA 配套雙倉活躍:題庫倉遷移到 PTOAS 0.70,基準倉加 SV2 分塊縮放與逐列分析(10-29)
- 3.7 觀察:T-HEAD 名下的 tilelang 分支倉視窗內有推送,分支名指向 PPU 的矩陣指令佈局(10-29)
- 四、社群、教程與活動
- 4.1 文件站把內建歸約的適用邊界寫進指令頁,與 redux 合入同步(10-29)
- 4.2 開發工具倉完成下沉跟蹤抽取後的三筆收尾,與上游在審工具化改動同向(10-29)
- 4.3 新倉 tilelang-knowledge 一天四筆:33 個運算元設計案例、20 張帶出處的知識卡、7 條最佳化模式(10-29)
- 4.4 圖解課配套倉與分散式佈局實驗倉同日新建(10-29)
- 4.5 新聞側:Google News 六組中英文查詢、Hacker News 與 arXiv 均無視窗內實質命中(10-29)
- 五、趨勢觀察
- 附、素材與核查說明
今日重點
今天的合入量從昨天的 24 筆回落到 12 筆,但結構更集中:昇騰專用倉佔五筆、主倉五筆、TileFoundry 與 TileOPs 各一筆。昇騰倉這一批全是演算法側運算元,而且都是能對上游模型直接用的那一類。GLM-5.3-Flash 稀疏注意力 indexer 的 kpool 壓縮運算元把逐維 softmax 池化、BF16 邊界、歸一化 Hadamard-128 旋轉與 E4M3 量化壓進一個純 Vector 的 kernel,每塊覆蓋 64 個 pool、兩個向量子核各管 32 個,宿主側零重排,直接按呼叫方的三維佈局做跨步搬運;E4M3 編碼用純整數位運算實現(round-half-even、含次正規、飽和到 448),不依賴 FP8 型別也不做近似。Chunk Scaled Dot KKT 是分塊門控線性注意力前向裡的 A = gate(K K^T),靜態四塊展開、零 GM 工作區、零宿主運算元,L0 與 L1 三例的匹配率都是 1.0。
主倉今天的主題是編譯與發射層。全 warp 歸約改用 redux.sync:主分支上四個獨立的三十二線 FP32 取大在 CUDA 13.1、sm_103a 上展開成 20 條 shuffle,改後變成四條 redux.sync.max.f32 且不含 shuffle;整數歸約在 SM80 以上走硬體歸約,FP32 的取大取小隻在 sm_100a、sm_100f、sm_103a、sm_103f 上可用,窄型別(含 CUTLASS 的 FP16、BF16)提升到 FP32,FP64 仍留 shuffle 路徑。作者明確不聲稱端到端加速,只給了指令級對照。autotuner 的併發編譯則是把整段持有的編譯鎖拆開:每份覆蓋配置用 dataclasses 複製出獨立的 JITImpl,只有共享第一階段鍵的候選才序列,除錯輸出的寫盤仍序列。作者給出的跨後端驗收是昇騰 1797 項、ROCm 3220 項、Metal 95 項、CUDA 5500 項。
另有一處值得記的細節:今天新開的反向索引誤編譯缺陷(#3705)是拿 T-HEAD 名下的 tilelang 分支倉做執行時復現的,這說明已有整建制團隊把這條上游程式碼當作自有硬體棧的底座在跟,而不只是當一個外部依賴。
一、核心專案進展
1.1 主倉:全 warp 歸約改走 redux.sync,四個 32 lane 的 FP32 取大從 20 條 shuffle 變 4 條硬體歸約(#3692)
合於 10-29 15:08,4 檔案 +337/-35。此前普通的跨執行緒歸約(如 T.reduce_max)即使整組就是一個對齊的滿 warp,也照樣展開成 shuffle 樹:主分支 837486a 上四個獨立的 32 lane FP32 取大在 sm_103a 配 CUDA 13.1 編譯出 20 條 shuffle、零條 redux,改後是四條 redux.sync.max.f32、零條 shuffle。做法是複用標量與批次全歸約裡已有的 warp_reduce 助手:呼叫方只判執行緒幾何(是否與 warp 對齊、單位步長、滿 warp),架構、資料型別與運算元分發交給助手。整數歸約在 SM80 及以上走硬體 redux,更小的整數提升到 32 位;FP32 的取大取小隻在 sm_100a、sm_100f 與 sm_103a、sm_103f 上可用;窄浮點提升到 FP32 後再進助手,FP64 不降精度、仍在 shuffle 路徑。子 warp、跨步、打包歸約以及不支援的運算元與目標一律回退 shuffle,跨 warp 的共享記憶體交換與柵欄不變。順帶修掉一處編譯失敗:自定義無符號整數 reducer 此前因內聯分發返回空型別而編譯不過,現在落到 shuffle 回退。PTX 迴歸斷言真實指令與 shuffle 缺席,執行時覆蓋有符號無符號邊界值、批次不清零、NVRTC 執行、非有限值、有符號零與 FP64 精確結果。作者明確不聲稱端到端加速。既有的一處遺留問題也在描述裡點明:批次歸約在 batch=4 每執行緒兩輸出、batch=8 每執行緒四輸出時存在 FP16 與 BF16 失敗,主分支原模板同樣復現,且失敗發生在下沉階段,與本筆無關。
來源:https://github.com/tile-ai/tilelang/pull/3692
1.2 主倉:autotuner 的獨立 JIT 配置改為併發編譯,鎖只保留在同鍵候選(#3476)
合於 10-29 16:24,2 檔案 +54/-46。主分支的 _make_jit_compile_func 在整段編譯期間持一把鎖,連只用預設編譯選項的配置也序列,前端、下沉與裝置編譯因此全部被排隊。改法是給每份覆蓋配置用 dataclasses.replace 構造候選 JITImpl,後端、目標、輸出選擇、日誌級別、編譯標誌與除錯選項保持原值;預設選項仍走原有惰性呼叫路徑,保留其核心快取與物件身份;只對共享第一階段鍵(由 get_tir 用的引數繫結器算出,保留位置引數、預設值與帶張量註解的工廠)的候選序列;除錯檔案的共享輸出仍在編譯內部序列,覆蓋 debug_root_path、IR 轉儲設定與佈局視覺化。驗收資料在跨後端流水線上:昇騰任務合併 3be11be6 與主分支 c58a86f0,Python 3.12.14、torch_npu 2.10.0、CANN 9.3.0,1797 項檢查透過,AscendC 示例 50 例全過、PTO 示例 44 例過 6 例跳過;ROCm 7.2 任務在 gfx942 上 3220 項透過、2765 項跳過;Metal 任務在 macOS arm64 上 95 項透過、5 項跳過;CUDA 任務在 H100 PCIe、cu130 上 5500 項檢查透過、491 項跳過,示例 76 過 14 跳;CuTeDSL 任務 75 過 17 跳 2 例意外透過。
來源:https://github.com/tile-ai/tilelang/pull/3476
1.3 主倉:昇騰緩衝區檢視的佈局與指標單位不再跨別名傳播(#3623)
合於 10-29 13:17,23 檔案 +1139/-438。由 T.view 或 T.reshape 造出的多個緩衝可以共用同一段儲存而要求不同佈局,此前統一按共享的資料變數解析佈局,於是可能選中錯誤的 Cube 對映或 DMA 地址,陳舊的分塊幾何還會丟掉版本偏移、或把帶型別的指標偏移按錯的單位解釋。本筆把佈局附著到單個 Buffer 檢視上,貫穿佈局推斷、分形儲存歸一化、多緩衝物化與複製和矩陣乘的下沉;檢視身份保留到佈局對映與緊湊 L0 間距應用之後,裸指標偏移按各自的元素單位重對映。跨儲存別名的佈局不再傳播,每個檢視用自己的註解、運算元推斷出的佈局或普通定址;同一個帶版本的分配上,Cube 檢視現在可以用不同的 MajorK 與 MajorMN 佈局,條件是邏輯步長稠密、基址偏移為零、運算元型別相容且填充後的儲存大小相等。多緩衝別名共用一個物理槽位大小,按各檢視的元素單位換算而不傳播索引對映;檢視只是重新解釋已有位元組,寫入方必須按消費者期望的表示寫入,需要選檢視給內建指令時可走訪問指標介面。
來源:https://github.com/tile-ai/tilelang/pull/3623
1.4 主倉:GM 到 L1 的複製與越界填充合併成一個排程任務,去掉多餘 MTE2 柵欄(#3543)
合於 10-29 11:39。全域性記憶體到 L1 的複製與其越界填充共同準備同一塊 L1 資料,此前填充可能被拆成獨立的排程任務,同步插入就會在複製與填充之間加一道多餘的 MTE2 管道柵欄。本筆把每個複製與其填充編進同一個昇騰任務,消費者等待的是包含填充在內的完整資料;FP8 MX 的 K 尾填充同樣併入並保留任務內的條件填充,已有的顯式任務作用域保持不變、不引入巢狀;原來的 MX 延遲填充佇列被刪除,每個運算元的填充緊跟自己的複製。47 項聚焦測試透過,覆蓋普通越界填充與 MX 尾塊兩類。
來源:https://github.com/tile-ai/tilelang/pull/3543
1.5 主倉:CI 把 ROCm 與 Metal 任務鎖回相容的 PyTorch wheel(#3706)
合於 10-29 21:38。ROCm 的 CI 透過索引變數暴露匹配的 PyTorch ROCm 源,但工作流級的「不安全最優匹配」策略仍讓更晚的 PyPI 候選勝出:#3534 的 ROCm 任務裡解析器選了 PyPI 的 PyTorch 2.15.0(CUDA 13.2),而不是可用的 2.14.1 加 ROCm 7.2 構建,官方在執行器上直接復現了這兩種解析結果。同一批 2.15.0 的鋪開也打斷了 Metal 任務:macOS 26.6.2 的 ARM64 執行器上,2.15.0 因為不暴露 Metal 2.0 而拒絕 Apple 半虛擬 GPU,同一映象上上次透過用的是 2.14.1。改法有三條:按配置的穩定 ROCm 版本設定後端變數,讓 uv 把 PyTorch 留在 ROCm 源上而其餘依賴仍用最優匹配;夜間 PyTorch 安裝改用首個索引,使夜間任務無法逃出配置的源;新增一個快速失敗的校驗步,在測試套件啟動前確認 PyTorch 是 HIP 構建、不是 CUDA 構建、並能訪問執行器 GPU。
來源:https://github.com/tile-ai/tilelang/pull/3706
1.6 新開缺陷三條:帶填充的共享緩衝走批次複製報形狀不符、負縮放迭代圖誤編譯反向索引、檢視位寬在 int32 溢位(#3704 / #3705 / #3707)
三條 issue 都指向靜態檢查的邊界。其一,#3704 報告把末維不是 16 對齊的全域性張量(例如 K=100)拷進帶填充的共享緩衝(BK 取 cdiv(K,16) 乘 16 得 112)時,一旦這支複製被路由到批次複製路徑,下沉直接報共享塊與全域性塊形狀不符(112 對 100),位置在 CUDA 複製的批次下沉處;逐元素路徑下的帶填充緩衝是正常的,而帶填充恰好是滿足矩陣乘對齊的常見寫法,任何與批次複製同用的核心都可能踩到。其二,#3705 報告打平緩衝的簡化在負縮放迭代圖上誤編譯反向索引,復現環境是 v0.1.15 與 T-HEAD 名下 tilelang 分支倉的原始碼構建,報告同時核對了上游主分支與對應 TVM 修訂裡受影響的拆分函式未變,屬上游共有路徑。其三,#3707 報告檢視與重塑在核對緩衝位寬時按 int32 相乘,超過 2 的 31 次方即迴繞:Qwen3.8 的 lm_head 權重為 248320 乘 2560 位元組的 NVFP4,共 5,085,593,600 位,位寬乘積迴繞成 790,626,304,隨後把靜態檢視位數轉 int32 常量時丟擲字面量超限。
來源:https://github.com/tile-ai/tilelang/issues/3704
來源:https://github.com/tile-ai/tilelang/issues/3705
來源:https://github.com/tile-ai/tilelang/issues/3707
1.7 三條修復在審:位寬改 int64、CPU 標量 GEMM 只清目標區域、關快取時繞過兩層程序內快取(#3708 / #3709 / #3710)
對應上面三條缺陷與昨天的快取語義議題,修復同日進審。位寬那筆把靜態維度乘積改為按 Python 整數相乘、符號化維度轉 int64,並把超出 int32 的字面量轉成 int64 常量,迴歸用例直接取該 issue 的緩衝(248320 乘 5120 的檢視形狀),主分支報錯、改後透過,符號化維度的同形檢視也覆蓋。CPU 那筆修標量矩陣乘在清零累加時清整塊底層緩衝的問題:對 4 乘 4 輸出緩衝把 7 初始化後呼叫 1 到 3 行列的子區域,原實現會把區域外 12 個元素一併歸零,改法只用輸出區域自身的偏移與範圍初始化,迴歸用例照此寫了全 1 的 2 乘 2。快取那筆修昨天記錄的 #3699:關閉核心快取只翻轉了快取狀態,而 jit 包裝器在真正編譯前還查兩個程序內快取,塊內普通呼叫會拿到同程序早先編好的核心;改法是呼叫時讀快取開關(同時覆蓋環境變數禁用),關閉期間跳過兩層程序內快取,出塊後恢復。測試從改前的 1 失敗 4 透過變成 5 全過。
來源:https://github.com/tile-ai/tilelang/pull/3708
來源:https://github.com/tile-ai/tilelang/pull/3709
來源:https://github.com/tile-ai/tilelang/pull/3710
二、多後端適配(昇騰 / 海光 / 摩爾執行緒 / 沐曦)
2.1 昇騰:GLM-5.3 稀疏注意力 indexer 的 kpool 壓縮運算元合入,單 kernel 純 Vector 路線(#1910)
合於 10-29 15:41。運算元對每個 pool(pool_size 為 4)把打分張量與位置編碼相加後做逐維 softmax 池化,池化向量過 BF16 邊界,再乘歸一化 Sylvester Hadamard-128 旋轉(除以 128 的平方根),最後按每個向量的 FP32 二次冪絕對最大值縮放做 E4M3 量化:E4M3 行與縮放因子寫進連續暫存區,由宿主介面卡散射進分頁快取(鍵快取用無符號位元組承載 E4M3,縮放快取原地寫 FP32)。實現是單 kernel 純 Vector 路線,開發者模式開啟自動同步、自動向量與矩陣核合併、記憶體規劃三個開關,每塊覆蓋 64 個 pool、兩個向量子核各處理 32 個獨立並行;宿主零重排,kernel 直接按呼叫方的三維佈局跨步讀取,位置編碼折入池化;E4M3 編碼是純整數位運算,round-half-even、含次正規、飽和到 448,不使用 FP8 型別也不做近似。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1910
2.2 昇騰:分塊門控線性注意力的 Chunk Scaled Dot KKT 前向運算元合入,靜態四塊展開、零 GM 工作區(#1803)
合於 10-29 15:01。新增分塊內積的縮放點積前向:對分塊門控線性注意力的前向計算 A 等於門口控的 K 乘 K 轉置。實現為單個編譯器裝飾的 kernel,靜態展開四個分塊做片上融合,零全域性記憶體工作區、零宿主運算元呼叫,開發者模式四個開關全開。測試分層為 L0、L1、L2 與邊界四檔並可命令列選檔,L0 覆蓋三種情形(帶門控、不帶門控、小規模帶門控),精度對照 PyTorch CPU 的 FP32 金標準,採用混合容差雙門。報告給出的測形為批 1、序列長度 32768、頭 32、每頭維度 128、分塊 64、開啟門控,輸入輸出 BF16、累加 FP32,L0 與 L1 各三例全部透過,匹配率 1.0000,最大絕對誤差在 0 到 0.25 之間。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1803
2.3 昇騰:W4A4 的 int4 乘 int4 GEMM 結果全錯問題修復合入,根因是型別尺寸謊報與專用轉置指令被靜默關閉(#1731)
昨天仍在審的這筆今天 15:08 合入,對應 #438。根因是三層錯誤疊加。其一是分形幾何按型別尺寸等於 1 推導,而半位元組打包型別每個元素只佔 4 位、每位元組裝兩個元素,於是 32 位元組的 C0 塊正確元素數是 64 而不是 32,所有複製步長、重複次數與位元組計算全部偏差一倍,資料在 L0A 與 L0B 裡堆到錯的位置。其二是 B 運算元的轉置載入被靜默丟棄:通用載入路徑對 zN 到 nZ 的轉換設了轉置標誌,但四位元專用指令硬編碼關閉轉置並忽略該標誌,B 永遠得不到矩陣乘指令期望的佈局;正確路徑是帶轉置的專用載入,而它在 CANN 標頭檔案裡沒有任何文件,只能用探針 kernel 在裝置上實測反推引數(一次重複搬完完整的塊,重複次數取 N 除以 64、源步長取 K 向上取整 16 後除以 64、目的間隙取 3),且該版本實際按 2048 位元組塊工作,與通用文件的 512 位元組分形單位不同。其三是向量與全域性記憶體之間的位元組計數同樣按每元素 1 位元組而非 0.5 計算,每行複製長度偏長一倍,混入相鄰行資料。計算指令本身始終正確,只是忠實地把擺錯的資料相乘,症狀是 99.6% 不匹配、輸出為負 501 與 627 而正確值是 641 與 1785。報告同時寫清了全域性記憶體的 int4 打包約定:每位元組兩個半位元組,低半位元組為偶數元素、高半位元組為奇數元素,第 r 行起始於 r 乘 N 除 2 的位元組位置,與 CANN 自家格式一致。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1731
2.4 昇騰:depthwise_conv_2d 示例重提,改用清單化測試通道以避開併發批次(#1920)
合於 10-29 15:43。這是 #1915 的重提:原筆合併後夜間全量 CI 失敗,經 #1919 回退,根因是 L0 最大形狀用例在八路併發的遺留指令碼批次下觸發未知錯誤;本筆把該用例瘦身四分之一,並把運算元註冊進持續整合的運算元測試清單走 Pytest 通道,與另外二十五個已遷移運算元同路線,徹底撤出併發批次,kernel 本體與原筆逐位元組一致。運算元本身支援 fp16、fp32、bf16 與任意核大小、步長、填充、膨脹,步長 1、2、4 走向量化區域路徑,其餘走標量內積兜底;實現是填充加摺積雙 kernel 加 1 乘 1 快路徑,純 Vector 路線,網格按幾何自適應不寫死核數。報告裡幾處實測約束值得記:填充階段耗時約等於語句條數乘 35 納秒、與位元組數無關,故末帶用截斷收尾;共享記憶體側的複製目標偏移若不是 32 位元組對齊會觸發向量核的 0x26 異常;同步用手寫冠軍模型(卷積段關自動同步),通道間用深度 1 的相鄰搬運與計算事件協議,深度 2 的協議或中途插入全柵欄都會死鎖。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1920
2.5 昇騰:新增運算元效能剖析 skill,用 msprof op 輸出定最佳化方向(#1931)
合於 10-29 17:17。新增一個運算元效能剖析的開發流程資產:在效能最佳化前用 msprof 的運算元級跑一遍,再按輸出判斷下一步最佳化方向。作者已在加法逐元素、歸約取小示例與動態量化三個運算元上驗證過流程可用。這類資產的價值在於把「先量後改」寫進流程,而不是每次重新摸索剖析口徑。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1931
2.6 昇騰在審:NPUIR 語言層三筆檢視與廣播修正(#1962 / #1960 / #1961)
三筆同日進審且互相依賴。其一,#1962 修本地統一緩衝的檢視描述符:橋接層按共享資料變數查緩衝,可能返回未應用描述符的分配記憶體引用,同型別偏移別名於是讀到分配區字首,不同型別別名則帶著分配的元素型別走進複製下沉,同型別重塑還可能帶著錯誤維度進入歸約;改法是引入統一的檢視獲取助手,同型別非恆等檢視用帶檢查的重解釋轉換,多元素型別訪問的分配用單塊位元組分配加帶型別的檢視操作保持共同儲存身份。其二,#1960 修檢視丟掉源偏移:檢視從資料變數重建緩衝時丟掉了元素偏移,偏移別名因此從分配偏移零開始而不是源首元素地址;改法對同型別保留偏移,位元組級型別變化時按元素寬度換算並要求可整除,不可對齊或不支援的轉換直接拒絕,不引入儲存或資料複製。其三,#1961 改一元向量廣播按訪問區域推斷:此前按原緩衝形狀推斷廣播維度,把 8 乘 8 源寫到 24 乘 8 目標的行 8 到 16 會失敗;改法改按生成的記憶體引用型別推斷,保留訪問區段的形狀與既有相容性檢查。三筆都附了獨立的乾淨基線構建與逐位驗證說明。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1962
2.7 昇騰在審:PTO 後端按位異或的無符號位元組工作區檢視修正(#1959)
進審於 10-29 16:36,關閉 #1721。PTO 後端裡按位異或的臨時工作區被當作原始一維位元組區傳給指令,而不是所需二維檢視:生成時把目的端的二維形狀複製到工作區的形狀資訊,卻沒標記為切片,於是對無符號位元組(工作區不做型別重解釋)解析出一維塊描述,指令執行時越界讀寫。工作區本身是一維區按源元素數分配,容量對對齊的二維塊是夠的,失敗在描述而非容量。修法是繼承目的形狀後強制構造二維檢視並標記為切片,測試補了無符號與有符號位元組的多組覆蓋(二維自動臨時、顯式一維臨時、一維、對齊非零偏移切片加哨兵、有符號位元組),並保留 16 位迴歸。
來源:https://github.com/tile-ai/tilelang-ascend/pull/1959
2.8 海光:HCU 兩筆在審有更新、社群出現 DCU 融合運算元倉;摩爾執行緒與沐曦視窗內無推送(10-29)
海光側仍在審:普通複製的共享記憶體縱向交錯特性分支(#30)與分散式那筆(#19)視窗內均有更新,主分支最新提交仍停在 10-27 的緩衝區分組特性。社群側同日出現一個 DCU 方向的倉庫,把海光平臺的融合運算元與對 TileLang 編譯器的改動一併保留,C++ 為主、約 15 MB,歷史提交裡可見注意力運算元支援與可變長支援,屬個人實踐倉而非海光官方。摩爾執行緒適配倉(tilelang-musa)與沐曦適配倉(tilelang-metax)視窗內沒有推送,最新分別是 09-30 與 10-26。
來源:https://github.com/tile-ai/tilelang-hygon/pull/30
來源:https://github.com/SealSoul117/TileLang_DCU_FuseCmpComm
三、生態與採用方(TileOPs / TileFoundry / TileKernels / FlashQLA / 第三方)
3.1 TileOPs:釋出映象合入,運算元庫與 TileFoundry 首次打進同一映象(#2793)
合於 10-29 11:18,9 檔案 +123/-49。執行器映象新增釋出目標:在開發映象之上安裝運算元庫與排程器的 PyPI 釋出版,並用兩者版本號命名;映象標籤為 tileops0.1.0 加 tilefoundry0.0.3 再加 cuda13.2、torch2.13 與 tilelang0.1.12 三段環境標識,釋出在容器映象倉庫的運算元庫釋出路徑下。附加指令碼會拉取映象並在 GPU 上跑棧檢查、基線檢查、打包測試與可選的基準檔案,因為輪子裡不帶測試與基準,故都從映象自身版本對應的標籤取。文件更新了快速開始,執行器文件補了構建、驗證與推送步驟。同時移除了 sgl-kernel 基線:不再有映象安裝它,相關基準此前一直在跳過。
來源:https://github.com/tile-ai/TileOPs/pull/2793
3.2 TileOPs 在審:GEMM 按功耗封頂重算柵格分組與常駐塊數,q-b 的 DRAM 讀從 168 MB 降到 94 MB(#2794)
進審於 10-29 21:44,2 檔案 +35/-12,承接背靠背呼叫的效能議題。兩處改動:其一,當 A 能落進一個二級快取分割槽而 B 不能時,柵格分組讓一個分組跨全部行,於是每個 B 塊只從記憶體讀一次而不是每個十六行分組讀一次,在 DeepSeek V3 預填充的 q-b 形狀(4096 乘 24576 乘 1536,NT 佈局)上記憶體讀從 168 MB 降到 94 MB;窄塊的分裂邊界仍按十六行分組約束,選擇不變。其二,常駐塊數在寬度均勻的雙協作組網格上改為按塊數、總塊數與處理器數算出向上取整的常駐數,而不是直接用處理器數,因為最忙的那塊無論如何都要跑同樣多的塊,而每一波能讀到更少的不同面板;乒乓、窄塊與混合寬度網格仍用處理器數。測試在 H200、CUDA 13.2、PyTorch 2.13、TileLang 0.1.12 上跑運算元測試,212 項透過、24 項跳過。
來源:https://github.com/tile-ai/TileOPs/pull/2794
3.3 TileFoundry:AMX 執行時運算元改成 39 個靜態 TIR 契約,屬破壞性變更(#279)
合於 10-29 15:59,破壞性變更(提交標題帶感嘆號)。動機是讓 AMX 有一份顯式、帶型別的運算元宣告,智慧體可以從同一份計算描述裡匹配與評估。做法是把 24 個 AMX 執行時運算元族描述成 39 個靜態註冊的契約版本,帶引數域、運算元與佈局模式、副作用集合,以及 M1 到 M4 的目標能力;新增 Python 評估器、原始 C 模擬器作為金標準,並補 59 個代表性原子用例覆蓋全部 39 個靜態運算元,高階與中階中間表示用例跑真實的 2、4、8、12 執行緒域並驗證候選的接受與拒絕;高階候選從運算元內的匹配宣告發現,保留針對目標的能力推斷。命名統一到 AMX 名稱空間,別名由註冊入口持有,註冊名由統一查詢給出;CPU 執行用執行緒拓撲,私有邏輯 AMX 暫存器狀態,執行緒數不是 AMX 引擎數。驗證為 46 項聚焦測試透過,含一處 57 個原始 C 結果的比對與兩項生命週期檢查。
來源:https://github.com/tile-ai/TileFoundry/pull/279
3.4 採用方:TileKernels 與 FlashQLA 主分支視窗內無推送(10-29)
兩個已知採用方的主分支都停在 10-14,視窗內無新推送:深度求索的運算元倉(2018 星)與通義側的 FlashQLA(742 星)。本日生態側的新增訊號主要來自社群自建倉而非這兩個內部使用方。
來源:https://github.com/deepseek-ai/TileKernels
來源:https://github.com/QwenLM/FlashQLA
3.5 昇騰側出現獨立的社群運算元倉 ops-tilelang,視窗內有推送(10-29)
該倉 10-20 建立、10-29 13:06 有推送,定位是面向昇騰 NPU 的社群高效能運算元倉,用統一方式管理 kernel 實現、自動化測試用例、標準工作負載與效能基線,目標是讓運算元資產可複用、可驗證、可持續演進。它掛在昇騰生態的分支組織名下而不是 tile-ai 組織裡,說明昇騰側的運算元供給正在從官方倉外溢位一層社群協作面;與官方 tilelang-ascend 的關係尚不明確,收錄時只作生態觀察,不視作官方動作。
來源:https://github.com/Atlas-Ascend-Fork/ops-tilelang
3.6 PTO-ISA 配套雙倉活躍:題庫倉遷移到 PTOAS 0.70,基準倉加 SV2 分塊縮放與逐列分析(10-29)
PTO-ISA 組織下的兩個配套倉視窗內都有提交。題庫倉把 PTO 的廣播載入遷移到 PTOAS 0.70,此前一週在做量化縮放符號的澄清與題面「指向自己的解而不是答案」的改寫;基準倉加了兩筆分析提交:SV2 的分塊縮放與逐列分析,以及一條依據(庫內絕對最大值指令的變體計數為 31.39,附執行時指令吞吐與屋頂線分析,明確不是統一緩衝口徑)。這兩個倉與官方倉的關係是教學與測量配套,值得作為昇騰側方法論的旁證。
來源:https://github.com/PTO-ISA/tilelang-puzzles-ascend
來源:https://github.com/PTO-ISA/tilelangGym
3.7 觀察:T-HEAD 名下的 tilelang 分支倉視窗內有推送,分支名指向 PPU 的矩陣指令佈局(10-29)
T-HEAD 組織(官網為開發者門戶域名)名下有一個 tilelang 的分支倉,2026-09-22 建立,是上游主倉的分支,公開 13 個分支,除主分支與 v0.1.11 到 v0.1.14 的版本映象外,還可見 PPU 矩陣共享佈局修復、PPU language 包、執行時編譯特性與若干缺陷修復分支;該倉主分支最近提交停在 10-22,但倉庫推送時間更新到 10-29 19:27(北京時間)。倉內沒有說明文件,也沒有釋出記錄,因此只能作為觀察線索:有整建制團隊在把 TileLang 作為自有硬體棧底座做分支開發,方向與 PPU 的矩陣指令佈局相關。本日 #3705 缺陷報告正是用該倉的原始碼構建復現的。
來源:https://github.com/t-head/tilelang-for-sail
四、社群、教程與活動
4.1 文件站把內建歸約的適用邊界寫進指令頁,與 redux 合入同步(10-29)
文件站視窗內有五筆自動更新,其中指令頁新增了一段歸約說明,與今天合入的硬體歸約改造對齊:CUDA 的跨執行緒歸約把與 warp 對齊、單位執行緒步長的滿 warp 組委託給與 warp 歸約同源的助手,助手在 SM80 以上的整數求和、取大、取小與按位運算上選硬體歸約指令,更小的整數提升到 32 位;標量 FP32 的取大取小隻在 sm_100a、sm_100f 與 sm_103a、sm_103f 上使用,標量 FP16 與 BF16 提升到 FP32;同一規則也適用於更大歸約的最後一段與未打包的批次歸約。子 warp、跨步、打包、FP64 以及不支援的型別運算元組合保留 shuffle 實現,跨 warp 的共享記憶體交換與柵欄不變。把回退邊界寫進文件,比只寫「支援硬體歸約」對使用者更有用。
來源:https://github.com/tile-ai/tilelang.github.io
4.2 開發工具倉完成下沉跟蹤抽取後的三筆收尾,與上游在審工具化改動同向(10-29)
獨立開發工具倉在 10-20 完成從主倉抽出的下沉跟蹤後,視窗內補了三筆收尾:避免跟蹤啟用時的迴圈匯入、把持續整合固定到支援插樁外掛的版本、在載入核心相容別名前先初始化 TileLang。這與上游在審的「把下沉跟蹤抽到開發工具倉」那筆同向,說明工具化已經不只是上游的分支實驗,而是形成了可獨立釋出的包。
來源:https://github.com/easy-tilelang/tilelang-devkit
4.3 新倉 tilelang-knowledge 一天四筆:33 個運算元設計案例、20 張帶出處的知識卡、7 條最佳化模式(10-29)
該倉 10-29 建立,當天四筆提交構成一套面向智慧體的知識底座:先是倉庫結構,隨後整理 33 個運算元設計案例(每條帶固定出處),再補 20 張帶出處的程式設計知識卡,最後文件化 7 條有證據支撐的 GPU 最佳化模式。倉庫自述的用途是讓智慧體面對一個運算元時能判斷哪些經驗適用、應當如何改程式碼、如何驗證、失敗後往哪裡找。這與 TileFoundry 的運算元契約、昇騰倉的剖析 skill 是同一類嘗試:把「怎麼改、怎麼驗」寫成可檢索的資產。
來源:https://github.com/qiuuuuu622/tilelang-knowledge
4.4 圖解課配套倉與分散式佈局實驗倉同日新建(10-29)
兩個個人倉同日建立:一個圖解課配套倉,自述用原創場景加動畫把每個知識點講透、隨影片釋出;一個分散式佈局實驗倉,提交說明是加入分散式的 TileLang 佈局實驗。加上今天有推送的模型移植倉與 Hexagon 後端倉(倉庫推送時間更新到 10-30 06:10,其主分支最近提交仍是 08-24 的依賴升級),社群在教育與方言後端兩個方向的自發投入都在繼續。
來源:https://github.com/zachshi-ai/tilelang-course
來源:https://github.com/Leewin239/tilelang-distributed-layout
4.5 新聞側:Google News 六組中英文查詢、Hacker News 與 arXiv 均無視窗內實質命中(10-29)
六組查詢覆蓋精確詞加兩日、中文加昇騰、中文加運算元、組織名加兩日、英文加 GPU 與中文加編譯器加七日。僅中文加昇騰返回一條,是臨床人工智慧進展的轉載稿,與運算元語言無關,其餘全部空返回。Hacker News 近四日的相關命中全部落在視窗之外(最近的是一條 10-02 的昇騰程式設計棧轉載討論)。arXiv 以 TileLang 檢索無新論文。新聞側連續多日零命中,日報主體繼續以倉庫與生態倉為準。
五、趨勢觀察
昇騰側的節奏已經從「搬後端」轉到「補運算元」。昨天記的是 NPUIR 後端併入主幹、驗收口徑是 504 項自測與 71 比 72 的示例;今天合入的五筆裡四筆是運算元與示例,而且是按上游模型的需求排的:GLM-5.3 的 indexer 壓縮、分塊門控注意力的內積、四位元量化的矩陣乘、量化視覺模型的深度卷積。這些運算元的共同點是都把量化與融合寫進了 kernel 本體(E4M3 純整數編碼、四位元打包約定、雙 kernel 加 1 乘 1 快路徑),而不是留給宿主側補救。同時兩筆順帶把工程口徑立起來:運算元效能剖析進流程、併發批次造成的不穩定用例改走清單化測試通道。這說明昇騰側的主要矛盾正從「能不能編出來」變成「運算元的數值與效能能不能被穩定驗收」。
主倉今天的三筆(硬體歸約、併發編譯、越界複製分組)分別對應指令選擇、編譯時延與發射質量,都是編譯器核心問題,而不是語言層特性。這與昨天的新增缺陷方向一致:越界、別名、位寬這類靜態檢查的邊界正在被逐個開啟。今天三條新缺陷裡有兩條(帶填充緩衝的批次複製、檢視位寬溢位)都在真實模型形狀上被觸發,說明 TileLang 的使用面已經深入到會碰到邊界條件的規模:248320 乘 5120 的 NVFP4 權重、K 不是 16 整數倍的注意力頭維度。
生態側今天最值得注意的是「fork 化」。昇騰社群出現獨立運算元倉,T-HEAD 名下有帶 PPU 分支的 TileLang 分支倉,PTO-ISA 把題庫與基準都做成了配套倉,另有開發工具倉把下沉跟蹤獨立發包。這些動作的共同前提是把 TileLang 當作自有硬體棧的底座來對待:需要按自己的指令集改佈局、按自己的測試體系做驗收、按自己的知識形態沉澱經驗。對一個領域專用語言來說,這是比星標數更硬的採用訊號,但也意味著上游要開始面對「分支與上游如何同步」的長期成本:今天 #3705 的缺陷就是分支使用者在真實構建裡先撞上的。
附、素材與核查說明
信源核查表
| 信源 | 核查結果 |
|---|---|
| GitHub tile-ai 組織逐倉(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / 文件站 / 夜間映象倉) | 視窗內有推送:tilelang、TileOPs、TileFoundry、tilelang-ascend、tilelang-hygon、tilelang.github.io、TileOPs.github.io、TileOPs-nightly;tilelang-metax(10-26)、tilelang-musa(09-30)、TileRT、tilescale、DeepStack、tilelang-puzzles 無視窗內主分支提交 |
| 合入真值檢索(is:pr is:merged) | 12 筆:tilelang 5、tilelang-ascend 5、TileFoundry 1、TileOPs 1 |
| 在審真值檢索(is:pr is:open updated) | 54 筆,含昇騰文件核查批次、NPUIR 三筆檢視修正與 TileOPs 的 GEMM 排程改動 |
| 新增 issue 核對 | 4 條:#3704、#3705、#3707 三條缺陷,另 1 條為昇騰倉每日測試機器人失敗通知 |
| 主分支提交級核查 | tilelang 5 筆、tilelang-ascend 5 筆、TileOPs 1 筆、TileFoundry 1 筆、文件站 5 筆;tilelang 與兩倉合入數一致 |
| 釋出頁與 tag | tilelang 最新仍為 v0.1.15(09-30 釋出),視窗內無新版本;TileOPs 最新為 v0.1.0(10-27);TileFoundry 最新為 0.0.4rc1(10-25);tilelang-ascend 最新為 v0.1.2.000(09-09) |
| 第三方與生態倉 | 10 個倉庫逐一核對:昇騰社群運算元倉、PTO-ISA 雙倉、開發工具倉、知識倉、圖解課倉、分散式佈局倉、DCU 融合運算元倉、模型移植倉、Hexagon 後端倉、T-HEAD 分支倉均有視窗內推送(Hexagon 主分支無新提交,僅倉庫推送時間更新) |
| 採用方 | TileKernels 與 FlashQLA 主分支停在 10-14,視窗內無推送 |
| Google News(中英六組) | 僅一條中文加昇騰的臨床人工智慧轉載稿,與主題無關,不收錄 |
| Hacker News | 近四日相關命中全部落在視窗之外 |
| arXiv | 以 TileLang 檢索無新論文 |
完整信源清單
- [1] tilelang 全 warp 歸約用 redux.sync(#3692)— https://github.com/tile-ai/tilelang/pull/3692
- [2] tilelang autotuner 併發編譯獨立配置(#3476)— https://github.com/tile-ai/tilelang/pull/3476
- [3] tilelang 昇騰緩衝區檢視佈局與指標單位(#3623)— https://github.com/tile-ai/tilelang/pull/3623
- [4] tilelang 昇騰越界複製與填充合組(#3543)— https://github.com/tile-ai/tilelang/pull/3543
- [5] tilelang CI 鎖定 PyTorch wheel(#3706)— https://github.com/tile-ai/tilelang/pull/3706
- [6] tilelang 帶填充共享緩衝批次複製缺陷(#3704)— https://github.com/tile-ai/tilelang/issues/3704
- [7] tilelang 負縮放迭代圖反向索引缺陷(#3705)— https://github.com/tile-ai/tilelang/issues/3705
- [8] tilelang 檢視位寬 int32 溢位缺陷(#3707)— https://github.com/tile-ai/tilelang/issues/3707
- [9] tilelang 位寬改 int64 修復(#3708)— https://github.com/tile-ai/tilelang/pull/3708
- [10] tilelang CPU 標量 GEMM 只清目標區域(#3709)— https://github.com/tile-ai/tilelang/pull/3709
- [11] tilelang 關快取時繞過程序內快取(#3710)— https://github.com/tile-ai/tilelang/pull/3710
- [12] tilelang-ascend GLM-5.3 kpool 壓縮運算元(#1910)— https://github.com/tile-ai/tilelang-ascend/pull/1910
- [13] tilelang-ascend Chunk Scaled Dot KKT 運算元(#1803)— https://github.com/tile-ai/tilelang-ascend/pull/1803
- [14] tilelang-ascend W4A4 int4 GEMM 修復(#1731)— https://github.com/tile-ai/tilelang-ascend/pull/1731
- [15] tilelang-ascend depthwise_conv_2d 示例(#1920)— https://github.com/tile-ai/tilelang-ascend/pull/1920
- [16] tilelang-ascend 運算元效能剖析 skill(#1931)— https://github.com/tile-ai/tilelang-ascend/pull/1931
- [17] tilelang-ascend NPUIR 本地 UB 檢視描述符(#1962)— https://github.com/tile-ai/tilelang-ascend/pull/1962
- [18] tilelang-ascend NPUIR 檢視源偏移(#1960)— https://github.com/tile-ai/tilelang-ascend/pull/1960
- [19] tilelang-ascend NPUIR 一元廣播推斷(#1961)— https://github.com/tile-ai/tilelang-ascend/pull/1961
- [20] tilelang-ascend PTO 按位異或工作區檢視(#1959)— https://github.com/tile-ai/tilelang-ascend/pull/1959
- [21] tilelang-hygon HCU 普通複製 LDS swizzle(#30)— https://github.com/tile-ai/tilelang-hygon/pull/30
- [22] TileLang DCU 融合運算元倉 — https://github.com/SealSoul117/TileLang_DCU_FuseCmpComm
- [23] TileOPs 釋出映象(#2793)— <ht
tps://github.com/tile-ai/TileOPs/pull/2793>
- [24] TileOPs GEMM 柵格分組與常駐塊數(#2794)— https://github.com/tile-ai/TileOPs/pull/2794
- [25] TileFoundry AMX 靜態 TIR 契約(#279)— https://github.com/tile-ai/TileFoundry/pull/279
- [26] 昇騰社群運算元倉 ops-tilelang — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [27] PTO-ISA tilelang-puzzles-ascend — https://github.com/PTO-ISA/tilelang-puzzles-ascend
- [28] PTO-ISA tilelangGym — https://github.com/PTO-ISA/tilelangGym
- [29] T-HEAD 名下 tilelang 分支倉 — https://github.com/t-head/tilelang-for-sail
- [30] tilelang 文件站 — https://github.com/tile-ai/tilelang.github.io
- [31] 開發工具倉 tilelang-devkit — https://github.com/easy-tilelang/tilelang-devkit
- [32] 知識倉 tilelang-knowledge — https://github.com/qiuuuuu622/tilelang-knowledge
- [33] 圖解課配套倉 tilelang-course — https://github.com/zachshi-ai/tilelang-course
- [34] 分散式佈局實驗倉 — https://github.com/Leewin239/tilelang-distributed-layout
- [35] Hexagon 後端倉 tilelang-hexagon — https://github.com/xwhzz/tilelang-hexagon
- [36] 深度求索運算元倉 TileKernels — https://github.com/deepseek-ai/TileKernels
- [37] 通義 FlashQLA — https://github.com/QwenLM/FlashQLA
- [38] tilelang 釋出頁 — https://github.com/tile-ai/tilelang/releases
- [39] TileOPs 釋出頁 — https://github.com/tile-ai/TileOPs/releases
- [40] TileFoundry 釋出頁 — https://github.com/tile-ai/TileFoundry/releases
- [41] Google News 檢索 — https://news.google.com/