調査期間: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 の 2 グループ、リポジトリ別にグループ化して集計)+ リリースページのタグ確認 + ドキュメントサイトのコミット確認 + 採用側リポジトリの確認(TileKernels / FlashQLA)+ コミュニティの新規リポジトリ検索 + Google News の中国語・英語による複数クエリ(プロキシ経由)+ Hacker News + arXiv + 第三者産業分析レポート。


本期インデックス

  • 今日の重点:TileFoundry 0.0.3 が正式リリース、命令レベルスケジューリングを agent に委ねる;TileOPs は1日で20件をマージ、MLA 逆方向と SM89 小規模パスを同日に补齐;メインリポジトリの8件のマージは「エラーは出ないが結果が信頼できない」コンパイラ欠陥に集中;昇騰 DeepGEMM アライメントは審査中からマージへ(10-21)
    1. コアプロジェクト進捗
      • 1.1 メインリポジトリ:バッチ式 finalize_reducer はこれまで最初の出力ブロックのみを帰約していた(#3070)(10-22)
      • 1.2 メインリポジトリ:BF16→FP32 ベクトル化変換をインライン PTX パッキング lane 変換に変更(#3595)(10-21)
      • 1.3 メインリポジトリ:展開パスは空の stepped ループを保持、ROCm の Select は lane ごとの下降に変更(#2887)(10-21)
      • 1.4 メインリポジトリ:ソフトウェアパイプラインは推移的コピー依存を保持、スレッドドメインが領域競合判定に参加(#2822)(10-21)
      • 1.5 メインリポジトリ:cast のソース/ターゲット/丸めの3組み合わせを事前検証、注釈の暗黙的破棄を拒否(#3350)(10-21)
      • 1.6 メインリポジトリ:arch なしの CUDA target を拒否し、sm_50 を暗黙的に補完しない(#2860)(10-21)
      • 1.7 審査中:シフトの符号性と TIRx ループスコープの修正を tvm fork とメインリポジトリでペアコミット(#3597 / tvm #87 / #86)(10-21)
    1. マルチバックエンド対応(昇騰 / 沐曦 / 海光 / 摩尔線程など)
      • 2.1 昇騰評価:SemiAnalysis の詳細レポートに TileLang Ascend 専章が設けられ、6リポジトリ中当時 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:調査期間内にプッシュなし、3ラインは既存バージョンで停止(10-21)
    1. エコシステムと採用側(TileOPs / TileFoundry / オペレータライブラリ)
      • 3.1 TileFoundry:0.0.3 をタグ付けリリース、命令レベルスケジューリング + リファクタリング分析 + 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)
    1. コミュニティ、チュートリアル、イベント
      • 4.1 コミュニティチュートリアル:中国語《TVM / TileLang Cookbook》が更新、jit から cubin まで解説(10-21)
      • 4.2 コミュニティ新リポジトリ:tilelang-wrapper-for-sail を開設、説明はまだなし(10-21)
      • 4.3 コミュニティツール:リモート自動チューニング agent リポジトリが継続プッシュ(10-21)
      • 4.4 ドキュメントとバージョン:ドキュメントサイトが7回自動更新、メインリポジトリの VERSION は依然 0.1.15(10-22)
    1. トレンド觀察

今日の重点

本ウィンドウで32件をマージ(TileOPs 20、tilelang 8、tilelang-ascend 3、TileOPs ドキュメントサイト 1)。さらに正式リリースが1件:TileFoundry が 10-21 10:26 に 0.0.3 を打った。三つのラインの重心はそれぞれ異なる。メインリポジトリの8件のマージはほぼすべて「エラーは出ないが結果が信頼できない」欠陥の整理——バッチ式リダクションが最初の出力ブロックしか処理しない、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 の節を特別に設け、昇騰サポートを得た6つのリポジトリのうち当時利用可能だったのは TileLang のみという判断を示した。


1. コアプロジェクト進捗

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 メインリポジトリ:展開パスが空の stepped ループを保持、ROCm の Select を逐 lane 下降に変更(#2887)

10-21 19:20 にマージ。ループ展開パスは従来、trip count が 0 の stepped ループを空殻として消し、後続パスが見るプログラム構造を変えていた;修正は展開過程中にこの種の空ループを保持するよう変更。同じコミットで 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 のみに依存し、デバイスアクセスはカーネル起動まで遅延されるためである。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. マルチバックエンド適配(昇騰 / 沐曦 / 海光 / 摩尔線程など)

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 メインリポジトリ昇騰:PTO 上の warp_reduce_max/min の整数符号性修正(#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 レーンで符号性の誤りを直接露呈させる。

出典: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 で、エンジンには表現できない。実測では 2 番目の burst が element 1 ではなく element 8 を読んでいた([127, 1151, 71, 199] のうち 1151 はまさに element 8 の値)。書き込み側に問題はなく、誤りは DMA の読み取り元アドレス列にある。修正は二半分:両側とも packed かつ total ≤ 65535 の場合は 1 本のフラットな単一 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 件の更新があり、3 つのグループに集中している。#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 に 3 回のプッシュがあり、内容は再利用可能な性能特定とチューニングの結論である。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 はもともと同じ 2 つの 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 のまま)の 3 リポジトリはいずれも今調査期間にコミットがなく、新規 tag もない。

出典:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa、https://github.com/tile-ai/tilelang-mlir-ascend


3. エコシステムと採用側(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 と再利用ウィンドウ、ループを意識したバッファライフタイム)、コスト(実行スコープごとに作業量とトラフィックを計上、各アクセス関係が何を読むかを宣言)、構造(定義ファクト、評価コンテキスト、スコープスナップショットを visit context に統一、投影の再利用により各コーパス用例を一度だけ分析)に再構築された。新チュートリアルでは 8192×5120 掛ける 5120×17408 の GEMM を辿り、H200 上の AtomSched 最適 kernel まで到達する。

性能面では 3 組の実数が示されている。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 外積として実装している。2 組の Qwen3.8-27B 実行はいずれの長さでも transformers と同一の先頭 token を生成する。既知の問題は 2 件。wheel が TileFoundry の CUDA ランタイムヘッダーをまだパッケージしておらず、インストールパッケージからは tf.compile で CUDA kernel をコンパイルできない(ソース checkout が必要、次バージョンで修正)。もう 1 件は性能タイムラインが 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)

4 件が一本の線でつながる。#2678(10-21 11:34)は MLADenseBwdOp と MLAVarlenBwdOp の 2 つの op 層を追加し、FA3 の _flash_attn_backward インターフェースに合わせた。#2680(10-21 20:58)は SM90 kernel を補完し、D_qk=192、D_v=128 向けに MLADenseBwdWSKernel と MLAVarlenBwdWSKernel の 2 組の warp 特化実装を提供した。#2688(10-21 23:28)は MLAVarlenPrefillWSFwdKernel の常駐 CTA のデータ取得順を変更した。従来はまず各 head に query tile を 1 つずつ割り当て、最初の tile が終わると head ごとに順番待ちするしかなかったが、今回は同一 head の query tile だけを一度に処理する方式に改めた。#2686(10-22 03:08)は仕上げのクリーンアップで、単一読者定数と helper を使用箇所へ移動し、挙動は変えていない。ドキュメントサイト #78 も同期して 2 つの新 op を Attention ページに書き込んだ。発端は check_api_pages.py がメインリポジトリでエクスポートされる 207 個の op のうち 2 ページが欠落していると報告したことにある。

ソース: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 の重みを保つ必要がある。ついでにベンチマークの歪みを1件修正した——チャンク化プリフィルのベンチマークがこれまで単一 query デコードとして測定されていた。融合 RoPE の末尾ケースでは KV 長を 271/1025 に調整し、分割実行をカバーできるようにした。

ソース:https://github.com/tile-ai/TileOPs/pull/2515

3.4 TileOPs:SM89 小規模パスの集中的な補完(#2681 / #2672 / #2669 / #2654)

4件のマージはいずれも L40S/4090 級の弱点を補うもの。#2681 は FP8 BMM の SM89 上のデフォルトタイルを 64×64 に変更——条件は結果 grid が SM あたり最大1ブロックであることで、大きな 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 は小さいトークン数に対して、カバーされる 8-token アキュムレータのプレフィックスのみを計算・転送し、132 SM の H200、N が 128 で割り切れる、133–225 出力タイル、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)

3件は同一の SM89 パスを巡るもの。#2674 は group48 に既存の tiled top-k プログラムを再利用させる:1 トークン 1 タイル(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)) に従ってからスコアリングタイルにアラインメントする。

ソース: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)

調査期間末の 2 件は依然審査中。#2689 は reduction 向け:edge 軸を折り畳んで一度の発射に収め、fold の末尾ロードを先に発射させる;#2690 は SM89 上で NSA top-k に narrower な列 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. コミュニティ、チュートリアル、イベント

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 ドキュメントとバージョン:ドキュメントサイトが 7 回自動更新、メインリポジトリの VERSION は依然 0.1.15

tilelang.github.io は調査期間内に7回自動更新され(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


5. トレンド観察

コンパイラ欠陥の性質が「計算ミス」から「ターゲット選択ミス」へ拡大。 今期のメインリポジトリの8件のマージのうち、4件が変更したのは「コンパイラが自分が何をしていると思っているか」である:arch なしの CUDA target が静かに sm_50 に退避(#2860)、cast の丸め注釈が分岐ごとに静かに破棄(#3350)、バッチ縮約が最初の出力ブロックのみ処理(#3070)、パイプライン計画が推移的依存を喪失(#2822)。これらの共通点はエラーを報告せず、特定のサイズや特定のループ形態でのみ信頼できない結果を出すこと——クラッシュよりも特定が難しく、プロジェクトが機能拡張よりも静かな正確性を優先していることも示している。

昇騰ラインのボトルネックはメモリモデルとランタイム境界へ移行。 32B 未満のコンパクト行間隔は DMA で表現できず、明示的中止しかできない(#1795);キャッシュ並行性はアトミック公開に変更し削除復元を補完(#1903);PTO の整数符号性はラッパー層で補完する必要(#3582)。これら3件はいずれも「オペレータ追加」ではなく「既存オペレータを境界条件でエラーにしない」ことであり、DeepGEMM ヒューリスティック整合のマージと合わせて、昇騰側は能力カバレッジからエンジニアリング信頼性へと転換しつつある。

分析機関が TileLang を昇騰可用性の判断基準とし始めた。 SemiAnalysis は「6つのリポジトリのうち当時利用可能だったのは TileLang のみ」で DeepSeek の昇騰サポートの完成度を定義しており、これはどの単独の技術ブログよりも TileLang の国産算力ナラティブにおける位置——実験的 DSL からデフォルトのキャリア層へ——をよく示している。同じレポートは逆方向のシグナルも示す——最適化はスーパーノードに集中し、x86 ホストサポートは依然不完全で、汎用性は性能に後れを取っている。

アテンションスタックの重心が推論フォワードから訓練バックワードへ移動。 TileOPs は1日で MLA プレフィルバックワードの op 層、SM90 warp 特化 kernel、データ取得順序最適化を補完し、さらに GQA バックワードのグループ KV head 処理を加え、昇騰側の先月以降のバックワードオペレータ補完(indexer_bwd、SiTU-GLU バックワード、WY fast バックワード)と完全に同方向である。訓練側の TileLang カバレッジは体系的に補完されつつある。

SM89 は「動く」から「チューニング」へ。 FP8 BMM タイル、W4A16 の M 上界、ページド DSA、NSA top-k の4ラインが同日に手を入れられ、いずれもシングルカード実測と「普遍的な優位性を構成しない」といった限定句を付している。コンシューマ向けおよびシングルカード推論シナリオが性能目標の一つとなり、同時に著者のベンチマーク基準に対する自己規制が明らかに厳しくなっている。

スケジューリングと正確性の分業が TileFoundry とメインリポジトリの間で明確化。 メインリポジトリは引き続きセマンティクスと正確性を厳格化し、TileFoundry 0.0.3 は「どの命令を選ぶか」を agent に、「セマンティクスを保証する」ことをコンパイラに委ね、H200 と Apple AMX の2つの例で再現可能な数字を示した。この分業が安定すれば、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 3 リポジトリとも調査期間内にコミットなし(最新はそれぞれ 10-16、09-30、10-19)、新規 tag なし
tile-ai/TileRT / tilescale / DeepStack 3 リポジトリとも調査期間内にコミットなし(最新はそれぞれ 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 調査期間内の産業報道 3 件はHuawei-DeepSeek の協同を巡るもの、SemiAnalysis の一文のみが TileLang 専章を含む;arXiv の最新関連論文は依然 07-24

補足説明:海光線のプッシュはデフォルトブランチにないため、判定方法はまず feat/ ブランチのコミットフローを取得し、次にリポジトリイベントインターフェースで PushEvent の ref と before→head を照合する;TileOPs-nightly の pushed_at 変動も同様の方法で除外する。審査中件数は各リポジトリの is:open 検索総数、調査期間内更新数は updated:>=調査期間起点 の検索総数とする。本調査期間起点(10-21 07:00)以前の半日以内に着地したコミットは前号レポートでカバー済みのため、本期では重複計上しない。

完全な情報源リスト