調研視窗：2026-11-02 07:00 ~ 2026-11-03 07:00（北京時間，24 小時；上一期截點為 11-02 07:00）。

信源：GitHub tile-ai 組織逐倉核查（29 倉庫，視窗內推送 9 個）+ 合入與在審真值檢索（合入 18 筆、在審更新 44 筆）+ 新增 issue 全量核對（8 條）+ 主分支逐筆提交核查（tilelang 5、TileOPs 7、TileFoundry 1、tilelang-ascend 4、tilelang-hygon 1、tilelang-lsp 1、文件站 5）+ 釋出頁與 tag 核查（5 倉）+ 採用方兩倉 + Google News 中英文查詢（走代理）+ Hacker News + arXiv。

本期索引

  • 今日重點：昇騰 950 的持久化 FlashMLA 核心落地（#3678，新增 1997 行、42 次提交）；向量化規劃一處真實迴歸修復（#3734）；#line 指令改為逐行重錨定（#3703）；TileOPs 單日 7 筆效能與測試基建合入；TileFoundry 把共享記憶體排序做成可證明的分析層（#283），並測出 main 的 lowering 在 sm_89 上 7 個程式錯 4 個；24 小時合入 18 筆、在審更新 44 筆、新增 issue 8 條
  • 一、核心專案進展
    • 1.1 主線：昇騰 950 的持久化 FlashMLA 核心，BF16 prefill 與緊湊 FP8 decode 雙路徑（#3678）（11-02）
    • 1.2 主線：cast 解耦被跳過後殘留的 512 位向量，SM100 上 ViT 前向編譯失敗（#3734）（11-02）
    • 1.3 主線：#line 指令改為逐行重錨定，偵錯程式對映可選密集模式（#3703）（11-02）
    • 1.4 主線：注意力示例按裝置共享記憶體上限推 tile，RTX PRO 6000 上由 2 過 7 敗變為 9 全過（#3702）（11-02）
    • 1.5 主線：CUDA float32x8 位轉換改用內建 bit cast，SASS 指令從 19 條降到 18 條（#3748）（11-02）
  • 二、多後端適配
    • 2.1 昇騰：WY Fast 前向運算元新增、wy_fast_bwd_split 反向分段遷移（#1827 / #1828）（11-02）
    • 2.2 昇騰：T.tile.fill 與 T.tile.clear 文件事實核查合入（#1634 / #1628）（11-02）
    • 2.3 昇騰在審：cython wrapper 快取符號形狀解析，承接宿主開銷這條線（#1967）（11-02）
    • 2.4 昇騰在審：z3 版同步從主線遷入，整倉遷主線同步推進（#1965 / #1975）（11-02）
    • 2.5 昇騰在審：歸約運算元兩個缺陷，首 Parallel 軸廣播改標量讀（#1978 / #1968 / #1979）（11-02）
    • 2.6 海光：gfx948 按 gfx946 相容處理，先讓既有指令可測（#32）（11-02）
    • 2.7 摩爾執行緒、沐曦、mlir-ascend 視窗內無推送（11-03）
  • 三、生態與採用方
    • 3.1 TileOPs：單日 7 筆合入，6 筆效能、1 筆測試基建（11-02）
    • 3.2 TileOPs：warp 專用 decode 的 softmax 與 QK/PV 重疊，478 次掃描呼叫幾何均值 0.952 倍（#2817）（11-02）
    • 3.3 TileOPs：MoE 按 token 收集預置換行，H200 上 387.8 TFLOPS（#2815）（11-02）
    • 3.4 TileOPs：16 位逐元素去掉帶範圍檢查的除法，8 位元組運算元一執行緒一向量（#2814 / #2813）（11-02）
    • 3.5 TileOPs：Norm 統計量改為減行內元素，消掉靠近 1e36 的溢位（#2811）（11-02）
    • 3.6 TileOPs：tileops_test 目標讓 tests/ops 脫離 CUDA 跑（#2816）（11-02）
    • 3.7 TileFoundry：共享記憶體排序做成可證明的分析層，83 個夾具計劃 52 拒絕 31（#283）（11-02）
    • 3.8 TileFoundry 在審：CUDA lowering 發出同步計劃並反向校驗，main 上 7 個程式錯 4 個（#285）（11-02）
    • 3.9 採用方：FlashQLA、TileKernels 視窗內無推送（11-03）
  • 四、社群、教程與活動
    • 4.1 昇騰同步正確性成為集中報障面：單日 8 條新 issue，兩名報告人分別系統性提交（11-02）
    • 4.2 跨行與轉置檢視的 SIMD 載入丟掉 DMA 依賴，根因落在邏輯區域重建（#3749）（11-02）
    • 4.3 新聞側：Google News 中英文、Hacker News、arXiv 三通道均無視窗內命中（11-03）
  • 五、趨勢觀察
  • 附、素材與核查說明

今日重點

這一天的主線只有一件事：TileLang 把 DeepSeek 的 MLA 解碼核心搬到了昇騰 950 上，而且是按編譯器可排程的寫法搬的。#3678 一次性加入 1997 行、20 個檔案、42 次提交，覆蓋 BF16 prefill 與緊湊 FP8 decode 兩條路徑，還帶可選的 FP8 與 FP4 額外 KV 快取、attention sink、非法索引與逐 query 獨立長度。實現方式值得注意：注意力迴圈、線上 softmax、反量化與排程全部寫在 TileLang 裡，靠 T.Stage 讓編譯器管理緩衝版本與同步，QK 與 PV 用同一塊 L0B 儲存環的兩個不同 layout 檢視，L1 的落位用 T.annotate_buffer_addresses 直接指定；prefill 在一個累加器分配裡分出四個不相交的輸出切片，歸一化算完一片就排空一片。為了讓這套寫法成立，編譯器側同時補了三處能力——L0C 到 UB 的複製保留 NZ 輸出、每個 AIV 流在自身流水線之前選定、顯式排程的階段在物理發射相位相同時保持原始碼順序。目標形狀寫得很窄：頭數 64、Dqk 與 Dv 都是 512、topk 為 64 的倍數。

第二條線是修一處真實的編譯器迴歸。#3734 處理的是 SM100 上的 ViT 前向：Cannot convert type float32x16 to CUDA type。路徑很長但每一步都清楚——向量化規劃器為一個 256 位的 BF16 儲存選了 16 個 lane，並推遲 cast 約束，指望 DecoupleTypeCast 把計算和訪存拆開；而 #3579 之前剛把「讀記憶體的守衛」從解耦裡排除（重複這些守衛會改變語義），於是最終合法化那一步仍在按推遲策略選寬度，殘留的 FP32 表示式被撐成 512 位向量，CUDA 後端表達不了。修法是把推遲約束做成規劃器的引數而非 IR 屬性，只有最後一步合法化傳 false，守衛的 FP32 轉 BF16 降到 8 個 lane，成功解耦的 BF16 複製仍保留 16 個 lane；地址 cast 也從資料向量型別約束裡摘出來，靠既有的訪問分析判合法性。

第三條線是除錯體驗。#3703 把 #line 指令從「每條 TIR 語句只錨定首行」改成可選的分級發射：級別 0 關閉、1 保持原狀、2 以上對每一行生成的非空行重新錨定，包括被展開成多行的語句和收尾大括號。實現放在 CUDA、HIP、CPU C、C host、Metal、Ascend 共用的 C 系輸出流裡，不動 vendored 的 TVM；沒有原始碼跨度的語句和函式顯式切到 <tilelang-generated>，使用者原文逐字保留以免破壞宏續行與字串字面量。新增環境變數 TILELANG_EMIT_LINE_DIRECTIVES 作為 Python 側的預設值，在快取查詢之前折進 pass 配置。本地 42 項透過、6 項跳過，CUDA 與 HIP 的用例在作者機器上被跳過，cuda-gdb 的實際行為尚未直接驗證——作者自己寫明瞭這一點。同日 LSP 倉跟進一筆，把 CUDA 除錯對齊到新的行對映。

運算元側的節奏沒有停，但方向變了。TileOPs 單日 7 筆合入裡有 6 筆效能、1 筆測試基建，全部來自同一作者；每筆仍然附 H200 實測表與清單迴歸門檻。其中最有價值的是 #2816：新增一個只在測試時存在的 tileops_test 目標，在 PrivateUse1 上為每個運算元算出共享的 workload 參考，讓 tests/ops 的冒煙測試不再必須站在 CUDA 上。改動 1281 行、59 個檔案，作者同時把邊界寫清楚了——比對共享參考只能證明後端無關的契約與可移植性，不能證明真實後端的數值正確性，那條線仍掛在 #2611。

最後是 TileFoundry 的一條暗線。#283 把「共享記憶體訪問應當如何排序」做成分析層：一個同步族按 lowering 的發射順序走每個函式，對每次共享記憶體訪問取指令、發射 mesh、運算元與位元組範圍，在衝突訪問前規劃 commit/wait 與作用域內柵欄。83 個排程夾具裡它規劃 52 個、拒絕 31 個（24 個用 TMA 複製、6 個用 wgmma，兩者的完成語義尚未宣告；1 個需要在跨執行緒邊界不等的迴圈裡插柵欄）。配套的 #285 在審，把計劃變成 CUDA lowering 並反向校驗：sm_89 單卡上能構建的 7 個程式裡 6 個與參考完全一致、synccheck 與 memcheck 全乾淨，而按 main 的方式 lower，同樣是這 7 個裡有 4 個結果錯誤；#283 的描述裡給出的直接症狀是 sm_89 的 E4M3 GEMM 最大偏 659。

一、核心專案進展

1.1 主線：昇騰 950 的持久化 FlashMLA 核心，BF16 prefill 與緊湊 FP8 decode 雙路徑（#3678）

合於 11-02 09:26，新增 1997 行、刪除 17 行，20 個檔案、42 次提交，是本視窗最重的一筆。核心面向昇騰 950，prefill 走 BF16、decode 走緊湊 FP8，可選 FP8 與 FP4 的額外 KV 快取，並支援 attention sink、非法索引與逐 query 的獨立長度。注意力迴圈、線上 softmax、反量化與排程全部用 TileLang 表達，靠 T.Stage 與編譯器管理的緩衝版本和同步驅動；QK 與 PV 共用一塊 L0B 儲存環的兩個不同 layout 檢視，L1 落位由 T.annotate_buffer_addresses 指定；prefill 在既有累加器分配內分出四個不相交輸出切片，並配獨立同步事件，歸一化完成一片即排空一片；decode 用整行輸出複製。編譯器側補三處支援：L0C 到 UB 的複製保留 NZ 輸出、每個 AIV 流在其流水線之前選定、顯式排程階段在物理發射相位打平時保持原始碼順序。目標為頭數 64、Dqk 與 Dv 均 512、topk 為 64 的倍數，pair gather 沿用已記錄的零 burst 程式碼生成繞法；隨附正確性與基準執行器、緩衝複用與 query 邊界迴歸、以及帶 FlashMLA MIT 宣告的 README。

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

1.2 主線：cast 解耦被跳過後殘留的 512 位向量，SM100 上 ViT 前向編譯失敗（#3734）

合於 11-02 09:05。在 e8225d7f5 上，面向 SM100 的 ViT 前向在受守衛的 FP32 轉 BF16 輸出儲存處報 Cannot convert type float32x16 to CUDA type：規劃器為 256 位 BF16 儲存選 16 個 lane 並推遲 cast 約束，期待 DecoupleTypeCast 把計算與訪存分離；#3579 讓解耦跳過那些會讀記憶體的守衛之後，後續 LegalizeVectorizedLoop 仍沿用同一策略繼續推遲，殘留的 FP32 表示式於是變成 512 位向量，CUDA 側無法表示。修法是把 defer_cast_constraints 做成規劃器的 C++ 引數（不是 IR 屬性），由呼叫方按階段選擇：LayoutInference 與 LowerTileOp 維持 true，最終合法化一律傳 false，並把算出的寬度作為顯式 vectorize_hint 前傳，避免重寫器用預設策略重新規劃。守衛的 FP32 轉 BF16 降到 8 個 lane，成功解耦的 BF16 複製仍保留 16 個 lane，守衛安全檢查保留；地址 cast 與訪問指標引數從資料向量型別約束中排除，連續地址降為標量基址加步長，由既有訪問分析判定合法性，int64 索引的對齊向量訪問與四 lane 原子用例不受影響。

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

1.3 主線：#line 指令改為逐行重錨定，偵錯程式對映可選密集模式（#3703）

合於 11-02 04:06，682 行新增、11 個檔案。此前的 #line 發射只錨定一條 TIR 語句生成的第一行，後續行會被編譯器邏輯行號推進，cuda-gdb 於是把生成的載入、運算與儲存歸到無關的 Python 行；合成語句與收尾大括號可能繼承已推進的對映，沒有原始碼跨度的函式還會繼承前一個函式的原始檔。本筆把 tl.emit_line_directives 擴成分級：0 或 False 關閉並保持預設，1 或 True 維持語句級，2 以上對每一行非空生成行重新錨定（含多行展開的語句與收尾大括號）。實現落在 C 系輸出流，覆蓋 CUDA、HIP、CPU C、C host、Metal 與 Ascend，未改動 vendored TVM；同時記錄語句與函式的原始碼跨度，無源區域顯式切到 <tilelang-generated>，使用者原文逐字保留以避免破壞宏續行與字串字面量。新增 TILELANG_EMIT_LINE_DIRECTIVES 作為 Python 側預設值，在快取查詢前折進 pass 配置，顯式配置優先（包括預設值非法時）。本地用 Metal 加 LLVM 構建驗證：行指令與 compile-only 兩組 42 項透過、6 項跳過；去掉指令後，向量化加法與條件選擇兩個夾具在級別 0、1、2 下產物完全一致。作者註明 CUDA 與 HIP 用例在本機被跳過，cuda-gdb 行為未直接測試。

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

1.4 主線：注意力示例按裝置共享記憶體上限推 tile，RTX PRO 6000 上由 2 過 7 敗變為 9 全過（#3702）

合於 11-02 06:59，345 行新增、13 個檔案。多個注意力示例的非調優路徑把 Hopper 尺寸的 tile 寫死，在 sm_86、sm_89、sm_120 上每塊動態共享記憶體上限約 99 KiB，啟動直接失敗：報錯原文是要求把允許的動態共享記憶體設為 163840 位元組，而裝置 0 每塊最多 101376 位元組。現在每個示例從 shared_memory_per_block_optin 推 tile，做法沿用既有的 GQA decode 示例（example_gqa_decode.py）；資料中心配置不變，因為每條分支只在低於上一檔 tile 需求時才觸發。約 99 KiB 裝置上的落點：attention sink 四個示例分別是 64×64 兩階段 128 執行緒（前向）與 64×32 一階段 128 執行緒（反向），example_gqa_bwd.py 前向 64×64、反向 64×32，example_mla_decode.py 取 BLOCK_N 32、BLOCK_H 16。後兩者新增 threads 引數：64 行或 16 行 tile 下 T.gemm 需要四個 warp 才能讓每個 warp 獨佔整塊 16 行，256 執行緒會直接 lower 失敗（分別報 acc_s 與 acc_s_cast 的 layout 衝突、以及 M=16 N=32 無合法 warp 劃分）。另外 flash-attn 2.x 沒有 sm_12x 核心，四個 varlen 示例原本在參考步驟就拋錯，現改為探測一次 flash_attn_varlen_supported，不支援則改用純 PyTorch 的 varlen_attention_ref（覆蓋 GQA、右下對齊的 causal 與 window_size）。驗證在 RTX PRO 6000 Blackwell（sm_120）上：CUDA 9 項全過（原為 2 過 7 敗），CuTeDSL 同樣 9 項全過；H100 有 227 KiB 可選上限，仍走原 tile，同樣 9 項透過。

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

1.5 主線：CUDA float32x8 位轉換改用內建 bit cast，SASS 指令從 19 條降到 18 條（#3748）

合於 11-02 09:58，20 行新增、3 行刪除，修 #3679 並取代 #3681。改動只用 CUDA：把 float32x8 的位重解釋從位運算序列換成 __builtin_bit_cast(uint64_t, make_float2(v, v))，生成的 SASS 指令數從 19 條降到 18 條。幅度不大，但屬於編譯器後端裡典型的「語義等價、產物更短」的一類收尾。

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

二、多後端適配（昇騰 / 海光 / 摩爾執行緒 / 沐曦）

2.1 昇騰：WY Fast 前向運算元新增、wy_fast_bwd_split 反向分段遷移（#1827 / #1828）

兩筆同在 11-02 02:33 合入，前後相差不到一分鐘。#1827 以「遷移並新增」的方式把 WY Fast 運算元搬到 Ascend NPU 側，#1828 遷移並補上 wy_fast_bwd_split 的分段反向。WY 表示法是線性注意力一類模型裡把遞推矩陣乘化簡成低秩形式的常用手段，前向與反向分段配合，等於把這條運算元鏈在昇騰上補齊到可訓練。兩筆的 PR 描述只有模板問候語，改動範圍以標題與提交為準，本報告不對內部實現細節作推斷。

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

2.2 昇騰：T.tile.fill 與 T.tile.clear 文件事實核查合入（#1634 / #1628）

同批合入的還有兩筆文件側的事實核查，分別針對 T.tile.fill 與 T.tile.clear 的 API 行為。#1628 的討論裡留了一句維護約定：這兩個介面比較接近，一邊有修改時請同步看看另一邊是否需要對應調整。這屬於持續進行的 API 事實核查批次，同期仍在審的還有 T.tile.bitwise_not（#1656）、T.tile.bitwise_and（#1651）、T.tile.bitwise_lshift（#1659）與 T.tile.bitwise_or（#1655）。

來源：https://github.com/tile-ai/tilelang-ascend/pull/1634、https://github.com/tile-ai/tilelang-ascend/pull/1628

2.3 昇騰在審：cython wrapper 快取符號形狀解析，承接宿主開銷這條線（#1967）

11-02 更新。#1967 在 cython wrapper 裡快取符號形狀的解析結果，目標是把每次呼叫的 dispatch 開銷繼續壓下來，對應缺陷 #1898。這條線與上一期合入的 #1948 是同一目標的兩個落點：#1948 把每次呼叫無條件構造 TVM Analyzer 的固定開銷改成惰性構造（910B3 上從約 450 微秒降到約 106 微秒），#1967 處理的是符號路徑上重複解析本身的成本。同族還有在審的 #3757，主線側對齊動態 ABI 與共享符號校驗：按 Var 物件而不是列印名做鍵，避免同名不同變數被當成同一符號，並把單例維度上的 stride 只當兜底、不讓它搶先決定取值。

來源：https://github.com/tile-ai/tilelang-ascend/pull/1967、https://github.com/tile-ai/tilelang/pull/3757

2.4 昇騰在審：z3 版同步從主線遷入，整倉遷主線同步推進（#1965 / #1975）

#1965 於 11-02 12:18 更新，把主線的 z3 版 insert sync 遷到 Ascend 倉。#1975 同日 03:28 更新，仍是把 TileLang-Ascend 整體遷到主線多後端架構上。兩條是同一收斂路徑的兩個層面：一條換同步演算法的實現，一條換倉的骨架。與之配套的主線側動作是 #3752（z3 rlimit 配置）在審，以及 #3735 提出在 thread sync 裡證明模組化索引的不相交性——上一期報告裡那筆 ThreadSync 回退留下的證明預算問題，正在從「配置求解資源」與「減少需要證明的形狀」兩頭同時處理。

來源：https://github.com/tile-ai/tilelang-ascend/pull/1965、https://github.com/tile-ai/tilelang-ascend/pull/1975

2.5 昇騰在審：歸約運算元兩個缺陷，首 Parallel 軸廣播改標量讀（#1978 / #1968 / #1979）

三筆都在 11-02 更新。#1978 修 AscendVidReduction 在行切分迴圈裡對整塊 [1,N] UB 錯誤地把 tile 運算元尺寸折半（對應 #1800）；#1968 把 threads=2 時 sigmoid 與 round 的計數折半（對應 #1897）；#1979 讓首 Parallel 軸上的廣播運算元改走標量讀，繞開 VEC 的 UB 對齊陷阱（對應 #1885）。三者都屬同一類問題：昇騰側向量化路徑對形狀與對齊的假設比 CUDA 側更硬，邊界條件一破就是編譯期或執行期的硬錯。

來源：https://github.com/tile-ai/tilelang-ascend/pull/1978、https://github.com/tile-ai/tilelang-ascend/pull/1968

2.6 海光：gfx948 按 gfx946 相容處理，先讓既有指令可測（#32）

合於 11-02 10:46。改動把 gfx948 當作 gfx946 相容，複用既有的 MLS atoms、MMAC、barrier 與編譯選項，目的是在接入 gather load 之前先讓現有 builtins 能在新型號上跑通。這是典型的兩步走：先把新架構號掛進現有路徑拿到可測基線，再補真正需要新指令的部分。該倉上一筆常規複製的 LDS swizzle 修正在上一視窗合入。

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

2.7 摩爾執行緒、沐曦、mlir-ascend 視窗內無推送（11-03）

tilelang-musa 最後推送停在 09-30，tilelang-metax 停在 10-26，tilelang-mlir-ascend 停在 10-31 08:21（早於本視窗起點）。本期的後端動態因此集中在昇騰與海光兩側，其中昇騰佔了合入 18 筆中的 4 筆、在審更新 44 筆中的多數。

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

三、生態與採用方（TileOPs / TileFoundry / TileKernels / FlashQLA）

3.1 TileOPs：單日 7 筆合入，6 筆效能、1 筆測試基建（11-02）

7 筆全部在 11-02 05:10 至 05:22 之間合入，作者同為 lcy-seso，覆蓋注意力、MoE、逐元素、歸一化、測試目標與 CI 契約六個方向。與上一期「一天 13 筆」相比數量減半，但結構變了：上一期是逐 shape 的調優掃尾，這一期出現了測試基建（#2816）與 CI 契約（#2818）兩類非效能改動，說明運算元庫開始為多後端與可維護性做結構性準備。每筆仍然附 H200 實測表，並宣告 manifest 裡沒有已測行回退。

來源：https://github.com/tile-ai/TileOPs/commits/main

3.2 TileOPs：warp 專用 decode 的 softmax 與 QK/PV 重疊，478 次掃描呼叫幾何均值 0.952 倍（#2817）

合於 11-02 05:21。改動把 warp 專用 GQA decode 的 softmax 與 QK、PV 矩陣乘重疊起來：機率留在暫存器裡，K 與 V 使用獨立 barrier，共享一個三槽環形緩衝。選擇條件是 D 為 128、64 個 key 的 tile、每個 KV 頭一行 tile、無 softcap，且「一 CTA 一 SM」的柵格能滿足佔用率與流長度界；不滿足的呼叫繼續走序列程式。九個規劃器邊界都在宿主側用固定裝置事實校驗，不需要真實分配注意力張量、也不要求 132 SM 的卡。實測環境 H200、CUDA 13.2、PyTorch 2.13.0+cu132、TileLang 0.1.12，用 CUPTI 裝置忙時：形狀 (16, 8, 4096, 64, 8, 128) 下 float16 為 0.074610 毫秒對 fa3 的 0.077490（1.039 倍，3.654 TB/s），bfloat16 為 0.074260 對 0.077220（1.040 倍，3.671 TB/s）；受影響的 478 次掃描呼叫記錄的幾何均值延遲為 main 的 0.952 倍。

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

3.3 TileOPs：MoE 按 token 收集預置換行，H200 上 387.8 TFLOPS（#2815）

合於 11-02 05:19。tight 佈局的 MoE 輸入改為按 token 收集，同一條 hidden 行在它的多條路由之間複用；token 填不滿 GPU 時把路由拆到不同塊上。同時刪掉 tight 掃描裡沒人用的專家序源索引分配與寫入。排程上新增一條規則：未融合、K 主序的 tight 組若平均行數大於 128 且不超過 192，選 H200 上的 ping-pong 64×256 tile；顯式給出組大小的呼叫仍走成本模型。實測 (4096, 256, 8, 7168, 2048) 上 bfloat16 為 7.4417 毫秒對 flashinfer-cutlass 的 7.6982（1.034 倍，387.844 TFLOPS），float16 為 7.5040 對 7.9869（1.064 倍，384.624 TFLOPS）；已記錄的路由與偶陣列掃描覆蓋兩種 dtype 與偏置模式，manifest 無超噪聲回退。

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

3.4 TileOPs：16 位逐元素去掉帶範圍檢查的除法，8 位元組運算元一執行緒一向量（#2814 / #2813）

#2814 合於 11-02 05:18，針對 bfloat16 的 Sigmoid 與 16 位 SiLUAndMul 改用四分之一縮放的倒數，繞開帶範圍檢查的除法，並把該輔助函式與 SiLU 共用；同時省掉指數運算裡用不上的次正規數守衛，融合門控輸出跨執行緒按相鄰 16 位元組向量存放，並保留 gate 接近負 88 時乘上無窮值的未舍入 SiLU 啟用。實測 bfloat16 Sigmoid 在 16777216 元素上 0.017504 毫秒對 torch-compile 的 0.018656（1.066 倍，3.834 TB/s），268435456 元素上 0.248545 對 0.262529（1.056 倍，4.320 TB/s）；SiLUAndMul 在 (2048, 14336) 上 0.041856 對 0.045440（1.086 倍）。#2813 合於 05:17，讓 8 位元組運算元在長啟動下按一執行緒一向量處理，改動前先用 TileFoundry 的 roofline 分析確認這類按位或運算元全部是訪存受限：按 4.8 TB/s 的 HBM 頻寬給出各 workload 的下界，hidden-state-prefill 的 int64 為 41.9 微秒、int32 為 21.0 微秒，llama-7b-ffn 的 int64 為 56.4、int32 為 28.2，cnn-feat-broadcast 的 int64 為 42.8，linear-bias 的 int64 為 14.0、int32 為 7.0。

來源：https://github.com/tile-ai/TileOPs/pull/2814、https://github.com/tile-ai/TileOPs/pull/2813

3.5 TileOPs：Norm 統計量改為減行內元素，消掉靠近 1e36 的溢位（#2811）

合於 11-02 05:10。根因寫得很直接：所有均值與方差核心都對原始行做 fp32 求和，和隨行的偏移量增長（約 N 乘 x ），靠近 1e36 的行直接溢位，全部輸出變成 NaN；即使沒溢位，均值也會在偏移量量級上舍入，x − mean 把這個舍入帶進每個元素，於是常數行算出的方差等於舍入的平方，歸一到正負 weight 而不是 bias。修法是給每個由行構成的 fp32 量加一條不變式：它只被行的擴散度約束，不被偏移量約束——實現上取「行內某個元素」作基準再相減，warp 主導的核用 warp 首元素、分片核用行的首元素。這條與上一期報告裡 #2812 列出的四個 Norm 運算元缺陷是同一類問題的正面修復。

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

3.6 TileOPs：tileops_test 目標讓 tests/ops 脫離 CUDA 跑（#2816）

合於 11-02 05:20，是本視窗單筆改動最大的一筆：1281 行新增、524 行刪除、59 個檔案，關閉 #2803。新增一個只在測試時存在的 tileops_test 目標，在 PrivateUse1 上為每個運算元計算共享的 workload 參考，tests/ops 的冒煙測試經由它執行；不支援的參考呼叫被顯式列成 unvalidated 跳過，只在 CUDA 上成立的排除項也在同一目標下校驗。運算元到 workload 的登錄檔透過 workloads/calls 共享，基準專用的繫結、複製記賬與輸入覆蓋留在 benchmarks/_cases；另補測試裝置生成器、卷積支援、修正測試能力標記，並讓參考實現遵守 GQA 反向的 scale 與空行語義。作者把侷限寫在顯眼處：這個目標測的是後端無關的契約與測試可移植性，比對共享參考並不能證明真實後端的數值正確性，#2611 仍然開著；沒有可用參考的呼叫繼續以 unvalidated 形式可見。同批的 #2818 把 dispatch 契約從 workload 資料生成裡拆開，為契約用例只例項化形狀、不生成也不校驗用不到的張量內容，並把五個宿主結構測試套件在 compile-contract-gate 裡只跑一遍。

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

3.7 TileFoundry：共享記憶體排序做成可證明的分析層，83 個夾具計劃 52 拒絕 31（#283）

合於 11-02 02:09。問題是實打實的：finalize 在共享記憶體訪問之間不發射 commit、wait 或 barrier，於是透過共享記憶體暫存的排程程式會競爭，卡片上 sm_89 的 E4M3 GEMM 相對 torch 最大偏 659。#267 定下的方向是「在 HIR 上排訪問順序的分析」加「按目標發射順序的 lowering」，並且只發最小正確形式、證不出來就拒絕；本筆交付的是分析部分。新增的同步族按 lowering 的發射順序走每個函式：對每次共享記憶體訪問，從 analysis/emission.py 取指令、發射 mesh 與運算元，從 IterationScope 取作用域，從記憶體族記錄的落位取位元組數；在衝突訪問之前規劃一次 commit 加一次 wait（針對未完成的 cp.async 寫）與一次覆蓋「同時持有兩者的最內層作用域」的柵欄，按它先執行的子節點的活躍位置編號。運算元透過新登錄檔 visitor_registry/completion.py 宣告效果如何完成，什麼都沒宣告的按 UNKNOWN 處理；該族拒絕 UNKNOWN 之後的訪問、返回時仍在飛行中的 UNKNOWN 訪問、在跨執行緒邊界不等的迴圈裡插柵欄、以及被保留呼叫且被調方持有共享記憶體的情形。預設沒有呼叫方請求該族，因此 75 個會 finalize 的排程夾具位元組與 main 完全相同；83 個夾具裡計劃 52、拒絕 31（24 個用 TMA 複製、6 個用 wgmma，尚未宣告完成語義；1 個需要線上程相關邊界的迴圈裡插柵欄），語料用例逐條斷言了這些拒絕。

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

3.8 TileFoundry 在審：CUDA lowering 發出同步計劃並反向校驗，main 上 7 個程式錯 4 個（#285）

11-02 12:40 更新。這筆是 #283 的執行端：CUDA lowering 按同步計劃發射步驟，並逐條檢查實際發射與計劃一致，不一致拋 LoweringError。語料上的結果：預設流程為 14 個程式加上同步、拒絕 31 個；這 14 個的 TIR 與 main 的差別只有新增語句，另外 30 個與 main 完全相同；經核心 visitor 後，main 會 finalize 的 75 個程式位元組一致。單張 sm_89 上跑了 7 個能構建且帶同步的程式：6 個與參考完全一致，且在 synccheck 與 memcheck 下全部乾淨；同樣這 7 個按 main 的方式 lower，有 4 個結果錯誤。作者同時報出 main 已有的兩個問題：GEMM 與 loop_outputs 上 racecheck 報同一 cp.async 的多個 lane 寫同一地址（HIR 裡的 copy_async 不接受拆分目標，於是每個執行緒拷整塊 tile），loop_outputs 還會把一個輸出寫到錯誤的元素上。

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

3.9 採用方：FlashQLA、TileKernels 視窗內無推送（11-03）

QwenLM/FlashQLA 最後推送停在 10-30 03:02，deepseek-ai/TileKernels 停在 10-14 05:49，兩倉在本視窗都沒有提交。TileKernels 自 10-14 起已連續多期無推送，採用方的可見活動本期全部落在 TileOPs 與 TileFoundry 兩個自家倉裡。

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

四、社群、教程與活動

4.1 昇騰同步正確性成為集中報障面：單日 8 條新 issue，兩名報告人分別系統性提交（11-02）

視窗內主倉新增 8 條 issue，編號集中在 #3740 到 #3749，全部指向昇騰 AscendC 與 PTO 兩個後端的程式碼生成與同步插入。其中 7 條由同一報告人提交，環境為 TileLang 0.1.15、PyTorch 2.10.0+cpu、torch_npu 2.10.0，主題是同一段程式在兩個後端上的行為分歧：生成程式碼在啟動哪些核上不一致（AscendC 出 mix(1,2)、PTO 出 vector）、enable_auto_schedule=True 時 AscendC 不插同步、關閉自動排程時生成成功而開啟後失敗且報錯資訊不清、同一份程式碼 AscendC 能生成而 PTO 失敗。這一組報告的共同點是都附了可復現的 prim_func 與逐行對照的生成程式碼，屬於可直接轉成迴歸用例的形態。加上在審的 #3735（thread sync 裡證明模組化索引不相交）與 #3580（謂詞與提升索引同步），昇騰側的同步正確性已經成為當前最集中的一條線。

來源：https://github.com/tile-ai/tilelang/issues/3747、https://github.com/tile-ai/tilelang/issues/3745

4.2 跨行與轉置檢視的 SIMD 載入丟掉 DMA 依賴，根因落在邏輯區域重建（#3749）

同批裡技術含量最高的一條，報告人環境為 main 分支 e8225d7f5、Ascend950PR、CANN 9.2.0。問題是兩處連續 SIMD 載入丟掉了對更早一次 GM 到 UB 複製的依賴：一是行主序載入從行中間開始並跨到下一行，二是訪問轉置檢視且物理連續維不是最後一個邏輯維。根因路徑寫得很細：生成的 tl.access_ptr 覆蓋物理載入範圍，而 MemoryAccessDetector::LogicalPointerRegion 重建出的邏輯區域漏掉了其中一部分，依賴分析於是認為輸入 DMA 與載入無關，AutoSchedule 把輸入複製挪到載入之後，InsertSync 也不發 MTE2 到 V 的 set/wait。最小復現裡 work.shape 為 (2, 128)，先 T.copy(A, work[1, 0:64])，再在 T.SimdVF() 裡 T.simd.vld(work[0, 64]) 讀 128 個連續半精度元素：生成的讀描述符是 tl.access_ptr(work[0, 64], 128, 1)，物理覆蓋 [64, 192)，其中 [128, 192) 正是剛被 DMA 寫過的部分，而重建出的區域是 work[0:1, 0:128] 即物理 [0, 128)，完全錯過被寫區間；行數按 ceil(128/128) 算成 1 行，起始列 64 沒被計入。這條與本視窗主力在補的 FlashMLA、PTO 向量化是同一塊地基，一旦命中就是靜默錯結果。

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

4.3 新聞側：Google News 中英文、Hacker News、arXiv 三通道均無視窗內命中（11-03）

六組 Google News 查詢（英文 TileLang、TileLang GPU kernel、tile-ai TileLang；中文 TileLang 運算元、TileLang 昇騰、沐曦或海光或摩爾執行緒）返回的最新相關報道是 10-21 的 SemiAnalysis 華為與 DeepSeek 一文，其餘集中在 09-29 至 10-01 的 DeepSeek 開源昇騰元件那一批，全部早於本視窗。arXiv 上 all:"TileLang" 的最新命中為 2026-07-24 的 TileSight，TileLang 與 Ascend 的聯合查詢無結果；Hacker News 的候選條目裡與 TileLang 相關的只有倉庫推送條目。候選清單 14 條中 12 條來自 GitHub 倉庫發現，2 條是無關的 Show HN 與 LSP 倉，本期資訊面完全由倉庫側構成。

來源：https://news.google.com/、https://arxiv.org/

五、趨勢觀察

這一天最值得記的判斷是：昇騰不再是「適配目標」，而是「首發目標」。FlashMLA 這種 DeepSeek 自家解碼路徑上的核心核心，直接以昇騰 950 為目標寫在主線裡，用的還是 T.Stage、annotate_buffer_addresses 這類為顯式排程準備的底層能力，而不是另開一個倉做移植。配合編譯器側那三處支援（L0C 到 UB 保留 NZ、AIV 流在流水線前選定、顯式階段按原始碼順序發射），可以看出昇騰路徑正在從「能跑」轉向「能按排程意圖跑」。

第二條線索是正確性的重心在下移。上一期主線為 TIRX 遷移回退了一整套 ThreadSync 正確性改進，這一期立刻出現兩條補位：主線 #3752 調 z3 求解資源上限，#3735 提出用模組化索引的不相交性減少需要證明的形狀；同時 #3749 把一處「依賴分析看不見物理載入範圍」的缺陷挖到了 MemoryAccessDetector::LogicalPointerRegion 這一層。把這些和 TileFoundry 那條暗線放在一起看——分析層能證明 83 個夾具裡的 52 個共享記憶體排序，而 lowering 按 main 的方式做會讓 7 個程式裡的 4 個算錯——同一個訊號出現了三次：TileLang 現在的主要風險不是功能不夠，而是「生成的程式碼看起來對、實際不對」，而團隊選擇的做法是把正確性做成可拒絕的證明，而不是靠補丁逐個堵。

第三條是測試與契約開始獨立成層。TileOPs 的 tileops_test 把冒煙測試從 CUDA 上摘下來，compile-contract-gate 把 dispatch 契約與 workload 資料生成拆開，兩者都在同一小時內合入；作者同時明確寫了「共享參考不能證明真實後端正確」。這種把「能測什麼」和「證明了什麼」分開表述的習慣，比多合幾筆效能補丁更能說明這個專案正在走出原型期。

最後注意一個反差：運算元側單日 7 筆效能合入、6 筆都帶 H200 實測表格，而新聞側三通道 24 小時零命中。TileLang 的對外聲量仍然集中在「DeepSeek 開源昇騰元件」這類事件上，日常進展幾乎只發生在倉庫裡。對跟蹤這個專案的人來說，這意味著日報的主體只能是提交與 PR，而不是報道。

附、素材與核查說明

信源核查表

信源 核查結果
tile-ai 組織（29 倉庫） 視窗內推送 9 個倉庫
主倉 tilelang 主分支 5 筆合入；新開 issue 8 條；文件站 5 筆自動構建
TileOPs 主分支 7 筆合入；站點 1 筆
TileFoundry 主分支 1 筆合入；1 筆在審新增
tilelang-ascend 主分支 4 筆合入；6 筆在審更新
tilelang-hygon 1 筆合入
tilelang-lsp 1 筆合入（對齊 CUDA 除錯行對映）
TileOPs-nightly 1 次自動推送
tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack 視窗內無推送
釋出頁（tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon） 視窗內無新版本
FlashQLA / TileKernels 視窗內無推送
Google News（中英文六組查詢）/ Hacker News / arXiv 無視窗內實質命中

完整信源清單

ileFoundry/pull/285>