TileLang デイリーレポート (2026-11-05)
調査期間:2026-11-04 07:00 ~ 2026-11-05 07:00(北京時間、24時間;前回の締め切りは11-04 07:00)。
情報源:GitHub tile-ai 組織のリポジトリ別調査(29リポジトリ、期間内プッシュ8件)+ マージ済み・審査中コミットの実データ検索(マージ済み12件、審査中更新42件)+ 新規issue 4件 + メインブランチのコミット別調査(tilelang 3、TileOPs 7、TileFoundry 1、tilelang-ascend 1、mlir-ascend 2)+ 非デフォルトブランチ調査(tilelang-hygonのfeat/hcu-scaled-castブランチ4件)+ リリースページ・タグ調査(5リポジトリ、期間内に新バージョンなし)+ 採用側2リポジトリ + Google News中国語・英語4組のクエリ(プロキシ経由)+ Hacker News + arXiv。
本期インデックス
- 今日の重点:10-14にマージされた昇騰PTOフレームワーク変更(#3512)が集中処理期間に入り、1日で互いに独立した5つのパッチが出現;TileOPsは1日で性能変更7件をマージし、連休後で最も密度の高い1日となった;メインリポジトリのマージ3件はすべて「計測基準」に関係し、L2フラッシュ書き込みがゼロ書き込みから読み取り専用に変更
-
- コアプロジェクトの進展
- 1.1 メインライン:CuTeDSLのTMAアーキテクチャゲートがsm_103aとsm_107aを受け入れ、B300とgr100で6つのテストケースが復旧(#3772)(11-04)
- 1.2 メインライン:L2フラッシュ書き込みがゼロ書き込みから読み取り専用リダクションに変更、do_benchからfast_flushパラメータを削除(#3534)(11-04)
- 1.3 メインライン:AscendC CIの自動チューニング供給テストケースが2台のデバイスを明示的に同期、非同期コピーエラーがarange失敗に見えなくなる(#3784)(11-04)
- 1.4 審査中:ThreadSyncのwarp一貫性証明にwarpインデックスで構成される組み合わせ述語を追加(#3787)(11-04)
- 1.5 審査中:T.copyの目的領域越境がコンパイル時エラーに変更、全コピーfragmentの包含証明にスレッド帰属検証を追加(#3786 / #3788)(11-04)
- 1.6 審査中:SM120のFP8 GEMMがblock-scaled MMAに切り替え、RTX 5090で8192立方が463から727 TFLOPSに(#3790)(11-04)
- 1.7 審査中:Profilerと自動チューニングの基準検証をグループでコミット、6件を同日提出(#3797 / #3798 / #3799 / #3716 / #3717 / #3718)(11-04)
- コアプロジェクトの進展
-
- マルチバックエンド適配
- 2.1 昇騰PTO:1件のアーキテクチャ変更が4種類のリグレッションを引き起こし、互いに独立した5つのパッチを同日コミット(#3792 / #3793 / #3794 / #3795 / #3796)(11-04)
- 2.2 昇騰:CubeとVectorのコア比が設定可能に、1:2固定を解除(#3783)(11-04)
- 2.3 昇騰:AutoScheduleの周期下限を厳格化、Z3求解時間が10.04%短縮(#3776)(11-04)
- 2.4 昇騰:AI Core異常ダンプがCANN 9.2.0と互換、バージョン要件を9.3から引き下げ(#1956)(11-04)
- 2.5 海光:
feat/hcu-scaled-castブランチ4件、E8M0スケーリング変換と64 bankのLDS読み取り計画(11-04) - 2.6 mlir-ascend:BishengIRツールチェーンアダプタをcontribに移行(#209)(11-04)
- 2.7 沐曦と摩尔線程は期間内にプッシュなし(11-04)
- マルチバックエンド適配
-
- エコシステムと採用側
- 3.1 TileOPs:1日で7件マージ、GQA varlenパッキング、ArgMax/ArgMinのスレッドあたり8ベクトル、デフォルトキャッシュ戦略(11-04)
- 3.2 TileOPs審査中:単項要素ごと演算の単一ウェーブ起動、FP8グループGEMMの連続レイアウト、GEMMのオプションbias(#2829 / #2824 / #2823)(11-04)
- 3.3 TileFoundry:FSDPとHSDPのトレーニングセマンティクスをHIR評価器に導入(#284)(11-04)
- 3.4 TileFoundry審査中:wgmma同期計画(#286)、命名統一のissueを別途起票(#287)(11-04)
- 3.5 採用側とサードパーティ:TileKernelsとFlashQLAはプッシュなし、3つのサードパーティリポジトリが24時間以内に活発(11-04)
- エコシステムと採用側
-
- コミュニティ、チュートリアル、イベント
- 4.1 新規issue 4件:warpサイズ64の証明、PTO符号付きアトミック、CuTeDSLのNVFP4量子化器(11-04)
- 4.2 ニュース側:4組のクエリ、HN、arXivいずれも期間内にヒットなし(11-04)
- コミュニティ、チュートリアル、イベント
-
- トレンド観察
- 付、素材と調査説明
今日の重点
この日の主线は「また何を統合したか」ではなく、「昨日の土台が検収され始めた」ことにある。10-14にマージされた #3512 は T.Parallel を PTO の VMI パイプラインへベクトル化したもので、Ascend バックエンドにおける構造的変更だ。整数型を符号なしセマンティクスの i* から明示的な符号付き si* へ変更し、SIMD の lowering を新たな PTO の検証・合法化・ベクトル化という三段フローに引き継がせ、明示的 SIMD 並列ループの判定も書き直した。三週間後、この変更は実際のオペレータ形状の上で四種類のリグレッションとして次々に踏み抜かれた。整数アトミック操作が pto.si32 を受け付けられずにコンパイル失敗し(#3789 報告、#3792 修正)、T.clear が T.SimdVF 内で浮動小数点のみ対応のブロードキャスト経路に落ち(#3795)、スカラー cast の境界事実が VMI メモリ合法化で早期に束縛され(#3793)、明示的 SIMD の並列ループがスカラー lane アクセスとみなされて拒否された(#3796)。さらに整数アトミックを符号なしキャリアへ適応させる一本(#3794)がある。五つのパッチはすべて主线に基づき、互いに依存しない。作者は説明の中で繰り返し同じ一文を書いている——「This PR is based directly on main and does not depend on the other PTO regression fixes」——これは意図的だ。五本をそれぞれ独立してマージ可能にし、過去のように巨大なロールバックパッケージに溜め込むのを避ける。検証も「通った」ではなく、具体的なケース数を伴うハードウェア実測となった。47、18、16、37 項目で、四つの独立オペレータ(SwiGLU、Top-k、Cache、Engram 逆方向)を代表ケースとしている。
第二は計測口径の自己修正であり、しかも二つのリポジトリが同じ日に同じことを行った。メインリポジトリの #3534 は do_bench の L2 フラッシュを「バッファへゼロを書き込む」から「バッファを一度読む」へ変更した。ゼロ書き込みは L2 に丸ごとダーティ行を残し、次ラウンドの計時中にこれらの行を書き戻す必要があり、キャッシュフラッシュのコストを被測定カーネルに摊分してしまう。 #3534 は cache.view(-1, 4096).sum(dim=1) で分割読み込みし、ついでに fast_flush という公開パラメータを削除した——このパラメータは本来 int32 か int8 でバッファを構築するかを選ぶだけのものだった。Ascend 側の msprof バックエンドはさらに徹底しており、元は SimtVF ループでゼロを UB に書き込んでから WTS_FV ストアしていたが、今は NORMAL_FV でグローバルから UB へ読み込む方式に変更し、型も int32 に固定した。TileOPs の #2822 は同じ発想の双子の変更で、さらにフラッシュ方法を記録し、方法が一致する場合にのみ履歴値と比較する。三つの変更を並べて見ると、一つのシグナルだ。性能数値は「バージョン間で比較可能」であることを要求され始めており、単に「今回が前回より速い」だけでは済まなくなっている。
第三はハードウェア前沿への補票だ。B300 が報告するターゲットアーキテクチャは sm_103a だが、CuTeDSL の TMA ヘルパー関数はホワイトリストでアーキテクチャを門前払いしており、B300 上の六つの CuTeDSL ケースが DSL に入る前に ValueError を投げていた。#3772 は sm_103a と gr100(cc 10.7)が報告する sm_107a を一緒にホワイトリストへ追加し、作者は B300 ノードで 8 通過 1 スキップ、gr100 で 13 通過 1 スキップを再現した。同じ調査期間に #3790 は別の道を取る。SM120 上の FP8 加 fp32 累加の GEMM はずっと SM89 の mma.sync アトミックを使っており、RTX 5090 では cuBLASLt の 68% しか出なかった。ue8m0 単位スケーリング付きの block-scaled MMA に置き換えたところ、8192 立方が 463 から 727 TFLOPS へ上昇し、出力は旧アトミックとビット単位で一致した。二つの性質は異なる——一つは「新しいカードを動かせるようにする」、一つは「新しいカードをフルに動かす」——だがどちらも同じことを指している。アーキテクチャ対応の境界は Hopper から Blackwell と国産アクセラレータの両端へ同時に拡大しつつある。
第四は TileOPs のコミットリズムだ。単日 7 本のマージのうち、5 本が性能、1 本が CI、1 本がドキュメント。GQA varlen は SM90 上で KV グループごとにリクエストをパッキングし、短いリクエストの key/value 読み取り回数を KV ヘッドあたり一回に削減した。ArgMax/ArgMin は一波で滞在できる前提の下で各スレッドに 16 バイトベクトルを八つ抱えさせ、(2048, 4096) の bf16 ArgMin は 8.704 マイクロ秒から 7.393 マイクロ秒へ短縮し、FlagGems の 8.543 を上回った。一元逐次要素と正規化系カーネルは「一度だけ読む」オペランドをデフォルトキャッシュ戦略へ戻し、本当に繰り返し読む箇所だけ evict-last/evict-first を残した。これらの変更の共通点は「より速い」ではなく「なぜ速いかを知っている」ことだ——どの変更も対照表を伴い、自分がまだ比較対象に劣っている行を明示している(Softmax の (32768, 4096) は 121.792 から 131.648 マイクロ秒へ後退した)。
1. コアプロジェクト進展
1.1 メインライン:CuTeDSL の TMA アーキテクチャゲートが sm_103a と sm_107a を受け入れ、B300 と gr100 上の六つのケースが復旧(#3772)
11-04 12:46(北京時間、以下同)にマージされ、今期最初のマージとなった。B300 が報告するターゲットアーキテクチャは sm_103a だが、CuTeDSL の TMA ヘルパー関数はアーキテクチャのホワイトリストを保持しており、リストにない値に遭遇すると DSL に入る前に直接エラーを投げ、そのエラーメッセージには許容される値として sm_90、sm_90a、sm_100a、sm_110、sm_120、sm_120a、sm_121、sm_121a が列挙される。fa69f7fa のメインライン上、nvidia-cutlass-dsl 4.7.0 と CUDA 13.3 を組み合わせた環境で、B300 では六つのテストケースがこれにより失敗した。test_tilelang_jit_cutedsl.py の tma_load パスが三つ(GEMM JIT、マルチストリームカーネル、動的シェイプが各一つ)、そしてもう一つの tcgen05 の CuTeDSL テストファイルの tma_store パスにある三つの tmem_copy_roundtrip バリアントである。DSL 自体は sm_103a を受け入れる——その tcgen05 オペレータは sm_100f ファミリー全体を受け入れる——ため、問題は純粋にこのゲートにある。この項目を追加した後、二つのファイルは B300(cc 10.3、CUDA 13.3、cutlass-dsl 4.8.1)上で 8 パス 1 スキップとなった。同じマージでは sm_107a も追加されている。gr100 の bring-up ボード(cc 10.7)がこの値を報告し、同じゲートに引っかかるためだ。両項目を追加した後、同じファイルは gr100 上で 13 パス 1 スキップとなり、その中には #3685 で導入された gather/scatter ファイルも含まれる。
特筆すべきは、作者が明示した検証の境界である。手元に sm_103 のハードウェアはなく、変更は別の報告者(0z5a、#3685 で最初に報告し bisect を行った)が提供した B300 マシンで検証されたものであり、CI は sm_103a をカバーしていない。このような「コードはマージ済み、検証はマシンを借りて実施」という状況はアーキテクチャ追趕期にますます増えていくため、読者はどの機能が実測の裏付けを持ち、どれが単にホワイトリストに追加されただけなのかを知る必要がある。
出典:https://github.com/tile-ai/tilelang/pull/3772
1.2 メインライン:L2 フラッシュをゼロ書き込みから読み取り専用リダクションに変更、do_bench から fast_flush パラメータを削除(#3534)
11-04 11:27 にマージされ、7 ファイルを変更、36 行追加、49 行削除。うち三箇所が実際の実装で、四箇所はそれに追随するテストである。元々は cache.zero_() であった。cache_size MB のバッファにゼロを書き込んで L2 を追い出す仕組みだが、問題はゼロ書き込み後にバッファがダーティ行で満たされ、次ラウンドの被測定カーネル実行時にこれらの行をメモリへ書き戻す必要があり、その書き戻し帯域が被測定カーネルの時間に算入されてしまうことだった——バッファが大きいほどこの偏差は顕著になる。現在は cache.view(-1, 4096).sum(dim=1, dtype=cache.dtype) に変更されている。4096 列ごとに分割して読み取り、行を L2 に読み込んでリダクション出力を生成するため、ダーティ行の問題は存在しない。分割はブロック全体のリダクションによるグローバル同期を避けるためであり、dtype を保持するのはバッファ全体の int64 変換を一度トリガーしないためである。この変更は同時に bench.py の estimate ループ、torch_bench.py の event、cupti、cudagraph の三パス、そして msprof.py の昇騰専用フラッシュカーネルにも適用されている——後者は元々 SimtVF でゼロを UB に書き込み、WTS_FV でグローバルに書き戻すという典型的な書き込み式フラッシュであったが、現在は NORMAL_FV でグローバルから UB へ読み込む方式に変更され、要素型も int8/int32 の二択から int32 に固定された。
公開インターフェースの変更に注意が必要である。do_bench の fast_flush パラメータが削除された。これは元々 int32 と int8 の二つのバッファ要素型を切り替えるだけのもので(int32 の方がバッファ構築が速い)、現在バッファは int32 に固定されたため、パラメータは意味をなさない。do_bench(..., fast_flush=False) を直接呼び出すスクリプトにとっては破壊的変更であり、リポジトリ内のテストは既に同期修正されているが、外部の利用者にはリリースノートで注意を促す必要がある。
出典:https://github.com/tile-ai/tilelang/pull/3534
1.3 メインライン:AscendC CI の自動チューニング供給テストケースが二台のデバイスを明示的に同期、非同期コピーエラーが arange の失敗に偽装されなくなる(#3784)
11-04 17:01 にマージ。現象は AscendC の CI ジョブにおいて、自動チューニングの set-inputs テストケースのうち [other_npu] というパラメータ化項目が、次の torch.arange() 呼び出し時に srclen=4, dstlen=4 を報告するというものだった——4 バイトの memcpy のサイズ不一致であり、これはまさにクロスデバイスコピープローブの形状である。真の原因はクロスデバイスコピーの失敗が非同期に報告され、エラーが浮上する位置が発生位置より一歩遅れるため、エラー報告箇所が完全に無関係な arange 上に見えていたことだった。修正方法は単純である。プローブの前後にそれぞれデバイス同期(現在のデバイスと npu:other)を追加し、非同期エラーを try ... except RuntimeError のスコープ内に落とし込む。検証はシングルマシンデュアルカードで当該テストケースファイルを実行し、四項目すべてパスし、かつクロスデバイスコピーが実際に非同期失敗した場合、[other_npu] のパラメータ化項目はスキップされ、テストケース全体が失敗することはない。
このコミット自体は小さいが、前号 #3546 の続きに当たる——あのコミットは自動チューニングの入力供給を現在のデバイスに基づいて解析するように変更し、验收時に昇腾デュアルカード上のリグレッションに言及していたが、ここで補うのは同一シナリオの CI 上での安定性である。同一機能を二つの調査期間にわたってそれぞれ一回ずつ修正していることは、昇腾側の自動チューニングが依然として「動くが境界が不安定」な状態にあることを示している。
出典:https://github.com/tile-ai/tilelang/pull/3784
1.4 審査中:ThreadSync の warp 一致性証明に、warp インデックスから組み立てた複合述語を追加(#3787)
このコミットが引き継ぐのは、前号 #3761、#3765 が残した同一領域である:ThreadSync は「バリア参加スレッドが warp 全体で一致している」ことを証明する必要があり、証明できなければ lowering を拒否する。トリガーとなる形状は 192 スレッド下の複合述語 tx < 160 and tx // 32 % 2 == 0 である——warp 0、2、4 は丸ごと参加し、warp 1、3、5 は丸ごと不参加で、分割される warp は一つもなく、本来なら証明可能なはずだが、メインラインは内部エラー「参加スレッドが warp 一致ではない」を直接投げる。作者は比較表を作成し、欠陥範囲を极めて正確に絞り込んだ:単項述語 tx // 32 % 2 == 0、tx % 64 < 32、tx < 160 はメインライン上でいずれも lowering できるが、「境界項と warp インデックスの剰余項を and で組み合わせた」この形態だけが失敗する;一方 tx < 168(実際に第 5 warp を分割する)は両側で正しく拒否される。根本原因は証明器のエンコーディングにある:スレッド番号を二つの仮想スレッドにリネームし、算術ソルバーに linear1 / warp_size == linear2 / warp_size から両側の述語が等しいことを導出させるが、FloorDiv と FloorMod のエンコーディングでは「上界 + 剰余」という組み合わせの下で二つの除算の関係を確立できない。修正方法は threadIdx.x を warp_index * warp_size + lane に書き換え、二つの仮想スレッドが warp インデックスを共有し、それぞれが自由な lane を持つようにする——まさに「同一 warp」という関係そのものである;これにより warp インデックスのみを通じてスレッド番号を読む述語は両側で同一式に潰れ、一致性が証明可能となり、本当に lane に依存する述語は依然として等しくならず、 #3786 に列挙された tx < 168 は引き続き拒否される——証明は緩和されておらず、見落とされていた一類型を補っただけである。
同一作者は同日に #3786 も提出した:T.copy の目的領域がバッファを超える場合、越境アクセスを生成するのではなくコンパイル期エラーに変更する。二つのコミットを合わせると同一の指向である——「コンパイラが黙って誤ったコードを生成する」を「コンパイラが否と言う」に変える。
出典:https://github.com/tile-ai/tilelang/pull/3787、https://github.com/tile-ai/tilelang/pull/3786
1.5 審査中:全複製 fragment の包含証明にスレッド帰属検証を追加(#3788)
ProveFragmentContains は「大 fragment が小 fragment を覆っているか」を判断する際に高速パスを持つ:大 fragment が完全複製であれば、直ちに真を返す。このパスはスレッド帰属を確認する前に戻ってしまうため、次のような誤判定が生じる:小 fragment がスレッド 0 から 127 を所有し、大 fragment が完全複製だがスレッド 0 から 95 のみを所有する場合、96 から 127 は全く覆われていないのに、関数は包含が成立すると言う;check_forward_index を有効にすると、完全複製だが物理スロットが異なる二つの fragment も受け入れられる。結果としてレイアウト推論が非互換のマッピングを互換と見なし、「エラーが静かに見逃される」類型に属する。報告者は #2399 を調べる際に偶発的に発見した。修正方法はまず物理インデックスを見て、次に小 fragment のスレッドマッピングを大 fragment の座標系に換算し(有界なスレッド範囲オフセットを考慮する必要がある)、算術アナライザーで複製範囲とスレッド範囲内に収まることを証明し、証明できなければ偽を返す。検証は A100 上で上流メインラインの新規ビルドに対して行い、テスト期にコンパイルされたブリッジを通じて実際の C++ 実装を直接呼び出し、スレッド数 128 対 96 の拒否、96 対 128 と 96 対 96 の受け入れ、物理スロット、非連続帰属、有界範囲、記号インデックス、部分複製をカバーする。
出典:https://github.com/tile-ai/tilelang/pull/3788
1.6 審査中:SM120 の FP8 GEMM を block-scaled MMA に切り替え、RTX 5090 上で 8192 立方が 463 から 727 TFLOPS に(#3790)
これは今回の性能面で最も重要な項目です。問題は命令選択にあります。SM120 上で FP8 と fp32 累加を行う T.gemm は、RTX 5090、8192 立方の形状で約 465 TFLOPS しか出ず、cuBLASLt は 681 です。tile やパイプライン段数を調整しても無駄で、ボトルネックは命令そのものにあります。TileLang が発行するのは SM89 時代の mma.sync ... f32.e4m3.e4m3.f32 で、この原子はこのカード上では半分の速度しか出ません。著者はレジスタのみを測定するプローブを一组作成しました。e4m3 から f32 へ mma.sync 経由で 512 TFLOPS、kind::mxf8f6f4.block_scale(単位スケーリング)経由で 1021、f16 累加の mma.sync 経由で 1022 です。差はすべて命令にあります。ue8m0 スケーリングが 0x7f(すなわち 2 の 0 乗)を取る場合、block-scaled と通常の MMA の結果は同じになるため、TL_SM120_BLOCK_SCALED_MMA_ENABLED が有効でターゲットが sm_120a/f、sm_121a/f の場合、4 つの fp32 累加 FP8 ディスパッチパスを block-scaled 原子に切り替え、その他のターゲット、f16 累加、明示的に非 a サフィックスの sm_120 はそのまま維持できます。ビット単位の一致は命令レベルで検証済みです。e4m3 と e5m2 の 4 通りの組み合わせ、ランダムなレジスタ内容に NaN と Inf のエンコーディングを含みます。
数値:e4m3 GEMM、bf16 出力、128×128×128 tile、3 段パイプライン、256 スレッドの条件下で、4096 立方は 408 から 673 へ(cuBLASLt 530)、8192 立方は 463 から 727 へ(681)、4096×14336×4096 は 449 から 696 へ(498)、256×8192×8192 は 348 から 606 へ(457)。唯一遅くなるのは M=16 で 128 行 tile の形状で、31.5 マイクロ秒から 33.5 マイクロ秒へ後退しますが、著者はこの形状は帯域制限を受けており、16 または 64 行 tile に変えれば速くなると説明しています。追加されたテストケースは PTX を直接検査します。sm_120a は block-scaled 原子を取得し、sm_89 は通常の原子を取得しなければなりません。
出典:https://github.com/tile-ai/tilelang/pull/3790
1.7 審査中:Profiler と自動チューニングの口径検証がグループでコミット、6 件が同日提出(#3797 / #3798 / #3799 / #3716 / #3717 / #3718)
6 件はすべて同一の著者(Afloat16)によるもので、テーマは「検証器と被検証対象の口径」を揃えることです。assert_allclose は現在、カーネルがテンソルを返すべき位置で None を返すことを拒否します。従来の実装はどちらか一方が None であれば比較をスキップしていたため、空を返すカーネルも「通過」できました(#3797)。assert_consistent は従来、戻り値に対して直接 zip を行っていたため、スカラーテンソルは TypeError を投げ、出力が 1 つ欠けたり多かったりしても黙って無視され、形状が異なれば行の切り詰めやブロードキャストでごまかされる可能性がありました。現在はまず出力リストに正規化してから数と形状を照合し、空テンソルも対象に含めます(#3798)。CUPTI ベンチマークの suppress_stdout_stderr は従来、Python ストリームに対して fileno() を呼んでいたため、ノートブックや StringIO にはファイル記述子がなくクラッシュしました。現在はネイティブ記述子 1/2 と Python ストリームをそれぞれリダイレクトし、ExitStack で例外パスでも確実に復元します(#3799)。残り 3 件は自動チューニング側です。目標辞書をキャッシュキーとする際に文字列表現を使っていたため、挿入順が異なる等価な辞書が異なるキーを生成していましたが、ソート後の JSON シリアライズに変更しました(#3716)。Tensor Core 検出は明示的な演算能力を優先し、次に明示的または現在のターゲットアーキテクチャを使い、最後にローカルデバイスを確認します。そうでなければ GPU なしのクロスコンパイルで非対応と誤報されます(#3718)。Profiler の動的形状置換は m + 1 のような複合次元を補い、欠落変数に対しては TypeError ではなく ValueError を返します(#3717)。
このグループの技術的な重要性は変更自体ではなく、それが露呈した形状にあります。検証器は長らく「通過すれば正しい」という薄弱な環節であり、自動チューニングのキャッシュキーやターゲット解決といった基盤インフラに偏差があると、「結果が再現できない」という形で現れ、原因特定が非常に難しくなります。
出典:https://github.com/tile-ai/tilelang/pull/3797、https://github.com/tile-ai/tilelang/pull/3798、https://github.com/tile-ai/tilelang/pull/3799
2. マルチバックエンド適配
2.1 昇騰 PTO:1 件のアーキテクチャ変更が 4 種類のリグレッションを引き起こし、相互に依存しない 5 つのパッチが同日コミット(#3792 / #3793 / #3794 / #3795 / #3796)
まず背景を明確にする。10-14 にマージされた #3512([Ascend][PTO] Vectorize T.Parallel to PTO VMI)は、並列ループのベクトル化を旧 SIMD チャネルから PTO の VMI パイプラインへ移行した。整数型は符号なしの i* から明示的な符号付き si* に変更され、Lowering は Verify、Legalize、Vectorize の三段階を通り、明示的 SIMD ループの判定も再実装された。これら五つのパッチはそれぞれ、それが実オペレータ上で踏み抜いた一箇所に対応している:
- #3792 整数アトミックのコンパイル失敗:
codegen_ptoは整数に対してpto.si32/pto.si64を発行するが、PTO のアトミックインターフェースは符号なしのi32/i64しか受け付けないため、T.atomic_add、T.atomic_max、T.atomic_minは整数バッファ上で直接コンパイル失敗する(issue #3789 に対応)。修正方法はtl.atomic_*の外側に互換ラッパーを被せることである:ポインタ要素型とオペランドを同幅の符号なしに書き換え、整数 min/max には正しい符号性を渡し、戻り値は元の型に戻す。テストケースは 15 項から 18 項に拡大され、符号付きと符号なしの 32/64 ビットの加減算・極値取得、負オペランド、前値返却、32 ビットを超えるカウンタをカバーする。 - #3794 符号付きキャリアの適応:#3792 と同源だが別の道を取る——
tl.atomic_*のセマンティクスは変えず、生成される PTODSL 呼び出し境界でポインタビューとオペランドを書き換え、アドレス空間を保持し、浮動小数点アトミックは元のパスを維持する。代表的なテストケースは SwiGLU の clamp カウント(bf16、hidden 2048、clamp 0.5、psum レイアウト)で、18 項のアトミックケースがすべて通過する。 - #3795 SIMD フィルの消失:
T.SimdVF内の int32T.clearが浮動小数点のみをサポートするブロードキャストパスに到達した。原因は新しい PTO パイプラインが旧 SIMD pass を実行しなくなったことであり、その pass はフィルと並列ループの両方を担当していた。修正方法はフィルのみを行うエントリを抽出し、既存のLowerFillを再利用して、バッファバージョン正規化の後、PTO 並列検証の前に配置し、元の vdup/vsts シーケンスを引き続き発行し、連続領域、アラインメント、整数ベクトル、およびマルチバッファバージョンスライスのチェックを保持する。16 項のテストケースが通過し、代表的なテストケースは Top-k(512 token、1 エキスパート、top-1)。 - #3793 スカラー cast の境界事実が早期にバインドされる:
u = uint32(i - 1)、iが[0, 63]にあるような VMI アドレスと無関係なスカラーは、メモリ合法化でエラーを投げる;性急なBindは過小評価された区間を残し、IntSet の cast 分析で再びエラーを投げる。修正方法は整数 cast の表現可能性を先にチェックしてから事実をインストールし、証明エンジンに安全なスカラー定義を、IntSet に保守的な数値区間を与え、cast 端点を直接求めることを避ける。47 項のテストケースが通過し、代表的なテストケースは Cache(block_size=64、n=65536)。 - #3796 明示的 SIMD 並列ループがスカラー lane として扱われる:Engram 逆伝播の初期化は Parallel ネストを使用し、その中で明示的 SIMD ストアが整数ベクトルブロック単位で進み、新しい lane レイアウト推論がそれをスカラーアクセスとして扱い拒否する(
M=2, N=32, L=128)。修正方法は明示的 SIMD の分類をレイアウト推論と並列ユニット発見の間で共有し、この種のネストは #3512 以前の汎用レイアウトパスを保持し、自動要素単位ユニットは新しい PTO lane パスを引き続き通す。37 項のテストケースが通過し、代表的なテストケースは Engram 逆伝播(512 token、hc 4、hidden 2048)。著者は同時に一つの制限を明記している:FP32/128 lane の明示的+自動混合テストケースは生成ソースレベルのカバレッジのみで、現在の PTOAS/Bisheng デバイスコンパイラはこの組み合わせで 6 KiB を超える VF スタック予算を超過し、デバイス実行結果はない。
これら五つを並べて見ると、昇騰側のエンジニアリングリズムが読み取れる:#3512 のような構造的変更が先に着地し、その後実オペレータごとに境界が露呈し、各パッチはハードウェア実測とテストケースカウントを伴い、かつ意図的に互いに独立を保っている。「まずフレームワークをマージし、後にリグレッションを清掃する」このやり方は前回のアーキテクチャ移行(#1975)でも一度使われており、今では通常フローとなっている——代償はメインブランチがこの種のウィンドウで「既知の数箇所が動作しない」状態にあり、外部ユーザーが踏む確率は低くないことである。
情報源:https://github.com/tile-ai/tilelang/pull/3792、https://github.com/tile-ai/tilelang/pull/3793、https://github.com/tile-ai/tilelang/pull/3794、https://github.com/tile-ai/tilelang/pull/3795、https://github.com/tile-ai/tilelang/pull/3796
2.2 昇騰:Cube と Vector のコア比が設定可能になり、1:2 に固定されなくなる(#3783)
昇騰の AI Core は Cube と Vector の 2 種類に分かれ、AscendC は mix(1,2)、mix(1,1)、mix(1,0)、mix(0,1) などのグループ分けをサポートしているが、TileLang はこれまで一律に 1:2 で起動しており、ハードウェアの自由度の上限を一つ固定してしまうことになり、将来コア比の異なるチップへの適応も難しくしていた。この変更はコア比を明示的またはデバイスから導出するパラメータに変える:T.MixedKernel(..., sids=1) または sids=2 は要求されたグループサイズに従って処理し、能力チェックを行う;省略または None を渡した場合は現在のデバイスの vector_core_num / cube_core_num を用い、デバイス情報が取得できなければ 2 に戻す;T.Kernel は元のインターフェースとコア種別判定を維持し、混合カーネルと識別された場合は準備済みのベクトルコア数を用い、単一コアのパスは変わらず、明示的な T.Vector(1/2) のグループ意味論は保持される。選択されたグループ数は sid ドメイン、タスク割り当て、同期、コード生成まで一貫して引き継がれ、デバイスの重複クエリと重複正規化を避ける;競合する宣言やデバイス能力を超える要求は直接拒否される。
出典:https://github.com/tile-ai/tilelang/pull/3783
2.3 昇騰:AutoSchedule の周期下界を引き締め、Z3 求解時間が 10.04% 減少(#3776)
AutoSchedule は二分探索でパイプラインの開始間隔(II)を求めるが、従来の下界は全タスクの II の最大値で緩く、二分の反復回数が多かった。この変更はリソースモデルの下で三つの必要条件を補った:単一タスク界(P >= max_i II_i を維持)、リソースクリーク界(資源が両々衝突するクリークを残差でソートして総和し、P >= sum_i II_i を得る)、自己依存界(S_u - S_u >= wait - d*P から P >= ceil(wait/d) を導出し、バージョン依存の距離はパイプライン段数で打ち切る)。三者の最大値で二分を初期化し、その他のロジックは変更しない。テスト環境は昇騰 950PR、CANN 9.3.0 週次更新版で、キャッシュを無効化し、ウォームアップ 1 回+測定 3 回を交互に実施:examples/ascend はすべてコンパイルが通りスケジューリング結果も変わらず、Z3 求解段階の累積所要時間が 10.04% 減少した。この種の「結果を変えず探索の起点だけを変える」変更はリスクが最も低く、その利益は全局に及ぶ——コンパイル期の求解時間はすべてのユーザーが支払うコストだからだ。
出典:https://github.com/tile-ai/tilelang/pull/3776
2.4 昇騰:AI Core 例外ダンプが CANN 9.2.0 に互換、バージョン要件を 9.3 から引き下げ(#1956)
この変更は tilelang-ascend リポジトリの ascendc_pto ブランチにマージされ(11-04 10:18)、今期の昇騰リポジトリで唯一のマージである。AI Core 例外ダンプ機能は当初 CANN 9.3.0 以上への依存を宣言していたが、作者は 9.2.0 で実現可能性を検証し三箇所を修正した:acldumpTensorInfo の必須フィールドを補完(addrType=ACL_DUMP_ADDR_RAW、placement、shapeNum、originShapeNum)。9.2.0 は addrType を強制検証するため、calloc によるゼロ初期化で得られる ACL_DUMP_ADDR_PTR は拒否される;空アドレスとゼロサイズのテンソルをスキップ。9.2.0 は tensorAddr が NULL であることを許容しない;登録インターフェースを非推奨の aclrtSetExceptionInfoCallback から aclrtExceptionInfoCallbackRegister に変更。ツール側の parse_exception_dump() は 9.2.0 が出力する .npy と旧版の .bin の両形式に対応し、中英文チュートリアルのバージョン要件も CANN 9.2.0 以上に統一した。まだ 9.2 を使っているユーザーにとって、これは「機能が使えない」から「使える」への変化である。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1956
2.5 海光:feat/hcu-scaled-cast ブランチの 4 コミット、E8M0 スケーリング変換と 64 bank の LDS 読み取りプランニング(11-04)
海光リポジトリのメインブランチは 11-02 で止まっており、今調査期間の動きはすべて機能ブランチ上にあり、4 コミットは 8 時間以内(北京時間 16:26 から 20:16)に集中しており、作者は海光の関係者(huangteng@hygon.cn と zhouweiwang@hygon.cn)である:
feat(hcu): add scaled_cast and unify GEMM LDS read planning:E8M0 スケーリングの FP8/BF8/FP4 変換を追加。スカラーとパックの2種類の lowering を実装し、src/tl_templates/hcu/scaled_cast.h(223行)とlower_scaled_cast.ccを新規追加。同時に GEMM 消費側のメタデータをマトリクス読み出し選択と LDS 戦略の間で共有し、単一レベルの LDS スライスをサポート。メイン変換ファイルの変更は650行。fix(hcu): align GEMM LDS strategies with 64-bank read phases:bank-word の幾何に基づいて plain-copy の swizzle を導出し、マトリクスと2種類のネイティブ pack サイズが共用する狭い b16 tile に明示的な逆変換を提供。混合または制限された位相の非同期消費側では互換性のある公共ストレージを優先。ネイティブ wrap を選択する際は完全な b128 読み出し位相を計上。12ファイル、1536行追加、112行削除、テストファイル5個を新規追加。fix(hcu): prefer mixed LDS wraps proven at most two-way:混合 LDS wrap に「最大2路」の証明を追加し、証明が成立する場合のみ優先使用する。fix(hcu): use physical read width for mixed wrap checks:混合 wrap チェックを物理読み出し幅に変更し、従来のkPackによる8バイト毎レーンの推論を廃止。
テーマは共有メモリの bank 競合と位相アラインメント——これは海光 DCU 上の GEMM 性能のクリティカルパスであり、当該バックエンドが「コンパイル可能」から「マイクロアーキテクチャに基づくチューニング」の段階に入ったことを示す。4件はいずれも機能ブランチ上にあり、PR は未作成。タイミングとしては TileLang メインラインと一度同期してからコミットする可能性がある。
情報源:https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-scaled-cast
2.6 mlir-ascend:BishengIR ツールチェーンアダプタを contrib へ移行(#209)
メインブランチへマージ(11-04 16:35)。BishengIR ツールチェーンのアダプタコードを元の位置から contrib ディレクトリへ移動。これは純粋な構造整理であり、「上流ツールチェーンの進化に伴い頻繁に変更が必要なグルーレイヤー」とコア lowering を分離し、今後の BishengIR アップグレード時の変更範囲を contrib 内に収束させる。同一ブランチには refactor(contrib): move BishengIR toolchain adapters into contrib という主体コミットもあり、マージコミットは2件目。この種の整理は単体では機能価値がないが、「アダプタ層」を独立したライフサイクルで管理する始まりである——昇騰ツールチェーンのバージョン反復は速く、グルーコードの可換性はエレガンスよりも重要だ。
情報源:https://github.com/tile-ai/tilelang-mlir-ascend/pull/209
2.7 沐曦と摩尔線程はウィンドウ内にプッシュなし(11-04)
tile-ai/tilelang-metax の最終プッシュは 10-26 15:08、tile-ai/tilelang-musa は 09-30 17:50 で停止。両リポジトリとも本ウィンドウ内にコミットがなく、レビュー中 PR の更新もない。国産4社のバックエンドのうち、昇騰と海光が本ウィンドウの2つの活発なラインであり、沐曦と摩尔線程は静穏期にある。前回レポートでも沐曦は同様にプッシュなしで、2ウィンドウ連続。摩尔線程は 09-30 以降動きがなく、今後のウィンドウで内部開発へ移行したか注視する価値がある。
情報源:https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-musa
3. エコシステムと採用側
3.1 TileOPs:1日で7件マージ、GQA varlen パッキング、ArgMax/ArgMin スレッド毎8ベクトル、デフォルトキャッシュ戦略(11-04)
7件すべてがウィンドウ内にマージされ、ペースとしては今四半期で最も密な1日。1件ずつ見ていく:
#2819 は SM90 の GQAVarlenFwdOp にパッキングされた warp 特化プログラムを追加:1リクエストの KV ヘッドグループが1 tile 内に位置で満たせる場合、リクエストを KV グループで tile を充填し、各 KV ヘッドが key/value を1回だけ読むようにする。パッキングを行うかどうかはコストモデルで判断し、「形態が節約を保証する量」が実測の9%単ヘッドオーバーヘッドを超える場合のみ使用。広い/狭いリクエストは同一起動内で交互に取得し、ホストリードバックなし、同期追加なし。デバイス上で統計された分割不可項が SM 数より少ない場合、単一 tile のパッキングリクエストは key スキャンを複数セグメントに分割し、fp32 ワークスペースでマージ。さらにメインラインの3つの欠陥を修正:sink 加算スコアスケーリングが約2.4を超えると可視 key のない行が NaN を返す、フォールバックパスがリクエスト最初の key より前に RoPE テーブルを読む、および負の対照が呼び出し元スタックフレームを参照サイクルに残しテスト間の CUDA graph が早期破棄される。テストノードは173から197へ増加。
#2828 行単位で分割された ArgMax/ArgMin において、「一波駐留」の前提で各スレッドに8つの16バイトベクトルを持たせ、64スレッドブロックをチューニング候補に追加した。駐留判定は SM のスレッド数、常駐 CTA 上限、レジスタによって共同で決定される。(2048, 4096) の bf16 ArgMin は 8.704 マイクロ秒から 7.393 に短縮(比較項目 FlagGems は 8.543、1.155 倍高速)、バックツーバック場面では 7.808 から 6.000 に短縮。f16 ArgMax は 6.816 から 6.400 に短縮(比較項目 9.344、1.46 倍高速)。12行のベンチマークのうち4行が高速化、0行が低速化した。
#2820 は「各要素を一度だけ読む」オペランドをデフォルトのキャッシュ戦略に戻した。要素ごと、rope、norm、softmax、リダクション系カーネルは一律に evict-last/evict-first を使わず、本当に繰り返し読む箇所にのみ残した(バイナリカーネルのブロードキャスト行、マルチバッチ・マルチヘッドで共有される rope の cos/sin 行、2パス RMSNorm の2回読む箇所、融合加算の最初の2パス)。同時に RMSNorm にペアカーネルと段階カーネルの2実装を追加し、warp 行 Softmax の CTA を SM の warp 数をちょうど満たすまで縮小、2D neox RoPE を128スレッドに変更し、アーキテクチャごとの常駐 CTA 上限を RESIDENT_CTAS に収めた。代償として Softmax の (32768, 4096) は 121.792 から 131.648 マイクロ秒に後退した(比較項目 quack は 126.080)。著者は表に正直に明記している。
#2822 は L2 フラッシュ書き込みのリードオンリー修正で、メインリポジトリ #3534 と同源。168項目のテストが通過し、さらにフラッシュ書き込み方法を記録し、方法が一致する場合のみ履歴の測定値と比較する。#2825 は単一再フィットのコストモデルで W4A16 の合法構成をランク付けし、split-K の CTA は自身の量化メタデータのみを一時保持し、共有メモリの記帳に起動マージンを算入する。#2826 は付随ドキュメントで、稠密および tiny-M のキャリブレーション契約(測定方式、フィット制約、ホールドアウト検証)を明記し、ランタイムの変更はない。#2827 は off-CUDA スイートを夜間ジョブに移動する。従来 preflight ではこの2ステップが ubuntu-latest で13〜18分かかり、他のチェックより10分遅く結果が出ていたが、うち tests/ops のスモーク実行はオペレータの参照実装を用いており、検査対象はデバイス非依存性とディスパッチであってカーネル結果ではない。移動後、preflight は収集チェック、破棄テストレポート、cuda_only 必失敗チェックのみを残し(約3分)、1856+2789のテストノードを夜間へ移行。リリースゲートは依然として夜間成功を要求する。
出典:https://github.com/tile-ai/TileOPs/pull/2819、https://github.com/tile-ai/TileOPs/pull/2828、https://github.com/tile-ai/TileOPs/pull/2820
3.2 TileOPs レビュー中:単項要素ごとの単一波起動、FP8 グループ GEMM 連続レイアウト、GEMM オプション bias(#2829 / #2824 / #2823)
#2829 は「本体に1〜2個の単一命令操作しかない」単項要素ごとの起動を単一波で実行させる。読み書きバイト数が L2 の4分の3以下、結果とオペランドが同幅、本体が最大2個でそれぞれ1命令に落ちる操作、という3条件を満たす場合のみ256スレッドブロックの一波に縮小する。本体内条件は TIR 変数上のノードカウントで判定し、cast、除算、数学呼び出し、ガード、または裸の変数上で評価できない本体は一律デフォルト起動を維持する。現在該当するのは relu、abs、neg、bitwise_not、および bool 上の logical_not。ついでに実際の問題を1件修正:ReLU が NaN 入力に対して0を返していたが、現在は NaN を返す(16ビットは torch のビットパターン、float32 は規範 NaN)。ベンチマークでは3行の ReLU と1行の LogicalNotFwdOp が高速化し、その他の構成は不変。著者は GPU 6 上のフラッシュ書き込み付き decode 行が単に引き分けであることも明記している。
#2824 は SM90 の FP8 グループ GEMM に連続レイアウトを追加(コンパクトとアラインの2種パッキング、セグメント末尾または行単位の2種エキスパートメタデータ)。SM89 は引き続きマスクレイアウトのみをサポートし、エキスパートサイズで偏らせた実測負荷と DeepGEMM 連続レイアウトとの比較を補完する。#2823 は GEMMFwdOp にオプションの bias[N] を追加し、稠密、GEMV、split-K の3経路をカバー、cuBLASLt BIAS との比較を付す。両者とも #2130 に列挙されたギャップに対応しており、TileOPs が「torch エコシステムと項目ごとに整合させる」リストに沿って進んでいることを示す。ランダムにオペレータを追加しているのではない。
出典:https://github.com/tile-ai/TileOPs/pull/2829、https://github.com/tile-ai/TileOPs/pull/2824、https://github.com/tile-ai/TileOPs/pull/2823
3.3 TileFoundry:FSDP と HSDP の学習セマンティクスを HIR 評価器へ(#284)
11-04 16:36 にマージ。従来の評価器は reshard をグローバルテンソルとして扱い、参加者ごとの貢献の差異を捨てていたため、訓練セマンティクス(特にシャードとレプリカが混在する HSDP)が HIR 層で正しく評価できなかった。この実装 #255 の訓練サブセット:参加者ごとのローカル値、GPU 座標、呼び出しとループ、gather/scatter、および軸に沿った Partial(sum) 還元は、いずれも共有の投影規則を通る;Reshard の別名証明は帰属に敏感な形に変更し、#170 の構造化ライフサイクルに接続し、タプル実引数はテンソルリーフから検証する。併せて modded-nanogpt に触発された 2 ブロック GQA 研究を追加:明示的な前方・後方、モーメンタム SGD、独立した PyTorch リファレンス、FSDP、2×2 と 4×4 の HSDP、2 種の常駐戦略をカバーし、集合通信グループ、ペイロードとストレージ占有を報告し、再利用可能な SVG を埋め込み、README の各表はコードから導出可能である。
検証:リファクタリング後のコアスイート 775 項目、バックエンドスイート 778 項目(13 スキップ、sm90 カード 2 枚)、10 組の HIR 設定が CPU 上で通過;2 種の常駐戦略は 4 枚の sm90 上で実際の 2×2 PyTorch HSDP と整合(更新を 2 回実行し、ローカル勾配、パラメータ、モーメンタムを比較)。
出典:https://github.com/tile-ai/TileFoundry/pull/284
3.4 TileFoundry 審査中:wgmma 同期計画(#286)、別途命名統一の issue を起票(#287)
#286 が扱うのは WGMMA の非同期性である:命令の発行後も共有メモリとレジスタを読み続ける可能性があるため、後続のアクセスはその効果の後に順序付けなければならない。方法は、各オペランド効果に完了条件を宣言し、WGMMA の fence、commit、wait をそれを必要とする HIR ノードに紐付け、lowering が命令の発行点でこれらのステップを消費する;後続スコープが集合待機を収められない場合、発行スコープ内で操作を完了させる。同時にレジスタ Reshard の別名を修正し、各要素がレプリカを保持する各スレッド上で同一のレジスタに留まるようにし、このストレージ同一性で依存関係を追跡する計画である。著者はリスクを明記している:行列命令の発行器はまだ接続されておらず、この変更は WGMMA の実行正確性を証明せず、性能も与えず、TMA の完了セマンティクスは別途扱う。
#287 はドキュメント向け issue で、コードベースに「1 つの概念に 5 つの名前」があることを指摘する(mesh の 1 メンバーが placement 133 箇所、program 110 箇所、instance 92 箇所、participant 44 箇所と呼ばれ、他に wave_units のような unit 付きの書き方がある)、しかも unit という 1 語が空型、ハードウェア計算ユニット、mesh メンバーを同時に指し、Python と CUDA ランタイムで local_layout のシグネチャも一致していない。この種の issue は通常すぐにコードを変えないが、コンパイラプロジェクトが「書ける」から「長く人に読まれる」へ進むための必经ステップである。
出典:https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/issues/287
3.5 採用側とサードパーティ:TileKernels と FlashQLA はプッシュなし、3 つのサードパーティリポジトリが 24 時間以内に活発(11-04)
2 つの採用側リポジトリは本調査期間にコミットがない:deepseek-ai/TileKernels は 10-14 05:49 で停止、QwenLM/FlashQLA は 10-30 03:02 で停止(最新リリースは依然 09-30 の v0.1.3)。注意すべきは、採用側がプッシュしないことは影響されないことを意味しない——前号 #3761 が修正したのはまさに TileKernels の mhc_pre_big_fuse カーネルの B300 におけるバリア占有の回退であり、コンパイラの変更が「受動的リグレッション」の形でそれらの形状に作用することを示している。
サードパーティ側にはリポジトリ発見で拾った活発な項目が 3 つあり、いずれも 24 時間以内に新規作成またはプッシュされた:sablin39/tilelang-cuda-skills(149 スター、TileLang の記述と CUDA ツールチェーンでのデバッグのためのスキルセットを提供)、xwhzz/tilelang-hexagon(2 スター、TileLang を Qualcomm Hexagon DSP に持ち込む)、Index4NekoMeteor/Tilelang-AttnRes(1 スター、TileLang で Kimi K3 のアテンション残差構造を実装)。3 つとも小さいが方向性は興味深い:1 つはツールチェーン側の開発体験、2 つは新目標アーキテクチャと先端モデル構造への移植であり、コミュニティが自発的に TileLang の境界を拡張する類の仕事である。
出典:https://github.com/deepseek-ai/TileKernels、https://github.com/QwenLM/FlashQLA、https://github.com/sablin39/tilelang-cuda-skills
4. コミュニティ、チュートリアルとイベント
4.1 新規 issue 4 件:warp サイズ 64 の証明、PTO の符号付きアトミック、CuTeDSL の NVFP4 量子化器(11-04)
メインリポジトリの3件。#3791 は「[BUG][ice-on-valid-code]」:ThreadSync は warp サイズが 32 の場合にのみ warp 一致述語を証明でき、同じ IR が thread_warp_size 64 では拒否される——これは今期審査中の #3787 と極めて密接に関係しており、証明器の warp サイズに対する仮定がハードコードされていることを示している。一方、AMD と一部の国産アクセラレータは 64 幅 warp である。#3789 は PTO バックエンドが整数の T.atomic_add に対して符号付きの pto.si32/si64 を発行し、検証に失敗することを報告しており、当日中に #3792 と #3794 という2方向の修正案が出ている。#3781 は CuTeDSL の NVFP4 block-scaled 量子化器がコンパイルできないことを報告:FP4 パッキングの再解釈と欠落した ceil/floor である。これに TileFoundry の #287(命名統一)を加え、本調査期間は合計4件の新規 issue であり、すべてが「内部実装の詳細がユーザー可視面に漏れている」タイプで、性能不満は1件もない。
情報源:https://github.com/tile-ai/tilelang/issues/3791、https://github.com/tile-ai/tilelang/issues/3789、https://github.com/tile-ai/tilelang/issues/3781
4.2 ニュース側:4組のクエリ、HN と arXiv はいずれも調査期間内にヒットなし(11-04)
Google News の4組のクエリ(英語 TileLang GPU kernel、"TileLang" OR "tile-lang"、中国語 TileLang 算子、TileLang 编译器)の中で、最新のヒットは352時間前(10-21 の SemiAnalysis による華為と DeepSeek の報道)で止まっており、残りは09-30の「DeepSeek が昇騰版 TileLang、DeepGEMM、DeepEP をオープンソース化」という報道の波(856時間前)に集中しており、すべて本調査期間より前である。Hacker News で tilelang を検索した直近24時間は0件。arXiv の all:tilelang の最新ヒットは依然として07-24の TileSight であり、調査期間内に新論文はない。GitHub リポジトリ発見で拾った9件の候補のうち、7件はリポジトリ自身、1件は Swarm World のような主題と無関係な Show HN、1件は無関係な2048系プロジェクトである。本調査期間のニュース側はゼロであることが確認でき、レポート本体はコミットと PR からのみ得られる——これはすでに複数期連続の安定した形態である。
情報源:https://news.google.com/、https://news.ycombinator.com/、https://arxiv.org/
5. トレンド観察
第一に、昇騰 PTO の清算パターンはすでに定着しており、その利益とコストの両方を記録する価値がある。#3512 は10-14にマージされ、3週間の間に実際のオペレータが次々とその境界を踏み越えており、本日の5件の補完はただ最新のバッチにすぎない。この「まずアーキテクチャを落とし、後からオペレータごとに回帰を清算する」リズムは #1975 のリポジトリ全体移行時にも見られた。利点はアーキテクチャの反復が速く、単一の変更が制御可能なこと、欠点はメインブランチがかなり長い期間「既知の不動作箇所が数か所ある」状態に置かれることである。5つのパッチが意図的に相互依存を避けているのは、メンテナがロールバックの粒度を意識的に下げていることを示す——これは大きなロールバックパッケージにまとめるより成熟しているが、同時により強力なリリース前回帰規律が必要であることを意味する。そうでなければ「独立してマージ可能」が「独立して漏れる」に変わってしまう。
第二に、性能測定が製品特性として扱われ始めている。メインリポジトリの #3534 と TileOPs の #2822 は同じ日に L2 フラッシュをゼロ書きから読み取り専用に変更しており、理由はいずれも「ダーティ行の書き戻しが次ラウンドの計時を汚染する」である。TileOPs はさらに進んで、フラッシュ方法を記録し、方法が一致する場合にのみ過去の読み値と比較する。これは小さな修正ではない:ベンチマーク数値があるバージョン間で比較不能であれば、「どれだけ速くなったか」というすべての結論を再計算しなければならない。測定方法そのものをバージョン管理に取り込むことは、性能エンジニアリングが信頼に足るものになる前提である。
第三に、オペレータライブラリの競争点は「オペレータ数」から「セレクタの品質」へ移っている。TileOPs が本日マージした W4A16 単一再フィット代价モデルと稠密 GEMM のキャリブレーション契約文書、さらに審査中の FP8 グループ GEMM 連続レイアウトと GEMM bias は、明確な一本の線を構成する:まず各形状の合法な配置を列挙し、次に統一された代价モデルで順位付けし、キャリブレーション過程を再現可能な契約として書き下す。この線が正しく引ければ、オペレータライブラリの性能は個別カーネルの手動チューニングに依存しなくなる。誤れば「ある形状は特に速く、他の形状は運任せ」に退化する。
第四に、アーキテクチャサポートの境界が両端で拡大している。一端は新しいカード:B300 の sm_103a、gr100 の sm_107a がホワイトリストに入り、SM120 の FP8 は block-scaled MMA に置き換わる。もう一端は国産アクセラレータ:海光は64バンクの LDS フェーズでマイクロアーキテクチャ級のチューニングを行い、昇騰はコア比や例外ダンプといったエンジニアリング詳細を補完している。この2つの拡大方法は作業量の構造がまったく異なる——新しいカードは主に「パスを通せるようにする」ことであり、国産カードはより「同じパスでハードウェアに近い効率を出させる」ことである。後者ははるかに高い実測密度を必要とし、まさに海光の4件のコミットでテストファイルの比率がこれほど大きい理由である。
最後のリスク面:1つのウィンドウ内で42件の審査中PR更新があり、そのうち12件が昇騰/PTOタグ付き、7件がCUDA/SM/CuTeDSLタグ付き、4件がProfilerに集中している。この構造は、リポジトリの変更圧力が「機能追加」から「境界修正」へ移行しつつあることを示している。ユーザーにとっての良いニュースは、欠陥が1つずつ指摘されていることだ。悪いニュースは、これらの欠陥の多くが検証器、キャッシュキー、証明器といった「普段は見えない」場所に落ちており、それらが顕在化する時にはすでに結果の信頼性に影響を与えていることが多いということだ。
付、素材と検証説明
情報源検証表
| 情報源 | 検証結果 |
|---|---|
| tile-ai組織(29リポジトリ) | ウィンドウ内で8リポジトリにプッシュ |
| メインリポジトリ tilelang | メインブランチ3件マージ;新規issue 3件;審査中更新38件(昇騰/PTO 12、CUDAとCuTeDSL 7、Profiler 4) |
| TileOPs | メインブランチ7件マージ;審査中更新3件;新規issueなし |
| TileFoundry | メインブランチ1件マージ;審査中更新1件;新規issue 1件 |
| tilelang-ascend | デフォルトブランチ ascendc_pto 1件マージ;審査中リストは匿名インターフェースの制限により取得できず、代わりにメインブランチとコミットの原子的情報源で検証 |
| tilelang-mlir-ascend | メインブランチ2件(#209マージを含む) |
| tilelang-hygon | メインブランチに新規コミットなし(11-02で停止);feat/hcu-scaled-castブランチ4件 |
| tilelang-metax / tilelang-musa | ウィンドウ内でプッシュなし |
| TileOPs-nightly | ブランチプッシュのみ、メインブランチに新規コミットなし |
| リリースページ(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) | ウィンドウ内で新バージョンなし、tilelangの最新は依然として09-30のv0.1.15 |
| FlashQLA / TileKernels | ウィンドウ内でプッシュなし |
| Google News(中英文4組クエリ)/ Hacker News / arXiv | ウィンドウ内で実質的ヒットなし、最新ヒットは352時間前 |
完全情報源リスト
- [1] tilelang CuTeDSL が sm_103a と sm_107a を受け入れ(#3772)— https://github.com/tile-ai/tilelang/pull/3772
- [2] tilelang L2 フラッシュを読み取り専用リダクションに変更(#3534)— https://github.com/tile-ai/tilelang/pull/3534
- [3] tilelang AscendC CI 自動チューニング供給用例の同期(#3784)— https://github.com/tile-ai/tilelang/pull/3784
- [4] tilelang ThreadSync が warp インデックス組合せ述語を証明(#3787)— https://github.com/tile-ai/tilelang/pull/3787
- [5] tilelang T.copy の転送先領域範囲外エラー(#3786)— https://github.com/tile-ai/tilelang/pull/3786
- [6] tilelang 全コピー fragment のスレッド帰属検証(#3788)— https://github.com/tile-ai/tilelang/pull/3788
- [7] tilelang SM120 FP8 block-scaled MMA(#3790)— https://github.com/tile-ai/tilelang/pull/3790
- [8] tilelang Profiler が出力欠落を拒否(#3797)— https://github.com/tile-ai/tilelang/pull/3797
- [9] tilelang Profiler の出力数と形状の検証(#3798)— https://github.com/tile-ai/tilelang/pull/3798
- [10] tilelang Profiler のリダイレクトストリーム出力抑制(#3799)— https://github.com/tile-ai/tilelang/pull/3799
- [11] tilelang 自動チューニングターゲット辞書のキャッシュキー(#3716)— https://github.com/tile-ai/tilelang/pull/3716
- [12] tilelang Profiler の複合シンボリック形状(#3717)— https://github.com/tile-ai/tilelang/pull/3717
- [13] tilelang NVCC の Tensor Core 検出(#3718)— https://github.com/tile-ai/tilelang/pull/3718
- [14] tilelang PTO 整数アトミックのコンパイル失敗(#3789)— https://github.com/tile-ai/tilelang/issues/3789
- [15] tilelang warp サイズ 64 の ThreadSync 証明(#3791)— https://github.com/tile-ai/tilelang/issues/3791
- [16] tilelang CuTeDSL NVFP4 量子化器(#3781)— https://github.com/tile-ai/tilelang/issues/3781
- [17] tilelang PTO アトミックの符号付きキャリア適応(#3792)— https://github.com/tile-ai/tilelang/pull/3792
- [18] tilelang PTO スカラー cast の VMI 合法化(#3793)— https://github.com/tile-ai/tilelang/pull/3793
- [19] tilelang PTO 整数アトミックの符号なしキャリア適応(#3794)— https://github.com/tile-ai/tilelang/pull/3794
- [20] tilelang PTO SIMD パディング lowering の復元(#3795)— https://github.com/tile-ai/tilelang/pull/3795
- [21] tilelang PTO 明示的 SIMD 並列ループの保持(#3796)— https://github.com/tile-ai/tilelang/pull/3796
- [22] tilelang 昇騰 Cube/Vector コア比設定(#3783)— https://github.com/tile-ai/tilelang/pull/3783
- [23] tilelang 昇騰 AutoSchedule 周期下限(#3776)— https://github.com/tile-ai/tilelang/pull/3776
- *
*[24]** tilelang-ascend 異常ダンプ CANN 9.2.0 互換(#1956)— https://github.com/tile-ai/tilelang-ascend/pull/1956
- [25] tilelang-hygon スケーリング変換と LDS 読み取りプランニングブランチ — https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-scaled-cast
- [26] tilelang-mlir-ascend BishengIR アダプタを contrib へ移行(#209)— https://github.com/tile-ai/tilelang-mlir-ascend/pull/209
- [27] TileOPs ショート varlen を KV グループごとにパッキング(#2819)— https://github.com/tile-ai/TileOPs/pull/2819
- [28] TileOPs ArgMax/ArgMin スレッドごと8ベクトル(#2828)— https://github.com/tile-ai/TileOPs/pull/2828
- [29] TileOPs オペランドを一度だけ読み取るキャッシュ戦略(#2820)— https://github.com/tile-ai/TileOPs/pull/2820
- [30] TileOPs L2 フラッシュを読み取り専用に変更(#2822)— https://github.com/tile-ai/TileOPs/pull/2822
- [31] TileOPs W4A16 単一リフィットコストモデル(#2825)— https://github.com/tile-ai/TileOPs/pull/2825
- [32] TileOPs 稠密 GEMM キャリブレーション契約ドキュメント(#2826)— https://github.com/tile-ai/TileOPs/pull/2826
- [33] TileOPs off-CUDA スイートを夜間ジョブへ移行(#2827)— https://github.com/tile-ai/TileOPs/pull/2827
- [34] TileOPs 単項要素ごと単一ウェーブ起動(#2829)— https://github.com/tile-ai/TileOPs/pull/2829
- [35] TileOPs FP8 グループ GEMM 連続レイアウト(#2824)— https://github.com/tile-ai/TileOPs/pull/2824
- [36] TileOPs GEMM オプション bias(#2823)— https://github.com/tile-ai/TileOPs/pull/2823
- [37] TileFoundry FSDP と HSDP を HIR 評価器へ(#284)— https://github.com/tile-ai/TileFoundry/pull/284
- [38] TileFoundry wgmma 同期プラン(#286)— https://github.com/tile-ai/TileFoundry/pull/286
- [39] TileFoundry 命名統一 issue(#287)— https://github.com/tile-ai/TileFoundry/issues/287
- [40] tilelang-cuda-skills(サードパーティスキルセット)— https://github.com/sablin39/tilelang-cuda-skills
- [41] tilelang-hexagon(サードパーティ Hexagon 移植)— https://github.com/xwhzz/tilelang-hexagon
- [42] Tilelang-AttnRes(サードパーティ注意力残差実装)— https://github.com/Index4NekoMeteor/Tilelang-AttnRes
- [43] tilelang リリースページ(最新 v0.1.15、09-30)— https://github.com/tile-ai/tilelang/releases
- [44] TileKernels(採用元、10-14 以降プッシュなし)— https://github.com/deepseek-ai/TileKernels
- [45] FlashQLA(採用元、10-30 以降プッシュなし)— https://github.com/QwenLM/FlashQLA
検証説明
本期の合入・審査中・issue・コミットの件数はすべて GitHub 検索と原子源の実測出力に基づくものであり、推計は行っていない。昇騰リポジトリと海光リポジトリの審査中 PR リストは、匿名インターフェースの割り当て上限に達したため取得できず、関連する結論はメインブランチのコミットとブランチレベルの原子源のみに基づいており、情報源核查表に注記済みである。tilelang-hygon の4件のコミットは PR 未作成の機能ブランチに位置し、ブランチレベルの動向であり、メインラインに合入済みであることを示すものではない。Google News、Hacker News、arXiv の3つのニュースチャネルは本調査期間中いずれもヒットがなく、そのためレポート本体はすべてコード側からのものである。