TileLang週報(2026-10-26 ~ 2026-10-30)
報告週期:2026 年 10 月 26 日(週一)至 10 月 30 日(週五),共 5 個工作日 素材來源:本刊每日 TileLang 社群動態日報 5 期(10-26、10-27、10-28、10-29、10-30),五期採集視窗合計 10-23 07:00 至 10-30 07:00(週一那期為 72 小時視窗,覆蓋上週末;週二至週五為 24 小時);另經 tile-ai 組織各倉庫 168 小時視窗補充核查、各倉釋出頁與標籤、VERSION 檔案複核(詳見附錄)
一、本週要聞
-
TileOPs 0.1.0 正式釋出,206 個運算元的首個正式版,公開 API 就此凍結(10-27):PyPI 製品於 10-27 07:45(北京時間)上傳,GitHub 釋出頁 13:15 掛出,是倉庫首個正式 tag。釋出說明給出的硬指標是三組:206 個運算元全部有 manifest 條目(簽名、測試與基準負載、roofline 公式,CI 逐條校驗);H200 上 2265 行夜間基準中 2262 行不慢於最快對照實現、幾何平均 1.37 倍,對照方包括 torch.compile、FlashAttention 2 與 3、FlashInfer、FlashMLA、vLLM、DeepGEMM、flash-linear-attention、FlagGems 與 QuACK;206 個運算元裡 195 筆帶 foundry 標籤的合入由 TileFoundry 生成核心並透過同一套 CI 進庫。釋出同時凍結公開 API,新增 API 相容性檢查指令碼,任何按上一版寫法呼叫而會中斷的 PR 直接失敗,此後 0.x 只增不改。已知缺口四項與路線圖三步同日公開,後者指向融合運算元、整模型與計算通訊融合。來源:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0、https://pypi.org/project/tileops/
-
昇騰 910 A2/A3 的 NPUIR 後端整條併入 mlir-ascend 主幹,驗收口徑首次寫成數字(10-28 合入):一筆 55 檔案 +21490/-1,Python 側是 tilelang/ascend910 包(backend、adapter、codegen、compiler、execution_backend、jit、pipeline、runtime、target 加 language 下的 kernel、copy、gemm、npuir 方言面),C++ 側是 src/ascend910/target 的開發者與專家兩套發射器、運算元層與 LowerNpuirBlock、NpuLoopVectorize、RewriteFDivToMul 三個 pass,執行時補 npu_launcher.h 與 npu_utils,target 別名接受 ascend910、ascend-910、a2、a3、npuir。作者給出的驗收是 pytest testing/npuir 504 項全透過、examples 72 例中 71 例透過且 0 超時 0 報錯,唯一失敗的一例在舊後端同機同樣失敗,屬環境與容差問題。來源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
-
主倉在審的 TIRx 重構把內建 TVM 換成外部依賴,規模從 1043 檔案擴到 1194 檔案(10-25 起,全周更新):方案是刪掉 3rdparty/tvm 及其子模組條目,把前端、原生張量與作用域 IR、下降 pass、後端流水線與執行適配整體遷到獨立構建釋出的 TIRx,並對未打補丁的上游原始碼釘版本(apache/tvm-tirx 9af5544d,依賴 apache-tvm-tirx==0.1.0a2 與 apache-tvm-ffi==0.1.15.post0),原先帶 +tilelang.1 字尾的整合補丁不再需要。週末到週中該 PR 三度更新,規模由 1043 檔案 +71428/-104176 增至 1194 檔案 +106648/-104497,是全周佇列裡體量最大的一筆,也決定後續 TVM 側修復如何進入主線。至本週五仍在審,未合入。來源:https://github.com/tile-ai/tilelang/pull/3666
-
主倉一週的正確性硬化集中在「編譯透過、執行不報錯、結果錯」這一族(10-26 至 10-30):典型條目是輸出分配流(#3649,三個介面卡接受呼叫方的 stream 卻在當前流上分配 out_idx 輸出,快取分配器按分配流排序複用,H100 上五次跑三次輸出損壞,修後歸零)、全 warp 歸約改走 redux.sync(#3692,四個 32 lane 的 FP32 取大從 20 條 shuffle 變 4 條硬體歸約)、切片填充與負索引的靜態邊界(#3352 的 T.fill 在長度 8 的緩衝上寫 As[6:12] 曾被放行,實際寫到 6 到 11;#3347 的 A[-9] 曾靜默迴繞成 -1)、T.ws(-1) 區域在零個執行緒上執行(#3354,存全部丟棄且無診斷)、ROCm 七個預設數學函式不再降到快速內建(#3621,MI300X 上 sin/cos 誤差從 7.8e-04 降到 6.6e-08)。來源:https://github.com/tile-ai/tilelang/pull/3649、https://github.com/tile-ai/tilelang/pull/3692、https://github.com/tile-ai/tilelang/pull/3621
-
昇騰側的運算元清單按上游模型結構逐個鋪開(10-26 至 10-30):一週內落地的包括對齊 Kimi K3 Stable LatentMoE 門控啟用的 situ_glu(#1891)、GLM-5.3 稀疏注意力 indexer 的 kpool 壓縮運算元(#1910,單 kernel 純 Vector 路線,E4M3 編碼用純整數位運算實現、不依賴 FP8 型別也不做近似)、分塊門控線性注意力的 Chunk Scaled Dot KKT 前向(#1803,靜態四塊展開、零 GM 工作區,L0 與 L1 各三例匹配率 1.0000)、W4A4 的 int4×int4 GEMM 結果全錯根因與修法(#1731,半位元組打包型別被當成 1 位元組,所有複製 stride 與 repeat 次數 2 倍偏差,症狀為 99.6% 不匹配),以及主倉 DeepGEMM 示例的 FP4×FP4 打包輸入(#3589)。來源:https://github.com/tile-ai/tilelang-ascend/pull/1910、https://github.com/tile-ai/tilelang-ascend/pull/1803、https://github.com/tile-ai/tilelang-ascend/pull/1731
-
海光一週四筆主幹動作,從「能跑」轉向「按硬體結構擺」(10-26 至 10-30):暫存器流水線加 warp 發散合入(#18,28 檔案 +5312/-64,跨兩週多收斂,作者給出的結論是 example_gemm 達到與 rocblas 相當的效能)、原生 B16 儲存與動態 WASP GEMM 合入(#28,6 檔案 +340/-230)、scale 緩衝拆成獨立雙 bank 並與 LDS bank 分開(#29,12 檔案 +1389/-314,本視窗軟體側最大的一筆)、普通複製路徑的 LDS swizzle 特性分支進在審(#30,18 檔案 +2000/-84,含五個新測試)。來源:https://github.com/tile-ai/tilelang-hygon/pull/18、https://github.com/tile-ai/tilelang-hygon/pull/29
-
釋出映象把運算元庫與 TileFoundry 首次打進同一映象(10-29):執行器映象在開發映象之上安裝 TileOPs 與排程器的 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
-
生態出現整建制的分支與社群倉,TileLang 被當作自有硬體棧的底座(10-26 至 10-30):T-HEAD 名下的 tilelang-for-sail 分支倉公開 13 個分支,除版本映象外可見 PPU 矩陣共享佈局修復、PPU language 包與執行時編譯特性;CANN 社群的 ops-tilelang 運算元倉 10-20 建倉後持續推送並已用七個運算元串出 DeepSeek V4.1 Flash 的模型片段;高通 Hexagon 後端單日 15 筆,把 HVX 向量化、VTCM 競技場、Crouton 佈局與 DMA 內建一路鋪到對映與歸約;另有 RISC-V 社群樞紐倉、PTO-ISA 題庫與基準雙倉、tilelang-knowledge 知識倉同日推進。來源:https://github.com/t-head/tilelang-for-sail、https://github.com/Atlas-Ascend-Fork/ops-tilelang
二、核心專案進展(主倉 tilelang)
版本與釋出
- 主倉最新標籤與釋出仍為 v0.1.15(09-30 釋出),全周無新標籤、無新發布頁條目,至本週五已滿 30 天。VERSION 檔案核查仍為 0.1.15。
- TileOPs v0.1.0 是本週組織內唯一的新正式版本(10-27),也是 09-30 以來繼 TileFoundry 0.0.3 之後的第二個新標籤。TileFoundry 最新停在 0.0.4rc1(10-25),尚未轉正。
- 適配線標籤全周未動:tilelang-ascend v0.1.2.000(09-09)、tilelang-metax v0.1.15(10-16)、tilelang-musa v0.1.15+musa.1(09-30)、tilelang-mlir-ascend 停在 v0.1.2.020 系列。
- 合入量與節奏:週一那期為 72 小時視窗(10-23 07:00 至 10-26 07:00),合入 78 筆(TileOPs 51、tilelang 10、TileFoundry 10、文件站 4、昇騰倉 1、海光倉 1、Sunrise 1),與上一期週報(10-19 至 10-23)在 10-23 一天上有重疊,故不與後四期相加。週二至週五四個 24 小時視窗合計合入 86 筆:tilelang 47、TileOPs 26、tilelang-ascend 5、TileFoundry 4、tilelang-hygon 3、tilelang-mlir-ascend 1;單日依次為 30、20、24、12 筆,呈前高後低的收斂形態。在審更新全周維持在 45 至 76 筆之間,新開 issue 分別為 19、11、7、6、4 條。
編譯器、語言層與執行路徑
- 全 warp 歸約改走硬體 redux(#3692,10-29 合入,4 檔案 +337/-35):此前即使整組就是一個對齊的滿 warp,普通跨執行緒歸約也照樣展開成 shuffle 樹。主分支 837486a 上四個獨立的 32 lane FP32 取大在 sm_103a 配 CUDA 13.1 編譯出 20 條 shuffle、零條 redux,改後是四條 redux.sync.max.f32、零條 shuffle。做法是複用標量與批次全歸約已有的 warp_reduce 助手,呼叫方只判執行緒幾何。整數歸約在 SM80 及以上走硬體 redux;FP32 的取大取小隻在 sm_100a、sm_100f 與 sm_103a、sm_103f 上可用;窄浮點提升到 FP32 後再進助手,FP64 不降精度仍在 shuffle 路徑。作者明確不聲稱端到端加速,只給指令級對照。
- autotuner 的獨立 JIT 配置改為併發編譯(#3476,10-29 合入,2 檔案 +54/-46):原實現整段編譯期持一把鎖,連只用預設編譯選項的配置也序列,前端、下沉與裝置編譯全部排隊。改法是給每份覆蓋配置用 dataclasses 複製出獨立 JITImpl,預設選項仍走原有惰性路徑以保留核心快取與物件身份,只對共享第一階段鍵的候選序列,除錯輸出的寫盤仍序列。跨後端驗收資料:昇騰 1797 項透過(Python 3.12.14、torch_npu 2.10.0、CANN 9.3.0,AscendC 示例 50 例全過、PTO 示例 44 過 6 跳)、ROCm 7.2 在 gfx942 上 3220 項透過、Metal 95 項透過、CUDA 在 H100 PCIe cu130 上 5500 項透過、CuTeDSL 75 過 17 跳。
- 昇騰緩衝區檢視的佈局與指標單位不再跨別名傳播(#3623,10-29 合入,23 檔案 +1139/-438):由 T.view 或 T.reshape 造出的多個緩衝可共用同一段儲存而要求不同佈局,此前統一按共享的資料變數解析佈局,可能選中錯誤的 Cube 對映或 DMA 地址。本筆把佈局附著到單個 Buffer 檢視上,貫穿佈局推斷、分形儲存歸一化、多緩衝物化與複製和矩陣乘的下沉;裸指標偏移按各自元素單位重對映。同一帶版本分配上,Cube 檢視現在可用不同 MajorK 與 MajorMN 佈局,條件是邏輯步長稠密、基址偏移為零、運算元型別相容且填充後儲存大小相等。
- GM 到 L1 的複製與越界填充合併成一個排程任務(#3543,10-29 合入):兩者共同準備同一塊 L1 資料,此前填充可能被拆成獨立排程任務,同步插入就會多一道 MTE2 管道柵欄。改後消費者等待的是包含填充在內的完整資料,FP8 MX 的 K 尾填充同樣併入並保留任務內條件填充,原來的 MX 延遲填充佇列被刪除。47 項聚焦測試透過。
- ROCm 與執行上下文的一組修正:exp、exp10、log、log2、log10、sin、cos 七個預設函式在 HIP 上不再無論 fast_math 設定都降到快速內建,八個帶下劃線的快速運算元新增最低下沉規則(#3621);自動目標改讀當前 HIP 裝置而非固定 0 號卡(#3594);NVRTC 與 CuTeDSL 按首個輸入張量所在裝置取當前流(#3359);profiler 未指定裝置時改用當前裝置,快取沖刷不再打到 0 號卡(#3357)。
- 語言層與下降的小修:TMA gather/scatter 的 Python 整型行索引轉為 int32,字面量寫法在 Blackwell 上不再編譯崩(#3365);執行期行程數的流水線 epilogue 直接展開,分頁解碼恢復可編譯(該缺陷由下游 TileOPs#1852 報出)(#3356);VectorizeLoop 不再靜默忽略迴圈步長([#3345](https://github.com/tile-ai
/tilelang/pull/3345));T.clamp 的 NaN 傳播下降被回退、回到 min/max 語義(#3628);T.Parallel 的 coalesced_width 統一轉 IntImm(#3374);TMA 描述符只記錄已校驗的名字(#3484);非 restrict 引數的 handle 字尾不再被裁掉(#3366)。
正確性與缺陷修復
- 流與快取身份:輸出在呼叫方流上分配(#3649)與 NVRTC、CuTeDSL 取流(#3359)之後,本視窗把自動調優快取鍵納入 compile_flags(#3338,同一核心換編譯選項不再複用舊二進位制),昇騰側把 Bisheng 構建身份納入二進位制快取鍵(#3691)。兩條獨立路線在同一周做同一件事:快取若只按原始碼與形狀做鍵,換編譯器就會誤命中。
- 同步與記憶體規劃:ThreadSync 把 address_of 記為寫,不透明呼叫後的柵欄補上(#3346,此前 T.address_of(B[0]) 形式在 extern 之後沒有柵欄,而 T.access_ptr(B, “w”) 形式有);T.ws(-1) 改為在所有執行緒上執行(#3354);跨儲存域命名屏障重號(#3643)與 sm_100 的 TMEM 釋放缺 CTA 屏障(#3645 / #3663 / #3664)在周初處理。週三浮出三筆 ThreadSync 錯位柵欄缺陷(#3680):柵欄被放進發散分支、分支裡的柵欄錯誤清掉外層掛起依賴、按執行緒計數迴圈的柵欄位置,配 #3682 與 #3681 兩筆在審修復。
- 靜態邊界與型別:切片填充按區域末端而非長度校驗、負索引在迴繞後仍為負時拒絕(#3352 / #3347);SIMT 原子區域的形狀校驗補上(#3465);稀疏 MMA 的 K 尾塊不再被靜默放過(#3367);atomic_addx4 的目標 dtype 收窄到 float*(#3369);profiler 輸出索引統一校驗(#3486);標量讀的索引係數不再混進 lane 數(#3453,行寬 32768 以上的編譯崩潰關閉);批次 all-reduce 按標量分支解析執行緒區間(#3341,batch>1 時 2048 個輸出全錯);int64 原子 max/min 與 AtomicStore 的 int64、bf16、fp16 崩潰(#3307);歸約的 FP8/FP4 非有限常量寫出非法字面量(#3189)。週四把六個假定冷編譯的測試改走 without_kernel_cache(#3695 / #3696)。
- 新開缺陷的指向:週四開出的三條都落在靜態檢查邊界——帶填充的共享緩衝走批次複製報形狀不符(#3704,末維非 16 對齊的全域性張量拷進帶填充共享緩衝,BK 取 cdiv 乘 16,下沉報 112 對 100)、負縮放迭代圖誤編譯反向索引(#3705,復現環境是 v0.1.15 與 T-HEAD 分支倉原始碼構建,屬上游共有路徑)、檢視位寬按 int32 相乘溢位(#3707,Qwen3.8 的 lm_head 權重 248320 乘 2560 位元組 NVFP4 共 5,085,593,600 位,迴繞成 790,626,304 後拋字面量超限)。同批還有 sm_120 上 13 個示例失敗、IfStmtBinding 指數級編譯、without_kernel_cache 不強制重編(#3701 / #3700 / #3699),以及 CuTeDSL 的 CUDA graph 捕獲為空圖(#3693)。週五三條修復進在審:位寬改 int64、CPU 標量 GEMM 只清目標區域、關快取時繞過兩層程序內快取(#3708 / #3709 / #3710)。
執行時、CI 與工具鏈
- CI 把 ROCm 與 Metal 任務鎖回相容的 PyTorch wheel(#3706,10-29 合入):ROCm 的 CI 透過索引變數暴露匹配的 PyTorch ROCm 源,但工作流級的「不安全最優匹配」策略讓更晚的 PyPI 候選勝出,解析到 PyTorch 2.15.0(CUDA 13.2)而不是 2.14.1 加 ROCm 7.2 構建;同一批 2.15.0 的鋪開也讓 macOS 26.6.2 的 ARM64 執行器拒絕 Apple 半虛擬 GPU。改法三條:按配置的穩定 ROCm 版本設定後端變數、夜間安裝改用首個索引、新增測試前快速失敗的校驗步(確認是 HIP 構建、非 CUDA 構建、且能訪問執行器 GPU)。
- 昇騰 CI 把 AscendC 與 PTO 拆成兩個矩陣作業,可分別重跑(#3687);CI 映象升到 cann9.3.0 周更版(#3630)。
- 貢獻流程治理四連(10-26,全圍繞評審成本):CONTRIBUTING.md 寫明 PR 描述應有問題、方案與取捨,證據貼在其支援的論斷旁(#3671);純文件 PR 跳過 CUDA、ROCm、Metal、Ascend 與 CuTeDSL 的構建與測試,白名單刻意收窄(#3672);提交技能加確認步,先把擬用標題給貢獻者確認(#3673);關掉評審機器人的自動摘要與文件字串覆蓋率檢查,要求把發現錨定在當前 diff 與實際呼叫路徑上(#3674)。
- 語言服務上線實驗性 CUDA 除錯(10-26):tilelang-lsp 推入首個除錯提交,新增 tlsp.debug 啟動器做 Python 行號對映、加 -G 與 -lineinfo 編譯、隔離快取並抓取生成源,把 Nsight 的 cuda-gdb 接進 VS Code,支援 Python 斷點、啟動前檢查與生成 CUDA 源的瀏覽。這是該倉自 8 月以來第一次實質提交。
- 文件站全周持續自動同步:10-26 八筆、10-27 六筆、10-28 兩筆,另有 TileOPs 文件站與本地化倉的常規更新。
評審佇列觀察
- 體量最大的一筆仍是 TIRx 重構(#3666,1194 檔案 +106648/-104497),對固定的上游基線打零補丁,是佇列裡唯一會決定「後續 TVM 側修復以何種節奏進入主線」的改動。
- 昇騰倉的 a2a3 排程器遷移佔 100 檔案 +18453/-666(#1951),是專用倉裡最大的一筆在審,說明下一階段變動量在排程層而不在運算元層。
- PTO 線的在審佇列一週內持續推進:T.Parallel 向量化、AutoSimVF 通道、批次三維 ND2NZ、寫計數多緩衝(#3512 / #3320 / #3689 / #3629),PTO 目標預設走 tvm_ffi 執行後端(#3516),另有三筆檢視與廣播修正(#1962 / #1960 / #1961)與按位異或的無符號位元組工作區檢視修正(#1959)。
- 在審更新數全周在 45 至 76 筆之間,週五回升到 54 筆,與合入量的回落形成對照:本週是「清賬周」而非「開單週」。
三、多後端適配(昇騰 / 海光 / ROCm / Hexagon)
昇騰(Ascend):後端併入主幹,運算元按模型鋪開
- 後端層面:910 A2/A3 的 NPUIR 路徑併入 mlir-ascend 主幹(#207,55 檔案 +21490/-1),模式開關 TILELANG_ASCEND_MODE 取 expert 或 developer、預設 expert;第二筆提交逐條對齊舊後端行為差異,重點是 bf16 合法化(HIVM 不接受 bf16 運算元,改為 fp32 計算後顯式轉換)與動態 extent 的複製下沉。
- 排程與定址:主倉新增 T.annotate_buffer_addresses,允許在排程後的昇騰流水線裡對 UB、L1、L0A、L0B、L0C 與 SS 各池內的分配做精確位元組定址,放入定址的緩衝須先用 T.annotate_buffer_versions 宣告固定份數,校驗覆蓋缺少或不支援的份數、重疊、未對齊、動態尺寸與容量溢位(#3607,14 檔案 +982/-26)。PerCoreTask 候選體允許內部 if/else,條件 L1 填充不再被整體拒絕(#3686);手寫流水線保留配置暫存器相位(#3624)。
- 運算元供給:一週內合入 situ_glu(#1891)、GLM-5.3 的 kpool 壓縮(#1910)、Chunk Scaled Dot KKT 前向(#1803)、W4A4 int4×int4 的根因修復(#1731)、depthwise_conv_2d 示例重提(#1920,改用清單化測試通道以避開併發批次);在審還有 GroupNorm 與 fla_chunk 兩個示例(#1607 / #1953 / #1952)、NPUIR 檢視與廣播三筆(#1962 / #1960 / #1961)。這些運算元的共同點是把量化與融合寫進 kernel 本體,而不是留給宿主側補救。
- 工程口徑:新增運算元效能剖析 skill,用 msprof op 輸出定最佳化方向(#1931);安裝器的 SHMEM 段雙修,空殼 wheel 不再假成功(#1943 / #1946);超過兩維的 GM 到 UB 複製按 DMA 行數摺疊前導維(#1805);Cython 包裝器每次呼叫重建分析器的問題(約 0.28 毫秒)先記錄(#1945)後由在審的惰性構造修掉(#1948)。
- 迴歸狀態:每日全量測試 10-25 報 2903 項全通,10-26 凌晨測試作業報失敗(#1944),10-27 與 10-28 連續兩日 2903 項全透過(#1950)。用例規模已足以當門禁,結論有效期以周計。
- 在審的相容性改動:Carver 按所選 SDK 讀 A3 資源預算,避免把 256×256 的 FP32 累加器放到只有 128 KiB 的 L0C 上(#1955);AI Core Exception Dump 的依賴降到 CANN 9.2.0 並補齊四個必填欄位(#1956);文件改用 T.Tensor、棄用 T.Buffer DSL 介面(#1957)。昇騰倉的文件事實核查四條單條評論最多 46 輪(#1628 / #1655 / #1656 / #1659)。
海光(Hygon):一週四筆,按硬體結構擺放
- 暫存器流水線可與既有的共享記憶體流水線協同,對應的提交與屏障指令按流水級插入,warp 數超過 4 時啟用 warp 發散(#18)。這一筆 10-08 開出、跨兩週多收斂,合入時是先並主幹再走分支的大合併,故行數遠大於同批其他改動。
- 原生 B16 儲存與動態 WASP GEMM(#28):保持 MLS 的 Let 引數發射順序、把 GEMM 的 warp 劃分傳播到矩陣儲存、拒絕打包的 B16 到 B32 輸出別名,同時讓 WASP 維度可動態並按軸推斷對齊。
- scale 緩衝的行規劃改為獨立雙 bank 分割槽、與 LDS bank 分開,同一套行對映同時交給分配、scale 複製與 MMAC 下沉三處,避免三處各自推一遍佈局(#29,12 檔案 +1389/-314)。這一筆接上 #28,海光在兩天裡完成三筆主幹合入。
- 普通複製路徑的 LDS 寫相位與 swizzle 策略進在審,配 b8 與 fp8 兩類資料寬度(#30,18 檔案 +2000/-84,新增五個測試)。
來源:https://github.com/tile-ai/tilelang-hygon/pull/18、https://github.com/tile-ai/tilelang-hygon/pull/28、https://github.com/tile-ai/tilelang-hygon/pull/29、https://github.com/tile-ai/tilelang-hygon/pull/30
ROCm:pip 版 SDK 與浮點語義
- ROCm 10 的 pip 版 SDK 支援合入(#3549,2 檔案 +106/-19):ROCm 10.x 的 rocm/pytorch 映象把 ROCm 以 Python wheel 形式釋出,既沒有 /opt/rocm 目錄樹也沒有 rocm-* 的 apt 包;get_rocm_arch() 此前硬編碼 /opt/rocm 字首,路徑不存在時回落到 gfx900,於是 MI355X(gfx950)上 rocwmma 的靜態斷言直接失敗,設 ROCM_PATH 也無效。修法是按呼叫方字首、find_rocm_path()、torch 的 gcnArchName、gfx900 兜底的順序解析,每次 rocminfo 探測加 10 秒超時,並新增基於 rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12 的 Dockerfile。作者在 MI355X 上實測架構探測從 gfx900 變為 gfx950。
- exp/sin/cos 的精度修正與自動目標、取流、profiler 三處上下文修正見第二章。ROCm 側的 FP4 轉換與 E8M0 型別、pip 版映象的架構探測仍在審(#3665 / #3549 相關配套)。
沐曦(MetaX)/摩爾執行緒(MUSA):視窗內無實質提交
- tilelang-metax 組織檢視有推送時間戳(10-26 15:08),但逐分支核對後無視窗內提交:預設分支 main 最新仍是 09-29 的版本號提升,dev 分支最新 10-16,0.1.14 與更早的版本映象未動,事件流裡也沒有視窗內推送事件。最近一次正式釋出仍是 10-16 的 0.1.15。
- tilelang-musa 全周無視窗內主分支提交,最近標籤仍是 09-30 的 v0.1.15+musa.1。摩爾執行緒與沐曦在本週只出現在核查表裡,不構成條目。
來源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa
第三方:Hexagon DSP 後端
- 10-27 單日 15 筆,把 HVX 向量化、VTCM 競技場、Crouton 佈局與 DMA 內建一路鋪到 Crouton 上的對映與歸約。至 10-29 主分支最新提交仍停在 08-24 的 CI 依賴升級,倉庫推送時間更新但主分支無新程式碼,故只作為社群投入強度的觀察,不計入上游進展。
- 另有 PPU 版零改動包裝器(執行時攔截 import 協議與 site 機制,不修改上游原始碼也不重新打包,無 PPU 對應物的 CUDA 特性顯式報錯而非靜默降級),是後端遷移的第三種形態。
來源:https://github.com/xwhzz/tilelang-hexagon、https://github.com/flytiger-eco/tilelang-wrapper-for-sail
四、生態與採用方
TileOPs:0.1.0 之後的第二週
- 釋出之後的合入節奏明顯放緩:週二 14 筆、週三 4 筆、週四 7 筆、週五 1 筆,重心從「把介面凍住」轉到效能與正確性的收尾。
- 效能側的四處硬數字:選擇器不再貼著 manifest 擬合,改為在忙時(device_busy_ms)上按 160 個稠密 shape、60 個 m 9 至 32 的 shape、264 個 m≤8 的 shape 重新標定,用 softmin 期望對數遺憾做座標下降,端到端忙時遺憾在留出的 manifest 上從 5.79% 降到 0.48%、m≤8 的留出集從 14.18% 降到 1.79%(#2784);非 NT 佈局改走轉置持久核心 coop2_swap,2112×7168×4096 TN 從 195.3 微秒降到 175.3 微秒,同機另一塊 H200 從 174.3 降到 158.5 微秒,均壓過 QuACK 的 187.9 與 167.0 微秒(#2789);cumsum 長行改兩段式種子掃描,(8,65536) 從 12.75 微秒降到 5.18 微秒並修掉 65535 tile 以上的發射失敗(#2790);extrema 歸約改用 PTX 的 max.NaN / min.NaN,amax (2048,4096) bf16 從 8.352 微秒回到 6.976 微秒,快於 FlagGems(#2792)。
- 背靠背口徑被單獨立項:週四新開議題指出單次呼叫下本實現領先 QuACK,但背靠背、功率封頂的訓練場景裡十條大行有六條落後(#2791);隨後在審的 GEMM 改動按功耗封頂重算柵格分組與常駐塊數,DeepSeek V3 預填充的 q-b 形狀上記憶體讀從 168 MB 降到 94 MB(#2794)。
- 正確性與能力補齊:D128 稠密前向拆 key,prefill-rope 從 31.04 微秒降到 17.09(#2777);D128 持久化前向按執行緒重寫並寫出 LSE(#2774);融合加法的 LayerNorm 行跨 CTA 拆分,精度缺陷另開新單(#2768 / #2773);偏移儲存三類修復與 SM90 分組 GEMM、輸出 ABI(#2761 / #2765 / #2762)。
- 釋出形態:釋出映象合入(#2793),運算元庫與 TileFoundry 首次打進同一映象;已知缺口中的 FusedAddLayerNormFwdOp 精度問題在周內仍未被修掉,路線圖第一步的融合運算元尚無合入。
來源:https://github.com/tile-ai/TileOPs/pull/2784、https://github.com/tile-ai/TileOPs/pull/2789、https://github.com/tile-ai/TileOPs/pull/2790、https://github.com/tile-ai/TileOPs/pull/2792
TileFoundry:從三包拆分到靜態契約
- 周初拆成 core / nvidia / apple 三個發行包,測試輸入只裝載一次(#268 / #269);sm_89 補上 E4M3 的 warp MMA,並把能力校驗從「只查 MMA」擴到每條指令,bulk copy、TMA copy 與 mbarrier 形式此前能透過驗證、到 codegen 或 ptxas 才失敗(#261 / #262 / #264)。
- 週中修 cp.async 目標佈局定址與 MLA 反向佈局(#263 / #260 / #257),以及整視窗按運算元本身下發(#273,sm_89 的 GEMM 恢復可編譯)與動態共享記憶體按核心實際用量定尺寸(#271,作者未採納「按目標容量給滿」的建議,理由是靜態與動態共享記憶體同計一個上限、每塊申請多少決定一個 SM 能放幾塊)。
- 週四合入一筆破壞性變更(#279):把 24 個 AMX 執行時運算元族描述成 39 個靜態註冊的 TIR 契約,帶引數域、運算元與佈局模式、副作用集合,以及 M1 到 M4 的目標能力;新增 Python 評估器與原始 C 模擬器作為金標準,補 59 個代表性原子用例覆蓋全部 39 個靜態運算元,高階與中階中間表示用例跑真實的 2、4、8、12 執行緒域並驗證候選的接受與拒絕。驗證為 46 項聚焦測試透過。
- 新開缺陷側記一筆:分析側把共享記憶體峰值算重、且不傳到 TIR(#272),與主倉 ThreadSync 那一族同源。
採用方與社群分支
- 兩個已知採用方全周靜默:DeepSeek 的 TileKernels(2018 星)與通義側的 FlashQLA(742 星)主分支都停在 10-14,視窗內無新推送。本週生態訊號主要來自社群自建倉,而非這兩個內部使用方。
- T-HEAD 的 tilelang-for-sail:09-22 建倉,是上游主倉的分支,公開 13 個分支,除主分支與 v0.1.11 到 v0.1.14 的版本映象外,可見 PPU 矩陣共享佈局修復、PPU language 包、執行時編譯特性與若干缺陷修復分支。主分支最近提交停在 10-22,倉庫推送時間更新到 10-29 19:27;倉內無說明文件也無釋出記錄,故只能作為觀察線索。值得注意的是週五第 #3705 號缺陷正是用該倉原始碼構建復現的——分支使用者比上游更早撞上真實構建問題。
- CANN 社群的 ops-tilelang:10-20 建倉,定位面向昇騰 NPU 的社群高效能運算元倉,統一管理 kernel 實現、自動化用例、標準工作負載與效能基線,README 明確只支援 Ascend 950 及以上、要求 CANN 9.3.0 以上並用社群 weekly 版。已用七個運算元串出 DeepSeek V4.1 Flash 的「Attention 輸出加 FFN 前處理」片段,HIDDEN=5120、MHC=4,B 從 1 到 256 實測 SuperKernel 收益 1.17% 到 2.75%。它與官方 tilelang-ascend 的關係尚不明確,只作生態觀察。
- PTO-ISA 的雙配套倉:題庫倉把 PTO 的廣播載入遷移到 PTOAS 0.70,此前一週在做量化縮放符號澄清與題面改寫;基準倉加 SV2 分塊縮放與逐列分析兩筆提交,依據是庫內絕對最大值指令的變體計數 31.39,附執行時指令吞吐與屋頂線分析,明確不是統一緩衝口徑。
- 分發版:Sunrise 的 0.1.15+sunrise.1.2.0 釋出 PR 於周初合入(#9),是組織內第二例下游分發形態;主倉的 TIRx 遷移若落地,這類分發版需要跟著改依賴宣告,是下週值得跟蹤的連帶影響。
- RISC-V 方向的社群樞紐倉:TongmingLAIC/tilelang-riscv 於 10-23 建倉、10-26 推送,定位是 TileLang 在 RISC-V 上的社群樞紐,聚合編譯器後端、AI 運算元、開發工具與文件。目前仍是空骨架,未見可執行運算元或後端落地。
來源:https://github.com/t-head/tilelang-for-sail、https://github.com/Atlas-Ascend-Fork/ops-tilelang、https://github.com/PTO-ISA/tilelang-puzzles-ascend、https://github.com/TongmingLAIC/tilelang-riscv
五、社群、教程與活動
- 知識資產成為本週社群側的新品類:tilelang-knowledge 於 10-29 建倉,當天四筆提交構成一套面向智慧體的知識底座——倉庫結構、33 個運算元設計案例(每條帶固定出處)、20 張帶出處的程式設計知識卡、7 條有證據支撐的 GPU 最佳化模式。倉庫自述用途是讓智慧體面對一個運算元時能判斷哪些經驗適用、應當如何改程式碼、如何驗證、失敗後往哪裡找。這與 TileFoundry 的運算元契約、昇騰倉的剖析 skill 是同一類嘗試。
- 教程與練習方向持續投入:tilelang-puzzles-ascend 做 PTO 對 Ascend 的量化 kernel 教程;中文 TileLang 運算元面試練習站 9 模組 18 題;圖解課配套倉自述用原創場景加動畫講透每個知識點並隨影片釋出;分散式佈局實驗倉同日新建。
- 文件與本地化:tilelang.github.io 全周隨主倉合入自動同步(10-26 八筆、10-27 六筆、10-28 兩筆),第三方本地化倉 tilelang-docs-l10n 常規更新,中文 RISC-V 社群樞紐倉完成首次內容提交,說明文件同步鏈路正常運轉。
- 討論側:TileScale 使用者詢問所用 TileLang 版本能否升級(#72)。該倉最後推送停在 08-25,文件與示例仍指向較早版本;這類提問通常出現在上游介面進入快速改名期之後——本週主倉在審的 TIRx 替換與 TileOPs 的 26 個運算元類改名,都會放大下游跟版成本。
- 新聞側全周零實質命中:Google News 的中英文六組查詢(覆蓋精確詞加兩日、中文加昇騰、中文加運算元、組織名加兩日、英文加 GPU、中文加編譯器加七日)在五個採集視窗裡只返回一條臨床人工智慧的轉載稿,與運算元語言無關;Hacker News 的相關命中全部落在視窗之外,最近一條是 10-02 的昇騰程式設計棧轉載討論;arXiv 以 TileLang 檢索無新論文。日報主體因此繼續以倉庫與生態倉為準。產業側僅 10-26 有一篇昇騰與 DeepSeek 生態差距收窄的報道,屬二次評述,不含新的技術資料。
六、趨勢觀察
釋出工程從「卡在最後一格」走到「映象化交付」。 週一那期日報還在寫 TileOPs 首發「文件就緒、製品未上傳」,PyPI 上最新仍是 2025-10-09 的 0.0.1.dev2;週二 07:45 製品上傳、13:15 釋出頁掛出,週四釋出映象就把運算元庫、TileFoundry 與三段環境標識(cuda13.2、torch2.13、tilelang0.1.12)鎖在一起。一週之內,這條線從「能不能發」變成「怎麼組合交付」。對一個 200 餘運算元的庫來說,可安裝性比再最佳化一個百分點更能決定外部採用速度,本週的答案是把安裝路徑也做成製品。
缺陷型別高度集中在「編譯透過、執行不報錯、結果錯」,且與硬體或上下文強相關。 輸出在錯誤的流上分配(只在非預設流且有排隊工作時損壞)、T.ws(-1) 在零個執行緒上執行、切片填充按長度而非末端校驗、負索引迴繞後仍為負、address_of 被當成讀——五條都在編譯期無警告、執行期無異常。這類缺陷的代價遠高於編譯錯誤,因為它們汙染的是「測試全過」這個前提本身;週三浮出的六個假定冷編譯的測試與週四的 without_kernel_cache 語義漏洞,恰好說明測試基礎設施也在同一片風險區。本週的應對是把邊界校驗前移到編譯期,方向正確,但#3704、#3705、#3707 三條新開缺陷說明這條邊界還遠未畫完。
昇騰從「搬後端」轉到「補運算元」,並第一次給出可復算的驗收口徑。 週三是 NPUIR 併入主幹(504 項自測、71/72 示例),週四到週五的五筆合入裡四筆是運算元與示例,且都按上游模型的需求排列:GLM-5.3 的 indexer 壓縮、分塊門控注意力的內積、四位元量化矩陣乘、量化視覺模型的深度卷積。共同點是把量化與融合寫進 kernel 本體(E4M3 純整數編碼、四位元打包約定、雙 kernel 加 1 乘 1 快路徑),而不是留給宿主側補救;同時把剖析 skill 與清單化測試通道立成流程。主要矛盾正從「能不能編出來」變成「運算元的數值與效能能不能被穩定驗收」,而 a2a3 排程器遷移的 100 檔案預告了下一階段的變動在排程層。
執行路徑做減法,後端清單做加法。 主倉這一週把全 warp 歸約收進 redux.sync、把 autotuner 的編譯鎖拆細、把 GM 到 L1 的複製與填充併成一個任務、把 buffer 檢視的佈局從共享變數收回到單檢視;與此同時後端清單繼續擴張——ROCm 補 pip 版 SDK 與 FP4、Hexagon 單日 15 筆、PPU 出現零改動包裝器、Intel AMX 在 TileFoundry 裡變成 39 個靜態契約。兩條方向並不矛盾:上游收斂執行路徑,下游各自打包分發,這與 TIRx 把 TVM 變成外部依賴是同一條思路。
「fork 化」是本週最硬的採用訊號,也帶來最實的同步成本。 T-HEAD 把 TileLang 當自有硬體棧底座做 PPU 分支、CANN 社群自建運算元倉、PTO-ISA 把題庫與基準都做成配套倉、RISC-V 出現社群樞紐倉——這些都不在 tile-ai 組織裡,卻都依賴語言層的穩定性,這正是 0.1.0 凍結公開 API 並做成 CI 門禁的意義。代價同樣清楚:週五第 #3705 號缺陷由分支使用者的真實構建先撞上,上游才補修;分支越多,這類「下游先發現、上游後修復」的路徑就越長,主倉需要開始回答分支與上游如何同步的長期問題。
口徑問題開始被當作產品問題處理。 TileOPs 週四把「單次呼叫領先、背靠背落後」單獨立項(#2791),週五的 GEMM 改動按功耗封頂重算柵格分組並把 DRAM 讀從 168 MB 降到 94 MB(#2794);選擇器也從貼 manifest 擬合改為按忙時校準。相比「某行快了多少微秒」,這類改動承認了訓練與推理的真實呼叫形態與基準不同,並把差異寫進排程策略。對一門要給大模型寫核心的語言,這比再堆運算元數量更能決定它是否被長期採用。
附錄:素材與核查說明
信源核查表
| 信源 | 核查結果 |
|---|---|
| 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-mlir-ascend、tilelang-hygon、tilelang.github.io、TileOPs.github.io、TileOPs-nightly;tilelang-metax、tilelang-musa、TileRT、tilescale、DeepStack、tilelang-puzzles 無視窗內主分支提交 |
| 合入真值檢索(is:pr is:merged) | 週一 72 小時視窗 78 筆;週二至週五四個 24 小時視窗合計 86 筆(tilelang 47、TileOPs 26、tilelang-ascend 5、TileFoundry 4、tilelang-hygon 3、tilelang-mlir-ascend 1),單日依次 30、20、24、12 筆 |
| 在審真值檢索(is:pr is:open updated) | 五個視窗分別為 76、49、45、48、54 筆,含 TIRx 重構、a2a3 排程器遷移、昇騰文件核查批次與 TileOPs 的 GEMM 排程改動 |
| 新增 issue 核對 | 五個視窗分別為 19、11、7、6、4 條,其中每日包含昇騰倉測試機器人的失敗通知 |
| 主分支提交級核查 | 各視窗合入數與主分支提交數一致;tilelang 單日 10、14、13、15、5 筆 |
| 釋出頁與標籤 | 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);tilelang-metax 為 v0.1.15(10-16) |
| 第三方與生態倉 | 昇騰社群運算元倉、PTO-ISA 雙倉、開發工具倉、知識倉、圖解課倉、分散式佈局倉、DCU 融合運算元倉、模型移植倉、Hexagon 後端倉、T-HEAD 分支倉、RISC-V 社群樞紐倉均有視窗內推送;Hexagon 主分支無新提交,僅倉庫推送時間更新 |
| 採用方 | TileKernels 與 FlashQLA 主分支停在 10-14,視窗內無推送 |
| Google News(中英六組) | 僅一條中文加昇騰的臨床人工智慧轉載稿與一條產業側評述,均不含新的技術資料,不收錄為條目 |
| Hacker News | 近四日相關命中全部落在視窗之外 |
| arXiv | 以 TileLang 檢索無新論文 |
主要引用
- [1] TileOPs v0.1.0 釋出頁 — https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0
- [2] TileOPs PyPI 製品 — https://pypi.org/project/tileops/
- [3] 昇騰 910 NPUIR 後端併入主幹(#207)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/207
- [4] TIRx 重構(#3666)— https://github.com/tile-ai/tilelang/pull/3666
- [5] 輸出分配流修復(#3649)— https://github.com/tile-ai/tilelang/pull/3649
- [6] 全 warp 歸約改走 redux.sync(#3692)— https://github.com/tile-ai/tilelang/pull/3692
- [7] autotuner 併發編譯(#3476)— https://github.com/tile-ai/tilelang/pull/3476
- [8] ROCm 數學函式精度修正(#3621)— https://github.com/tile-ai/tilelang/pull/3621
- [9] 昇騰緩衝精確定址介面(#3607)— https://github.com/tile-ai/tilelang/pull/3607
- [10] 昇騰 kpool 壓縮運算元(#1910)— https://github.com/tile-ai/tilelang-ascend/pull/1910
- [11] W4A4 int4×int4 根因修復(#1731)— https://github.com/tile-ai/tilelang-ascend/pull/1731
- [12] 海光 scale 緩衝雙 bank(#29)— https://github.com/tile-ai/tilelang-hygon/pull/29
- [13] 海光暫存器流水線(#18)— https://github.com/tile-ai/tilelang-hygon/pull/18
- [14] ROCm 10 pip 版 SDK 支援(#3549)— https://github.com/tile-ai/tilelang/pull/3549
- [15] TileFoundry AMX 靜態契約(#279)— https://github.com/tile-ai/TileFoundry/pull/279
- [16] TileOPs 釋出映象(#2793)— https://github.com/tile-ai/TileOPs/pull/2793
- [17] TileOPs 選擇器按忙時校準(#2784)— https://github.com/tile-ai/TileOPs/pull/2784
- [18] T-HEAD 分支倉 — https://github.com/t-head/tilelang-for-sail
- [19] CANN 社群運算元倉 — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [20] Hexagon 後端倉 — https://github.com/xwhzz/tilelang-hexagon
- [21] RISC-V 社群樞紐倉 — https://github.com/TongmingLAIC/tilelang-riscv
- [22] tilelang-knowledge 知識倉 — https://github.com/qiuuuuu622/tilelang-knowledge
說明:本報告事實、版本號與日期均取自週期內五期日報原文及其對應的一手來源,未作推測性補充;週一那期日報的採集視窗為 72 小時(覆蓋上週末),其合入數含 10-23 一天,與上一期週報存在重疊,故本報告未將其與後四期相加。週期內五個工作日日報齊全,無缺失。