調査期間：2026-11-03 07:00 ~ 2026-11-04 07:00（北京時間、24時間；前回の締め切りは11-03 07:00）。

情報源：GitHub tile-ai 組織のリポジトリ別調査（29リポジトリ、期間内プッシュ9件）+ マージ済み・審査中真值検索（マージ済み8件、審査中更新48件）+ 新規issue全量確認（7件）+ メインブランチのコミット別確認（tilelang 5、TileOPs 1、TileFoundry 1、VS Code拡張機能 1、ドキュメントサイト 3）+ リリースページとtag確認（5リポジトリ、期間内新バージョンなし）+ 採用側2リポジトリ + Google News 中英文検索（プロキシ経由）+ Hacker News + arXiv。

今号のインデックス

  • 本日の重点：ThreadSync がストレージドメインをまたいで部分バリアを再利用し、#3643 の正確性修正時に導入された性能後退を補填（#3761）；3384 行の遺留 TVM パーサーブランチを整理（#3736）；自動チューニングの入力供給を現在のデバイスごとに変更（#3546）；TileFoundry 同期計画が正式に CUDA lowering へ着地（#285）；24 時間で 8 件をマージ、審査中更新 48 件、新規 issue 7 件
    1. コアプロジェクト進展
      • 1.1 メインライン：ThreadSync がストレージドメインをまたいで部分バリアの再利用を許可、B300 上のバリア占有が 12 から 16 ブロック/SM に回復（#3761）（11-04）
      • 1.2 メインライン：遺留の TVMScript AST とパーサーブランチを削除、一度に 3384 行を削除（#3736）（11-03）
      • 1.3 メインライン：自動チューニングの入力供給を現在のデバイスに変更、昇騰と MPS 上のチューニングが利用可能に回復（#3546）（11-03）
      • 1.4 審査中：条件内のスカラー cast がデータベクトル幅を圧縮しなくなり、昨日の #3734 を承継（#3759）（11-03）
      • 1.5 審査中：2 箇所のユーザー可視エラーを明確なエラー報告に変更（#3765 / #3767）（11-03）
      • 1.6 審査中：padding 付き共有 tile が全体コピーできない場合にフォールバック、2-CTA の TMA バイトを対等 leader で記帳（#3738 / #3766）（11-03）
      • 1.7 審査中：B300 の sm_103a を 2 ラインで適配、tcgen05 の int8 選択と CuTeDSL の TMA アーキテクチャゲート（#3770 / #3772）（11-03）
    1. マルチバックエンド適配
      • 2.1 昇騰：3D GM から L1 へのバッチ ND2NZ コピーをマージ、1 回の記述子でバッチ全体をカバー（#3689）（11-03）
      • 2.2 昇騰審査中：TileLang-Ascend リポジトリ全体をメインラインのマルチバックエンドアーキテクチャへ移行、2887 ファイル（#1975）（11-03）
      • 2.3 昇騰審査中：copy_pipe_to_l1 の同期注釈を実行エンジンごとに分割、A2/A3 上で 10/10 再現する誤結果を修正（#1981）（11-03）
      • 2.4 昇騰審査中：初回 Parallel 軸ブロードキャストをスカラー読み取りで処理、Chunk O フォワード移行、MHA サンプル（#1979 / #1866 / #1879）（11-03）
      • 2.5 mlir-ascend：Codex ワークフローとタスクハンドオーバーを harness ブランチへマージ（#208）（11-03）
      • 2.6 海光：Distributed 審査中 PR に更新あり、メインブランチに新規コミットなし（#19）（11-03）
      • 2.7 沐曦、摩尔線程は調査期間内にプッシュなし（11-03）
    1. エコシステムと採用側
      • 3.1 TileFoundry：同期計画を CUDA lowering に落とし込み逆検証、メインラインへマージ（#285）（11-04）
      • 3.2 TileFoundry 審査中：wgmma 同期計画、FSDP と HSDP 訓練セマンティクスを HIR 評価器へ（#286 / #284）（11-03）
      • 3.3 TileOPs：分散統合を最初の部分量から計算、定数入力の分散が NaN を出さなくなる（#2821）（11-03）
      • 3.4 TileOPs 審査中：GEMM にオプション bias を追加、FP8 グループ GEMM が連続レイアウトをサポート（#2823 / #2824）（11-03）
      • 3.5 TileOPs 審査中：W4A16 と稠密 GEMM のセレクタを単一代価モデルで再フィット（#2825 / #2826）（11-03）
      • 3.6 TileOPs 審査中：短 varlen を KV グループごとにパック、L2 flush を読み取り専用に変更（#2819 / #2822）（11-03）
      • 3.7 採用側：FlashQLA と TileKernels にプッシュなし、ただし #3761 は TileKernels の fused カーネルを修正（11-04）
    1. コミュニティ、チュートリアル、イベント
      • 4.1 単日 7 件の新規 issue：1 件のプロダクションレベル CUDA graph 問題、4 件は同一の外部報告者から（11-03）
      • 4.2 昇騰 950 の vcvt 飽和セマンティクスと mlir-ascend 日次テスト（#3762 / #1980）（11-03）
      • 4.3 ドキュメントサイト 3 回更新、VS Code 拡張を vsix に改名（11-03）
      • 4.4 ニュース側：3 チャネルで調査期間内のヒットなし、外媒の昇騰訓練報道は口径不一致（11-03）
    1. トレンド観察
  • 付、素材と検証説明

本日の重点

この日の主线は「性能を取り戻す」修正である。前の調査期間にマージされた #3643 は、実際の正確性の問題を解決した。個別に呼び出される ThreadSync("shared") と ThreadSync("shared.dyn") が、互いに交わらないスレッドグループに同じ名前付きバリア ID を割り当てる可能性があった。修正方法は、後の呼び出しが前の呼び出しで既に使われた ID を保持するようにするものだが、この保持ルールは同時にもう一つの事柄を妨げていた——同じスレッドグループが二つのスコープ間で本来同じ ID を再利用できるということである。代償は直接 TileKernels の mhc_pre_big_fuse カーネルに降りかかった。このカーネルはブロックごとに 96 スレッドで、スレッド 32 から 95 が二段式 cp.async パイプラインを実行し、非同期待機後の同期が第一パスで ID 3 を取得し、同じ 64 スレッドのバッファ依存同期が第二パスで ID 4 を取得する。余分なその ID により、コンパイル産物が宣言するバリアリソースが 4 から 5 に変わる。B300 上で隠れ次元 1280 を実行すると、バリア占有上限が 16 ブロック/SM から 12 ブロック/SM に低下し、レジスタ、共有メモリ、同期命令数はいずれも変化しなかった。#3761 の対処法は、ThreadSync が順序付きのスコープリストを受け入れるようにし、CUDA パイプラインを ThreadSync(["shared", "shared.dyn"]) に変更することである。同期は依然として元のスコープ順序で計画・挿入されるが、一部のバリアの ID は二つのスコープがともに処理された後に初めて統一的に割り当てられる。スコープを跨いで同じ ID を共有する前に、全く新しいアナライザで参加述語が候補スレッド範囲全体で成立することを証明する。分岐と Bind コンテキストは ConstrVisitor が提供し、スレッド関連の述語はデフォルトで成立するのではなく証明義務として扱い、証明済みの稠密グループは包围盒体積を直接用いて部分スレッド数の列挙をスキップし、未証明のグループはスコープをキーに残し参加者カウントも一緒にキーに入れ、明示的に番号付けされたバリア（より早い ThreadSync が割り当てた ID を含む）は引き続き保持される。変更は 401 行、4 ファイルに及ぶ。

著者は説明の中で非常にクリーンな帰属実験を行っている。コンパイル産物で宣言されたバリアリソース数を単独で変更し、いかなる実行可能命令も動かさなければ、同じ遅延を再現できる。この種の実験は記録に値する——「性能後退」を感覚から特定可能な単一変数に変え、修正の受け入れ基準も明確にした。すなわち、大まかにスコアが速くなることではなく、リソース宣言が 4 に戻り、占有が 16 ブロック/SM に戻ることである。

第二条は供給側の減算と口径の統一である。#3736 は tilelang/language/ast/、tilelang/language/parser/ と tilelang/language/tir/entry.py をまとめて削除した。これらは初期のローカル TVMScript パーサフォークの残したモジュールで、約 3380 行の Python であり、リポジトリ内に呼び出し元はもうない。言語層は現在 tilelang.language.tir.common を通じて上流の tvm.tirx.script.parser インターフェースを使用し、AST とステートメントのオーバーライド登録は tilelang.language.overrides にあり、T.prim_func と T.macro は tilelang.language.eager からエクスポートされる。実際の変更は 1 行追加、3384 行削除、10 ファイルで、同時に docs/conf.py の *language/ast* に対する無視ルールを削除した。公開言語層のエクスポートは方言テストとフロントエンドテストが担保する。#3546 は自動チューニングのデバイス口径を修正した。set_autotune_inputs がインストールする供給クロージャは、呼び出し点から決して渡されないパラメータ上でベースデバイスを解決していたが、実際に渡されるのはカーネルパラメータリストであった。そのため CUDA または CPU のフォールバック分岐が、CUDA のないマシン上でキャプチャした各テンソルを CPU に移動し、チューニングのたびにランチャーのデバイス不整合で失敗していた。現在クロージャは get_current_device() を使用し、デフォルトのテンソル供給と同じ戦略とした。CUDA の動作は変わらず、昇騰と MPS は現在のデバイス上で供給するように変更し、その他のバックエンドは依然として CPU にフォールバックする。受け入れはデュアルカード Ascend950 上で行った。三つの昇騰自動チューニング回帰は修正前全滅、修正後全通過し、メインリポジトリの testing/python/autotune/ は 12 項目通過、28 項目スキップ、さらに 3 項目がこのマシンに CUDA がないため失敗し、著者は明確に CUDA CI での検証に委ねている。

第三条は昇騰側である。GM から L1 へのコピーを「1 ディスクリプタ 1 行列」からバッチ全体に拡張した。#3689 は [B, M, N] 領域を一つの T.copy として書けるようにし、lowering はバッチ ND2NZ ディスクリプタを生成し、ソース側のバッチと行の stride を保持し、宛先側のバッチ間隔は物理レイアウトから導出し、転置バッチは DN2NZ で処理する。GM 境界を超えるソースバッチは直接スキップし、宛先ストレージをそのまま保持し、M/N の末尾パディングは合法なバッチにのみ作用する。既存の二次元および単一 matrices の動作は変わらず、実装は AscendC バックエンドのみを対象とし、L1 宛先が規範的な行列レイアウトであることを要求し、PTO は今回の範囲外である。検証は Ascend950DT と CANN 9.2 上で行った。50 の AscendC サンプルテストと 1733 の昇騰単体テストが通過し、うち 28 のデバイスケースがバッチと M/N 末尾ブロック、バッチ全体の越境、非連続ソース stride、宛先ストレージがより大きい場合、矩形転置をカバーする。CUDA と PTO のケースは除外され、CI イメージが使用する CANN 9.3 はローカルでテストされておらず、この点は著者が説明に記している。変更は 579 行、10 ファイルに及ぶ。

4つ目は TileFoundry の隠れた実装です。前回のレポートではまだ「審査中」でしたが、今回の調査期間でマージされました：#285 により、sync 分析が計画したコミット、待機、フェンスが実際に CUDA lowering によって発行され、発行された内容と計画が一致するかどうかが項目ごとに逆検証され、不一致の場合は LoweringError がスローされます。コーパス上での結果は、デフォルトフローで 14 個のプログラムに同期が追加され、31 個が拒否されました。この 14 個の TIR とメインラインの違いは追加された文のみで、別の 30 個はメインラインと完全に同一です。単一の sm_89 上でビルド可能かつ同期を含む 7 個のプログラムを実行したところ、6 個がリファレンスと完全に一致し、synccheck と memcheck はすべてクリーンでしたが、同じ 7 個をメインラインの方法で lower すると 4 個で結果が誤りとなりました。作者が動機で示した直接的な症状は、sm_89 の E4M3 GEMM の最大偏差 659 です。このコミットは 788 行追加、214 行削除、25 ファイルで、「分析が証明できる」から「発行は証明と一致しなければならない」へと進めました。

1. コアプロジェクトの進展

1.1 メインライン：ThreadSync がストレージドメインを跨いだ部分バリアの再利用を許可、B300 上のバリア占有が 12 から 16 ブロック/SM に回復（#3761）

11-04 00:50 にマージ、401 行追加、72 行削除、4 ファイル。背景は前回の調査期間の #3643 です：ThreadSync("shared") と ThreadSync("shared.dyn") を別々に呼び出すと、互いに交差しないスレッドグループに同じ名前付きバリア ID が割り当てられる可能性があり、これは実際の正確性の問題です。修正方法は、より早い呼び出しで既に使用された ID を保持することでした。副作用として、同じスレッドグループがスコープを跨いでも再利用できなくなり、TileKernels の mhc_pre_big_fuse カーネルではブロックあたり 96 スレッドの余分な割り当てが発生しました：スレッド 32 から 95 の 2 段階 cp.async パイプラインが 1 回目で ID 3、2 回目で ID 4 を取得し、コンパイル生成物が宣言するバリアリソースが 4 から 5 に増え、B300 の隠れ次元 1280 のバリア占有上限が 16 ブロック/SM から 12 ブロック/SM に低下しましたが、レジスタ、共有メモリ、同期命令数は全く変化しませんでした。本コミットでは ThreadSync が順序付きのスコープリストを受け入れるようにし、CUDA パイプラインを ThreadSync(["shared", "shared.dyn"]) に変更しました：同期は元の順序で計画通りに挿入され、部分的なバリア ID は両方のスコープの処理が完了した後に統一的に割り当てられます。スコープを跨いで 1 つの ID を共有する前に、新しいアナライザで参加述語が候補スレッド範囲内で恒真的に成立することを証明します。分岐と Bind コンテキストは ConstrVisitor によって提供され、スレッド関連の述語は仮定ではなく証明義務として扱われ、証明済みの密なグループはバウンディングボックス体積で計算され、部分的なスレッド数の列挙はスキップされます。未証明のグループはスコープをキーに残し、参加者数をキーに入れます。明示的に番号付けされたバリア（より早い呼び出しで割り当てられた ID を含む）は引き続き保持され、交差しないスレッドグループは依然として分離されます——#3643 の修正はロールバックされていません。文字列形式の旧来の呼び出しは互換性が維持され、CuTeDSL も同じ CUDA pass パイプラインを通ります。

出典：https://github.com/tile-ai/tilelang/pull/3761

1.2 メインライン：遺留の TVMScript AST とパーサーブランチを削除、一度に 3384 行削除（#3736）

11-03 13:05 にマージ、1 行追加、3384 行削除、10 ファイル。削除されたのは tilelang/language/ast/、tilelang/language/parser/、tilelang/language/tir/entry.py です——初期のローカル TVMScript パーサーフォークの残骸で、リポジトリ内に呼び出し元はありません：言語層は tilelang.language.tir.common を通じて上流の tvm.tirx.script.parser インターフェースを使用し、TileLang 自身の AST と文のオーバーライドは tilelang.language.overrides に登録され、T.prim_func と T.macro は tilelang.language.eager からエクスポートされ、tilelang/language/tir/ の下では common、exports、ir、op の 4 つのモジュールのみが消費されています。この孤立したパーサーフォークを保持するリスクは具体的に書かれています：誰かが直接 tilelang.language.parser.operation を import すると、tirx.PrimExpr と tirx.IterVar 上に互いに衝突するオペレータを登録してしまいます。本コミットでは同時に tilelang/language/tir/__init__.py 内の prim_func への不要な再エクスポート、および docs/conf.py 内の *language/ast* に対する無視ルールを削除しました（*parser* は tilelang/language/overrides/parser.py がまだ使用しているため保持）。受け入れ基準は公開エクスポートが不変であることで、test_tilelang_language_dialect.py と言語フロントエンドのテストでカバーされています。

出典：https://github.com/tile-ai/tilelang/pull/3736

1.3 メインライン：自動チューニングの入力供給が現在のデバイスを使用するように変更、昇騰と MPS 上のチューニングが再び利用可能に（#3546）

11-03 13:19 にマージ、153 行追加、12 行削除、3 ファイル、#3545 を修正。set_autotune_inputs がインストールする凍結供給クロージャは、「呼び出し点で決して渡されない」あるパラメータからベンチマークデバイスを解析していた——供給プログラムの契約によれば、実際に渡ってくるのはカーネルパラメータリストである——そのため、CUDA のないマシンでは、CUDA または CPU のフォールバック分岐が各キャプチャテンソルを CPU へ移動し、チューニングの各ラウンドがランチャーのデバイス不一致で失敗していた。現在クロージャは get_current_device() を使用し、デフォルトテンソル供給と同じ戦略をとる。CUDA 上では動作は変わらず、マルチカードのベンチマーク worker は依然としてカードごとにコピーする。昇騰と MPS 上ではキャプチャテンソルがそれぞれの現在のデバイスで供給され、チューニングが可能になる。その他のバックエンドは依然として CPU にフォールバックする。追加された昇騰のユースケースは testing/ascend/autotune/ に置かれ、デコレータ経路、AutoTuner のベンチマーク供給デバイス、およびクロスデバイスでキャプチャされた入力の 3 つのケースをカバーし、デュアルカード Ascend950 上で修正前は全敗、修正後は 3 項目すべて通過した。メインリポジトリの testing/python/autotune/ は 12 項目通過、28 項目スキップ、さらに 3 項目（eager モード、スカラー入力、入力付き）が同機に CUDA がないため失敗し、作者は CI の CUDA runner での検証に残すと注記している。

出典：https://github.com/tile-ai/tilelang/pull/3546

1.4 審査中：条件内のスカラー cast がデータベクトル幅を圧縮しなくなる、昨日の #3734 を受ける（#3759）

11-03 14:09 更新、148 行追加、29 行削除、2 ファイル。これは昨日マージされた #3734 の仕上げである：#3734 は DecoupleTypeCast の後でデータ cast の幅制限を正しく厳格化したが、「スカラー条件内の cast」もデータベクトル幅を制限するようになった——b5cb02633 上では、ガードされた BF16 コピーが条件内の int64 cast のために 16 レーンから 4 レーンに圧縮される可能性があった。たとえその条件が 16 レーン上で完全に一致していてもである。修正方法は、プランナーが既に持つ条件一致性解析を cast 幅の豁免に再利用し、if 文、tirx.if_then_else、およびスカラー述語の Select 条件をそれらの節から分離して個別に訪問することである。条件内の cast はもはやデータ型幅の制限を課さず、節内の cast は依然として課し、レーンごとのベクトル述語を持つターゲット上では条件 cast の制約を保持する。#3579 のメモリガード動作と #3734 が修正したガード付き FP32 から BF16 への制限はともに変わらない。回帰は 8 レーンと 16 レーンの一致ガード、3 種類の条件形式、節内の FP32 cast、ベクトル述語と丸めビットグループをカバーし、9 つの int64 ガードユースケースは修正前に失敗し、最終的な合法化と cast 丸めの 2 組のスイートは計 74 項目通過した。

出典：https://github.com/tile-ai/tilelang/pull/3759

1.5 審査中：2 箇所のユーザー可視エラーを明確なエラー報告に変更（#3765 / #3767）

いずれも 11-03 17:10 前後に更新され、同じ報告者によるコミットである。#3765 は #3653 に対応する：単一スレッドインデックス境界ガードによる部分バリアでは、その境界が warp サイズの整数倍のとき、warp 全体のスレッドが条件判断に一致する。しかし汎用算術証明は、外側分岐がビット演算 & を使用する際の thread_binding < 448 を見落とし、ThreadSync が直接中止した。本コミットはスレッド次元が整数 warp 数のときこの状況を明示的に受け入れ（同時に境界がその次元と等しいか超える場合もカバー）、その他の条件は依然として既存の証明経路を通る。 #3767 は #3005 に対応する：fragment のクロススレッド分布で、同じ T.Parallel 本体內で構造の異なる 2 つのインデックスを使って同じ fragment を読み書きする場合（例えば buf[i] と buf[(i + 1) % N] のようなクロスレーンテンプレート）、クロススレッド交換なしには lower できず、以前は ParallelOpNode::RecordBufferAccess が StructuralEqual の ICHECK に遭遇し、コンパイラ内部アサーションとして現れていた。現在は ValueError に変更され、2 つのインデックス式をエラーメッセージに書き込み、fragment をこのようにアクセスできないことを説明する。両方とも「内部アサーションをユーザーが読める診断に置き換える」ものであり、利用者への体感的な影響は変更規模より大きい。

出典：https://github.com/tile-ai/tilelang/pull/3765、https://github.com/tile-ai/tilelang/pull/3767

1.6 審査中：padding 付き共有 tile が全体コピーできない場合にフォールバック、2-CTA の TMA バイトは対等 leader で計上（#3738 / #3766）

2件とも 11-03 13:47 から 13:31 の間に更新された。#3738 は #3704 を修正する：padding 後の共有 tile より小さいグローバル領域から T.copy を行う場合（例えば K=100 を BK=112 にコピーする）、ClassifyWarpSpecializedCopy と SelectCopyInstForLowering によって TMA の一括コピーと判定され、その後 lowering でハード失敗し、s_extent == g_extent (112 vs. 100) を報告する。原因は、選択ロジックが最後の次元のアラインメント、dtype のペアリング、グローバル stride のみをチェックしており、padding 後の範囲が一括コピーを構成できるかをチェックしていなかったことである。本コミットはこの判定基準を追加し、満たさない場合は通常のコピーパスにフォールバックする。#3766 は cluster シナリオにおける TMA バイトカウントを修正する：cta_group::2 の TMA load はバリアアドレスのピアビットをクリアし、一対の CTA のバイトはすべて偶数 CTA のバリアに到達するが、cluster バリア付きの tma_copy は元々 rank 0 に対してのみ cluster_size × tile_bytes を期待していた——(2,1,1) と (1,2,1) では偶然一致するが、(4,1,1) と (2,2,1) では rank 0 は 2 份しか受け取らず 4 份を期待し、さらにバイトが rank 2 のバリアに到達する。2件とも新機能ではなく、「境界条件が破られると内部エラーを報告する」タイプの欠陥である。

情報源：https://github.com/tile-ai/tilelang/pull/3738、https://github.com/tile-ai/tilelang/pull/3766

1.7 審査中：B300 の sm_103a 適配の2つのライン、tcgen05 の int8 選択と CuTeDSL の TMA アーキテクチャゲート（#3770 / #3772）

B300 が報告するターゲットアーキテクチャは sm_103a であり、2つの審査中 PR がそれぞれそれがもたらす2つの障害点を処理する。#3770 は #3769 に対応する：tcgen05.mma.kind::i8 はアーキテクチャ専用の sm_100a と sm_110a にのみ存在し、sm_103a（B300）上にはなく、ファミリー可搬の sm_10xf にも属さないが、AllowTcgen5Mma は TargetIsSm100（arch 100 から 110）に対して一律に int8/uint8 オペランドを受け入れ、カーネルは順調に lowering を通過した後 ptxas で Feature '.kind::i8' not supported on .target 'sm_103a' を報告する。本コミットは TileLang がこれらのターゲットで kind::i8 を選択しないようにする。#3772 は CuTeDSL の TMA アーキテクチャゲートを処理する：B300 の sm_103a は DSL に入る前に拒否され、エラーは sm_90/90a/100a/110/120/120a/121/121a を列挙するが 103a を含まない。メインライン fa69f7fa に nvidia-cutlass-dsl 4.7.0 と CUDA 13.3 の組み合わせを加えた場合、B300 上で6つの CuTeDSL テストが失敗し、tma_load の複数の JIT ケースを含む。同じバッチには #3768 で SM100 flash attention 逆方向サンプルを修正するものもある：T.tcgen05_gemm は現在 mbar を必須のキーワード専用引数に固定し、暗黙的な待機を行わなくなったが、サンプルは依然として同期 GEMM の書き方で呼び出しており、2つのスクリプトはコンパイル前に missing 1 required keyword-only argument を報告する。作者は同時に examples/flash_attention_sm100/ ディレクトリに test_*.py がないことを指摘しており、これが問題がずっと発見されなかった原因である。

情報源：https://github.com/tile-ai/tilelang/pull/3770、https://github.com/tile-ai/tilelang/pull/3772、https://github.com/tile-ai/tilelang/pull/3768

2. マルチバックエンド適配（昇騰 / 海光 / 摩尔線程 / 沐曦）

2.1 昇騰：3D GM から L1 へのバッチ ND2NZ コピーがマージ、1つの記述子でバッチ全体をカバー（#3689）

11-03 17:07 時点で、579 行追加、39 行削除、10 ファイル。これまで昇騰の GM から L1 へのコピーは 1 つの DMA 記述子で 1 つの行列しか処理できず、複数バッチをコピーするには for b in T.serial(B) でバッチごとに呼び出す必要があった。本コミットでは [B, M, N] 領域を 1 回の T.copy で記述できるようにし、lowering がバッチ ND2NZ 記述子を生成する。ソース側のバッチ stride と行 stride を保持し、宛先側のバッチ間隔は物理レイアウトから導出し、転置バッチは DN2NZ で処理する。既存の行列軸境界処理を再利用する。すなわち、GM 境界を超えるソースバッチはスキップされ、宛先ストレージはそのまま維持され、M/N の末尾ゼロ埋めは有効なバッチにのみ適用される。既存の 2 次元および単一行列の動作は変わらない。実装は AscendC バックエンドのみを対象とし、L1 宛先側が正規の行列レイアウトであることを要求する。PTO のサポートは今回の変更には含まれない。検証環境は Ascend950DT と CANN 9.2 で、50 個の AscendC サンプルテストと 1733 個の昇騰ユニットテストが通過した。うち 28 個の BF16/FP32 デバイスケースがバッチおよび M/N の末尾ブロック、バッチ全体の越境、非連続ソース stride、宛先ストレージがソースより大きいケース、矩形転置コピーをカバーする。作者は CUDA と PTO のケースを除外したこと、CI イメージが使用する CANN 9.3 はローカルで検証していないことを注記している。これは前号でマージされた昇騰 950 FlashMLA カーネルと同じ方向性の補完であり、アテンションと GEMM のバッチ転送はいずれも 1 回の記述子でバッチ全体をカバーする必要があり、バッチごとの呼び出しでは本番形状で相当な発行オーバーヘッドが生じる。

出典：https://github.com/tile-ai/tilelang-ascend/pull/3689

2.2 昇騰レビュー中：TileLang-Ascend リポジトリ全体をメインラインのマルチバックエンドアーキテクチャへ移行、2887 ファイル（#1975）

11-03 16:34 更新、変更規模は 607317 行追加、97595 行削除、2887 ファイルに達し、今号で最大の変更量であり、組織全体のレビュー中 PR の中で最も追跡する価値のある一本でもある。目標は提案 #1911 に従い、TileLang-Ascend をメインラインのマルチバックエンドアーキテクチャへ移行することである。既存の昇騰コンポーネントをメインラインの対応ディレクトリの ascend/a2a3 下へ移す（tilelang/ は tilelang/ascend/a2a3、src/ は src/ascend/a2a3）。元の Ascend C と PTO A2A3 バックエンドはメインラインで独立して登録され、メインラインの Ascend、PTO A5 と直交し、pass と codegen は互いに干渉せず、少量のインフラのみを共有する。両者は target="ascend" と "pto" を共用するが、A2A3 バックエンドは arch="dav-2201" 識別子（NPU 名から自動認識）を付帯してこれで振り分ける。A2A3 フロントエンドプリミティブは import tilelang.ascend.a2a3.language as T に変更し、A5 の import tilelang.ascend.language as T と併存し、オペレータの記述方法は元リポジトリと変わらない。pass と codegen は TileLang-Ascend から移行したバージョンを直接使用して元のコンパイルフローを維持し、実行バックエンドはメインラインの cython を再利用しつつ workspace インデックスなど昇騰固有の適配を補う。examples と testing はそれぞれ examples/ascend_a2a3 と testing/ascend_a2a3 に配置し、CI は当面元リポジトリの ci_cd.yml を使用し、新パスに適配後も動作することを確認している。作者は説明の中で TODO を率直に記している。3 つのディレクトリ下のファイルと元ファイルとの差異はまだ収束させる必要があり、静的チェックとフォーマットの CI は復旧待ち、ドキュメントと Agent 関連は移行待ちであり、「ファイルが多く、Agent の一括操作ではどうしても漏れが生じる」と注記し、レビューの協力を求めている。移行中の不確実性をそのまま PR 説明に開示するこの做法は、「完了」という結論を待つよりもレビューに有利である。

出典：https://github.com/tile-ai/tilelang-ascend/pull/1975

2.3 昇騰レビュー中：copy_pipe_to_l1 の同期注記を実行エンジンごとに分割、A2/A3 で 10/10 再現する誤結果を修正（#1981）

11-03 12:36 更新、#1780 に対応。根本原因は明確に書かれている。A2/A3 の PTO クロスコア V から C への通信は GM を中継し、三段の転送はそれぞれ別のエンジンに属する——AIV の MTE3 が UB から GM への送信、AIC の MTE2 が GM から L1 への受信（TPOP 内部の TLOAD）、AIC の MTE1 が L1 から L0A を行う。一方 copy_pipe_to_l1 は operation_config.h 内で PIPE_MTE3 と标注されており、自動同期パスはそれゆえ copy_pipe_to_l1 と copy_l1_to_l0a の間に set_flag(PIPE_MTE3, PIPE_MTE1, …) と対応する wait_flag を生成する——しかし待つのは AIC 自コアの MTE3 であり、AIC の MTE3 キューは空でイベントは即座にセットされ、MTE1 の L1 読み取りは MTE2 の書き込み完了前に放行される可能性があり、L0A は古いデータを取得し、MMA の結果が誤る。誤った标注は #1238 で導入され、同じ表の中で copy_pipe_to_ub と copy_l0c_to_pipe はどちらも実行側エンジンで标注されており、この一条だけが送信側視点を踏襲している。再現証拠は二つ：issue 作者が生成ソース内の PIPE_MTE3 を PIPE_MTE2 に変更したところ 10/10 すべて精密に恢復し、ローカルでは最新主干上で 10/10 回の実行エラー、単回最大 16384/16384 個の要素が不一致。修正方法は issue の提案に従い名前を分割し直接置換しない：copy_pipe_to_l1 の标注を PIPE_MTE2 に変更（A2/A3 の GM から L1 へのロード）、新たに copy_pipe_to_l1_bind を追加し PIPE_MTE1 と标注（A5 直結パス上で L1 の準備完了を待ちアドレスをバインド）、copy_pipe_to_ub と copy_pipe_to_ub_V の既存の先例を踏襲し、A2/A3 を直しつつ A5 を誤らせることを避ける。同源の問題は WAR 方向もカバーする：MTE1 がまだ読み取り中に、MTE2 の TLOAD が同じ L1 バッファを上書きする可能性がある。

出典：https://github.com/tile-ai/tilelang-ascend/pull/1981

2.4 昇騰審査中：初の Parallel 軸ブロードキャストがスカラー読み取りに、Chunk O 前向き移行、MHA 示例（#1979 / #1866 / #1879）

三件とも 11-03 午前に集中更新され、昇騰オペレータのマイルストーン的な一批に属する。#1979 は #1885 に対応：失敗族は一つのパターンに帰着できる——最初の T.Parallel 軸（または外側の serial 変数）でインデックスされた一次元 UB オペランドが二項式の左側に現れる。issue の収縮パターン dk_i = sum_j v[j] * A[j, i] はまさにこの形態であり、21 点二分マトリクスの中で 5 つの FAIL はすべてこの族に属し、帰約方向やブロードキャストの書き方とは無関係；根本原因は AscendLowerParallelToVector が二次元並列ネストを行ループに降下させる際、エントリ HandleSimpleCase が右側オペランドに対してのみスカラーアクセス検出を行い、左側がベクトルパスに進んだ後 UB アライメントトラップに衝突する；本笔は首軸のブロードキャストオペランドもスカラー読み取りに進める。#1866 は Chunk O 前向き（chunk ベースのゲート付き線形注意力 O 前向き、GDN）を昇騰 NPU に移行し、開発者モードと四項 pass 設定（自動 C/V マージ、自動 C/V 同期、自動同期、メモリプランニング）を使用、単一の @tilelang.jit カーネルと固定 20 コアの grid-stride ループで、オンチップ融合、ゼロ GM workspace、ゼロ宿主 ATen オペレータを実現。#1879 は MHA 前向き示例を追加：BSHD レイアウト、Expert モードの三分派 MMA 流水を、wrapper、共有層、三形態モジュールに分けて 5 つの Python ファイルに分割、wrapper は T.* コードを一切含まず、メタデータ検証と JIT 分派のみを行う；is_causal は右下角アライメントマスクを採用し、全マスク行は sumexp + 1e-30 保護を経て出力 0；cann-bench 20 用例の精度はすべて通過し、torch-npu ベースラインと比較して平均加速比 0.97 倍、総合得点 73.13（ローカル 910B2 実測口径）。0.97 倍はまだ純収益がないことを意味するが、示例として「昇騰上で完全な注意力オペレータをどう書くか」の骨格を示している。

出典：https://github.com/tile-ai/tilelang-ascend/pull/1979、https://github.com/tile-ai/tilelang-ascend/pull/1866、https://github.com/tile-ai/tilelang-ascend/pull/1879

2.5 mlir-ascend：Codex ワークフローとタスク引き継ぎを harness ブランチにマージ（#208）

11-03 17:50 にマージ済み。2715 行追加、39 行削除、66 ファイル。マージ先ブランチは main ではなく harness（メインラインの最新コミットは依然として 10-31 の ci: remove internal cache-service references のまま）。この点を引用する際は「メインラインにマージ」と書かないこと。内容はオペレータ harness に最初の Codex ワークフローを追加するもの：プロジェクト初期化時に反復質問を用いる、既存オペレータリポジトリを接続する、永続化コンテキストから開発・チューニング・分析を復元する；3 つのスキルは .workspace/ 配下の編集可能な Markdown 記録を共有し、タスク引き継ぎでは「現在の候補」と「検証済みの最良候補」を区別し、失敗・ロールバックした実験を保持し、サブエージェントには焦点を絞ったコンテキストを渡し、メインエージェントが共有状態を維持する；ローカルインストーラはスキルを登録し、プロジェクトエントリを AGENTS.md に追記する（既存内容を保持、.workspace/ を無視、Git プロジェクトを worktree ルートに解決）。README、依存関係、再利用説明、ロードマップはいずれも TileLang 先行・昇騰をスコープとする設定を記述している；ワークフローは GSD を参考にし、出典リビジョンと MIT 声明を保持しており、NPU 実行アダプタと npm 配布はまだ TODO に残っている。このコミットはメインリポジトリ・昇騰リポジトリの Agent 関連ディレクトリと同一事項の 3 つの側面である：「エージェントがオペレータを書く」プロセスをリポジトリ内の再利用可能な資産として定着させること。

出典：https://github.com/tile-ai/tilelang-mlir-ascend/pull/208

2.6 海光：審査中 PR に更新あり、メインブランチに新規コミットなし（#19）

tilelang-hygon は調査期間内にプッシュ記録がある（11-03 11:06 北京時間）が、メインブランチに新規コミットはなく、最新は依然として 11-02 の #32（gfx948 を gfx946 互換として扱い、まず既存命令をテスト可能にする）。プッシュは審査中の #19 からのもの（タイトルは Distributed のみ、説明は空）、11-03 11:06 更新。時刻とタイトルから分散関連の適配を進めていると見られるが、PR 説明に内容がないため、本レポートでは実装詳細について推測しない。同リポジトリの最近の通常ペースは 2〜3 日ごとに 1 件の HCU（海光計算ユニット）側修正：10-30 の #31 は MLS 境界と永続 WASP 同期を変更、#30 は通常コピーの LDS swizzle を補完、10-27 の #29 は scale バッファの bank 分割を実施。

出典：https://github.com/tile-ai/tilelang-hygon/pull/19

2.7 沐曦、摩尔線程は調査期間内にプッシュなし（11-03）

tilelang-musa（摩尔線程）の最終プッシュは 09-30 09:50、tilelang-metax（沐曦）は 10-26 07:08 で、両リポジトリとも本調査期間にコミットはなく、審査中 PR の更新もない。本期のバックエンド動向はしたがって昇騰と海光の両側に集中している：マージ済み 8 件のうち 1 件が昇騰（#3689）、1 件が mlir-ascend（#208、harness ブランチにマージ）、審査中の 48 件の更新のうち昇騰リポジトリが 15 件、メインリポジトリが 23 件を占める。

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

3. エコシステムと採用側（TileOPs / TileFoundry / TileKernels / FlashQLA）

3.1 TileFoundry：同期計画が CUDA lowering を実現し逆方向検証、メインラインにマージ（#285）

11-04 00:19 にマージされ、788 行を追加、214 行を削除、25 ファイル。sync 分析は共有メモリアクセスに必要なコミット、待機、フェンスを計画するが、finalize はこれまで一つも発行していなかったため、共有メモリに一時退避するスケジューラには競合が存在し、カード上では sm_89 の E4M3 GEMM が torch に対して最大 659 ずれる形で現れた。本コミットは分析ステップをそれらが先行する HIR ノードに紐付け、SyncMetadata としてタグでインデックス化する。ノード自身の文のタグは空、InsertSlice が加えるコピーは base と update、ループキャリーコピーは ("entry", スロット) と ("backedge", スロット) を用いる。コア lowering はあるノードの文を lowering する前に新しいフック before(owner, ctx, tag) を呼び、CudaLowering がそこでそのノードのステップを発行する。schedule finalize が sync 分析の実行を担う。契約は四つの仕様ドキュメントに書き込まれ、うち分析仕様 §1.2.5 が SyncMetadata、タグ、規則、および分析が依然として拒否する情形を記述する。リスク面も著者が明記している。コーパス内でメインライン上 finalize できるすべてのプログラムは依然 finalize 可能で、文が増えるだけである。例外は hir/staged_dynamic_inner で、その lane は「行程数が lane に等しい」ループ内で同じ共有タイルを読み書きし、それらを合流させるフェンスが存在しない。また共有メモリに触れる保留呼び出しは「未知バイトへの一回のアクセス」として処理され、そのような二つの呼び出しがスレッドグリッドのない領域にある場合、TIR が配置できないフェンスを必要とし、バイトが互いに交差しなくても分析はそれらを拒否する。CopyAsyncTensor、CopyReduceAsyncTensor と wgmma にはまだエミッタがない。

出典：https://github.com/tile-ai/TileFoundry/pull/285

3.2 TileFoundry 審査中：wgmma 同期計画、FSDP と HSDP 訓練セマンティクスの HIR 評価器への導入（#286 / #284）

#286 は 11-04 02:28 に更新され、1907 行を追加、310 行を削除、52 ファイル。これは #285 の直接の続きである。WGMMA は発行後も共有メモリとレジスタを読み続ける可能性があり、CUDA 同期計画は後続アクセスをこれらの作用の後に並べる必要がある。本コミットは各オペランド効果の完了セマンティクスを宣言し、WGMMA のフェンス、コミット、待機をそれらを必要とする HIR ノードに紐付け、lowering が命令の発行点でこれらのステップを消費し、後続スコープが集合待機を収められない場合はその発行スコープ内で当該操作を完了する。同時にレジスタ Reshard の別名を修正し、要素がコピーを保持する各スレッド上で同じレジスタに留まるようにし、このストレージ同一性でレジスタ依存を追跡する計画である。著者はこのリスクを明記している。行列命令エミッタは依然使用不能で、本コミットは WGMMA の実行正確性や性能を確立せず、TMA 完了セマンティクスは依然別途処理される。#284 は 11-03 14:34 に更新され、訓練セマンティクスを扱う。従来の評価器は reshard をグローバルテンソルとして扱い、異なる参加者の寄与を漏らしていた。本コミットは参加者ローカル値、GPU 座標、呼び出しとループ、gather/scatter および軸ごとの Partial(sum) 解析を評価し、Reshard の別名証明に所有権セマンティクスを持たせ、#260 と #170 の構造化ライフサイクルを拡張する。併せて modded-nanogpt に触発された二ブロック GQA 研究を追加し、明示的な順伝播・逆伝播、モーメンタム SGD、独立した PyTorch 参照を含み、FSDP と 2×2、4×4 の二種の HSDP 常駐戦略をカバーし、集合通信グループ、ペイロード、ストレージを報告する。受け入れ数値は、リファクタ後のコアスイート 775 項、バックエンドスイート 778 項（二カード sm90 上で 13 項スキップ）、十組の HIR 設定が CPU 上で通過、二種の戦略が四カード sm90 上で実 2×2 PyTorch HSDP と二更新ステップを実行し、ローカル勾配、パラメータ、モーメンタムがすべて整合する。

出典：https://github.com/tile-ai/TileFoundry/pull/286、https://github.com/tile-ai/TileFoundry/pull/284

3.3 TileOPs：分散の統合が最初の部分量から起算され、定数入力の分散がもはや NaN を出さない（#2821）

11-03 18:04 にマージされ、43 行を追加、4 行を削除、3 ファイル。欠陥は分散の統合が空のアキュムレータから始まることである。有限の大きな平均を空のアキュムレータに統合する際、二乗差がゼロ重みを掛ける前に既にオーバーフローし、定数入力の分散が NaN になっていた。影響を受けるのは、保留軸の前後にそれぞれ一段のリダクションがある場合の var、std、var_mean の呼び出しである。修正は統合を最初の真の部分量から初期化し、後続の部分量は既存の Chan 漸化式を踏襲する。行リダクション、ディスパッチ、出力レイアウト、公開 API はすべて不変。受け入れ：三件の新規回帰が元のカーネルで失敗し、修正後に通過、候補は 155 項の回帰テストと 24 項の追加数値チェックを通過、tests/ops/test_reduce.py は 155 から 158 に増加。三項の自動チューニングテストは除外された。理由は以前のベースライン試行が 60 秒のテスト上限を超えたためである。この種の「空アキュムレータ + 大きなオフセット」の数値の落とし穴はリダクション族では珍しくなく、記す価値があるのは、保留軸の前後で同時にリダクションする形状でのみ露見し、通常の整軸リダクションでは検出できない点である。

出典：https://github.com/tile-ai/TileOPs/pull/2821

3.4 TileOPs レビュー中：GEMM にオプションの bias を追加、FP8 グループ化 GEMM が連続レイアウトをサポート（#2823 / #2824）

#2823 は GEMMFwdOp.forward にオプションの bias[N] を追加した。bias は float32 の累積結果に加算され、出力は1回だけ丸められ、torch.nn.functional.linear の仕様と一致する。各稠密構造はそれぞれのエピローグで bias を加算する——warp 専用（TMA と行ストレージ）、coop2（176/192 幅と狭い tile を含む）、coop2_swap、ping-pong、coop2s、simple、swap_ab、GEMV はそれぞれ cp.async カーネルを伴い、split-K 経路はリダクション内で1回だけ加算する。#2824 は MoEGroupedGEMMFP8FwdOp がマスク付きブロック以外にも連続したエキスパート行（コンパクトパックまたはアラインメントパック、セグメント末オフセット physical_psum または行ごとのエキスパート id 付き）を受け入れるようにし、セマンティクスを deep_gemm.m_grouped_fp8_gemm_nt_contiguous に合わせた。SM90 カーネルはセグメント末ごとに各エキスパートの実行セグメントを走査し、コンパクトな行ごと id はプロデューサとコンシューマの分割前に二分探索でセグメント末に変換され、アラインメントされた行ごと id は行 tile ごとにブロック単位で進み、先頭行からその tile のエキスパートを読み、コンパクトエキスパートの最後のブロックが block_m 行に満たない場合は行ごとに格納し、決して次のエキスパートの行に書き込まない。SM89 は連続レイアウトを拒否する。マスク付き呼び出しは以前と同じプログラムを構築し、35 個のマスク構築 TIR はメインラインとバイト単位で一致する。ベンチマーク側では Qwen3-30B-A3B の EP8 シェイプを3行追加し、実際のルーティングから測定したグループサイズを用いた（このモデルは 128 エキスパート、top-8、48 層のうち偏りの中央付近にある第 44 層を選び、EP8 で最も多くの行が割り当てられた rank の 16 エキスパートを選ぶ）。prefill 行は 4096 トークンをルーティングし（gate/up と down）、decode 行は 128 シーケンスの最終トークンを取る。

出典：https://github.com/tile-ai/TileOPs/pull/2823、https://github.com/tile-ai/TileOPs/pull/2824

3.5 TileOPs レビュー中：W4A16 と稠密 GEMM のセレクタを単一のコストモデルで再フィッティング（#2825 / #2826）

両方とも #2133 のサブタスク 2 に属し、それぞれ 11-03 20:54 と 20:54 に更新され、一方は W4A16 を、もう一方は稠密を扱い、後者はドキュメントのみを変更する。#2825 は _select_config を _legal_configs 上で _config_cost の argmin を取るように変更し、tile 優先、最も広い K の同点ウィンドウ、joint_switch_margin の3つの例外ルールを削除した。_config_cost には stage_fill 項を新設した——より深いリングは各 K tile のレイテンシをより多く隠せるが、フィリング自体は遅くなる。この項がない場合、K=4096 でモデルは 256 幅4ステージを 512 幅2ステージより前に配置するが、後者は実測で最大 11% 高速である。mma_big、iter_ws、launch_ms はデータがそれらを固定するのに不十分なため削除された。定数はベンチマークの SM クロック下でビジー時に再フィッティングされ、フィッティング形状はリスト外の形状を取り、リスト自体はホールドアウト検証として残される。_Calibration のドキュメント文字列にはスキャン、フィッティング、チェックの3ステップが明記され、中立定数から再実行すると出荷値が再現できる。カーネル側にも実質的なメモリ節約がある。split-K の CTA は K 全体ではなく自分自身のスライスの scale と zero グループのみを一時保存し、スライスアラインメント時にはベクトル化ロードを用いる。_smem_bytes と _legal_configs は split ごとに共有メモリを計算するが、メインラインでは各 split の CTA が K 全体のグループを一時保存していた。もう一つの詳細は _ConfigSpace.smem_bytes がオプション上限より 1 KB 低い位置に置かれていることだ——ちょうど上限に達すると、17 個のスキャンされた設定が起動に失敗する。アロケータが動的共有メモリを切り上げる一方、バリアは静的なためである。#2826 は稠密セレクタの定数の再フィッティング方法（スキャン、フィッティング、チェックの3ステップ、ソフトアニーリングとホールドアウト検証を含む）をドキュメント文字列に書き込み、コードは変更しない。

出典：https://github.com/tile-ai/TileOPs/pull/2825、https://github.com/tile-ai/TileOPs/pull/2826

3.6 TileOPs レビュー中：短い varlen を KV グループごとにパック、L2 flush を読み取り専用に変更（#2819 / #2822）

#2819 は 11-03 23:37 に更新され、#2496 の一部であり、SM90 上の GQAVarlenFwdOp にパッキングされた warp 専用プログラムを追加した。1 つのリクエストがその KV ヘッドグループでビット単位で 1 つの tile を埋める方が tile を節約できる場合、キーと値は KV ヘッドごとに 1 回だけ読み取られる。パッキングカーネルの拒否ロジックは、コストモデルを用いて呼び出しのバッチ、パッキング後の query 総数とグループサイズを衡量し、形状が保証する節約が実測の 9% の単一ヘッド tile オーバーヘッドを超える場合にのみ引き受ける。広い item と狭い item の 2 種類は同じ発射の中で交互に取得され、ホストへの読み戻しは行わず、追加の同期も加えない。デバイス上で数えられた未分割 item が SM 数より少ない場合、単一 tile のパッキングリクエストはキースキャンを複数段に分割し、fp32 ワークスペースを介してマージする。ウィンドウ側にはさらに 2 件の実測收益がある。ウィンドウ行は 1% から 9% を節約し、インタリーブ RoPE 行は 3% から 6% を節約する。同時にメインライン上の 3 つの問題を修正した。sink 付きで score scale が約 2.4 より大きい場合に可視キーのない行が NaN を返す問題（単一ヘッド warp 専用と FP8 warp 専用の 2 つのプログラム）、フォールバックの query 回転がリクエスト最初のキーより前の行で先に RoPE テーブルを読む問題、および拒否された負の制御用例が呼び出し元フレームを参照サイクルに残すことで、あるテストの CUDA graph が別のテストのキャプチャ期間中に破棄される問題。テストノードは 173 から 197 に増えた。#2822 は 11-03 19:01 に更新され、bench_kernel の L2 flush を「int8 バッファに対する zero_()」から「プライベートな int32 バッファを読む、サイズは報告された L2 の 2 倍」に変更し、バッファは固定シードの CPU 生成器によって一度に抽出される。CUDA 上では Triton カーネルでこの読み取りを行う（要素ごとに 1 回のロード、16 KB ごとに 1 ブロックで 1 回の atomic_max）。torch 自身のリダクションを使わないのは、それを flush として扱うと次の Marlin w4a16 呼び出しが 1.3 から 1.5 マイクロ秒遅くなるからである。收益は H200、1500 MHz、CUPTI カーネル時間（3 回の実行それぞれ 30 回の呼び出しの中央値）上で、2048×11008 f32 の Sum は 27.81 から 21.38 マイクロ秒に、1024×11008 f32 の Ge（torch-compile）は 31.31 から 25.67 に、1024×4096 f32 は 6.30 から 5.47 に低下した。旧 flush は約 120 MB のダーティ行を残し、その書き戻しが次の計時呼び出しに落ち込むが、新 flush 自体は 71 から 445 個の L2 セクター（2 から 14 KB）しか書かない。ついでにメインラインのバグを 1 つ修正した。L2 サイズが未知の場合にログは 256 MB と書くが 512 MB を割り当てていた。

出典：https://github.com/tile-ai/TileOPs/pull/2819、https://github.com/tile-ai/TileOPs/pull/2822

3.7 採用側：FlashQLA と TileKernels はプッシュなし、ただし #3761 が修正したのは TileKernels の fused カーネル（11-04）

QwenLM/FlashQLA の最後のプッシュは 10-30 03:02 に止まり、deepseek-ai/TileKernels は 10-14 05:49 に止まり、両リポジトリとも本調査期間中にコミットはなく、採用側に新しいコードはない。ただし今期のメインラインのトップニュースは採用側と直接関係がある。#3761 で修正されたそのパフォーマンス回退のトリガーは、まさに TileKernels の mhc_pre_big_fuse カーネルであった。同じカーネルが B300、隠れ次元 1280 上で 25% のバリア占有予算を失った。これは、採用側のカーネルがすでにコンパイラ側の回帰プローブになっていることを示している。TileKernels がコードをプッシュしないこの数週間、そのカーネルは依然としてメインラインの変更を通じて受動的に「テストされ続けて」いる。エコシステムを追跡する人にとって、これは注目に値する逆方向のシグナルである。採用側の沈黙期間中でも、コンパイラの変更が彼らの形状で問題を起こす可能性はある。

出典：https://github.com/QwenLM/FlashQLA、https://github.com/deepseek-ai/TileKernels、https://github.com/tile-ai/tilelang/pull/3761

4. コミュニティ、チュートリアルとイベント

4.1 単日 7 件の新規 issue：1 件のプロダクショングレードの CUDA graph 問題、4 件は同一の外部報告者からのもの（11-03）

調査期間中、メインリポジトリで新規 issue が 7 件追加されました（番号 #3762 から #3775、ほかに昇騰リポジトリの自動デイリーテストレポートが 1 件）。技術的難易度が最も高いのは #3775 で、報告者は本番負荷を抱えるチームです。環境は pip インストールの tilelang 0.1.12、Python 3.12.3、H200 2 枚 + NVLink、ドライバ 580.178.04、CUDA 13 コンテナ、torch 2.14.1+cu130、サービス側は単一プロセス TP2（各デバイスに 1 つのワーカースレッド、それぞれ独自の CUDA コンテキスト）の sglang、モデルは GLM 系 MoE に DSA アテンションを加えたもの、経路上の TileLang カーネルは mhc_pre_big_fuse_with_norm、mhc_post、DSA prefill カーネルです。実行バックエンドを TILELANG_EXECUTION_BACKEND=nvrtc に強制すると（デフォルトの auto は tvm_ffi に解決されます）、tilelang の JIT カーネルが CUDA graph キャプチャ中に実行されることでキャプチャが破損し、しかも破損は tilelang の发射（発行）より下流の無関係なカーネルに現れます。バイト単位で同一の設定を 3 回再現すると、3 回とも最初の失敗点が異なりました。1 回目は cudaError700（無効な TMA ディスクリプタハンドル）が TMA ディスクリプタの初期化で送出され、その後無関係な発行点で反響し、1 回目は cublasGemmEx が CUBLAS_STATUS_EXECUTION_FAILED（純粋な BF16 入出力）を報告し、1 回目はホストフレームワークが JIT で生成した要素単位カーネル内での不正メモリアクセスでした。3 回とも decode の CUDA graph キャプチャ中（バッチサイズバケット 8）に発生し、いずれもサービス起動後約 100 秒で、/proc/environ とクリーンな対照との差分はその環境変数 1 つだけでした。これは最も特定が困難な類の問題です。現象はランダムで、失敗点は漂流し、誘因は環境変数 1 つだけで、報告者は 3 回の実行の対照条件を漏れなく記載しています。

残り 4 件は同一の外部報告者によるもので、環境は統一してメインブランチまたは 0.1.15 pip パッケージ、WSL 上の Linux、Python 3.13.5、RTX 4060 Laptop（sm_89）、ドライバ 566.26、CUDA 12.6 ランタイムで、いずれも再現スクリプトと完全な出力が添付されています。#3774 は 16×16 の共有メモリ置換が行シフトのみを実行し、列オフセットを失っている問題です（256 要素中 223 個が誤り、かつ「列シフトを失った」結果と完全に一致）。一方、同じ書き方の 32×32 は正しく動作します。#3773 は均一な while ループ内で 3 回の共有メモリコピーに GEMM を加えると不安定な結果になる問題です（直列ループ対照は 100 回で 0 誤り、while 版は 100 回で 21 回誤り、10 種類の異なる出力が出現）。#3763 はネストした整数の切り捨て除算 (x // -2) // 4 が int16、64 要素、32 スレッドで Let cannot bind the same var to two different values を報告する問題です。これら 3 件に #3775 を加えると、「コンシューマー向けカード + WSL + pip インストール」という組み合わせの集中的な障害報告面を構成します。残り 2 件：#3769 は B300 上の tcgen05 の int8 GEMM が ptxas に拒否される問題（審査中の #3770 による修正あり）、#3762 は昇騰 950 上で T.simd.vcvt が sat=False でも飽和してしまう問題で、詳細は次項のとおりです。

情報源：https://github.com/tile-ai/tilelang/issues/3775、https://github.com/tile-ai/tilelang/issues/3774、https://github.com/tile-ai/tilelang/issues/3773

4.2 昇騰 950 の vcvt 飽和セマンティクスと mlir-ascend デイリーテスト（#3762 / #1980）

#3762 は昇騰側のユーザーによりコミットされました。環境は Ascend950PR（NpuArch 3510 / dav-3510）、CANN 9.2.0-beta.2、ドライバ 25.7.rc1、torch_npu 2.12.0、ソースからインストールした TileLang（10-21 の e84ce3de）です。問題は T.simd.vcvt(src, "uint8", sat=False) が int32 ベクトルに対して sat=False を暗黙に無視し、常に飽和することです。負の値は 0 に、255 を超える値は 255 になり、エラーも警告も出さず、呼び出し側はラップアラウンドセマンティクスを使っていると思い込んでしまいます。ドキュメント文字列はこの状況を明確に記載しています——sat=False のとき範囲外の値はラップアラウンドする——期待値は -1 が 255、-128 が 128、300 が 44 ですが、実測では -1 が 0、-128 が 0、300 が 255 です。影響は int8 量子化の終盤のような f32 から int8 への変換経路に及びます。負の中間値が一度でも現れると、セマンティクスがラップアラウンドからゼロへの切り捨てに変わってしまいます。#1980 は昇騰リポジトリの自動デイリーテストで、11-03 09:35 の回は 3030 項目すべて合格、合格率 100%、ワークフロー実行番号 41553942797 です。この 2 つの数字を並べて見ると興味深い点があります。昇騰リポジトリのデイリーテストは満点のグリーンですが、デイリーテストでカバーできないセマンティクスの差異（例えば sat パラメータの実際の挙動）は依然として人手による報告に頼らざるを得ません。

ソース：https://github.com/tile-ai/tilelang/issues/3762、https://github.com/tile-ai/tilelang-ascend/issues/1980

4.3 ドキュメントサイトが3回更新、VS Code拡張がvsixに改名（11-03）

ドキュメントサイト tilelang.github.io は本調査期間内に Update docs コミットが3回（北京時間 11-03 13:11、13:28、11-04 00:59）あり、メインライン合流に伴う自動ビルドの定例更新で、内容は本期間に合流した数筆——#line の段階マッピング、ThreadSync スコープリスト、昇騰バッチコピーのドキュメント該当箇所——に対応する。TileOPs.github.io は本期間プッシュがなく、最新は依然 11-01 の「TileOPs メインラインに追随：オーバーラップとアラインメント保証、ビルド説明」である。VS Code 拡張リポジトリは 11-04 02:19 に「ドキュメントを更新し vsix に改名」の1筆があり、11-02 の「CUDA デバッグを TileLang 行マッピングに合わせる」に続くものだ：前期にメインラインへ合流した行単位アンカリングが生成コードの行番号の來源を変えたため、拡張側のデバッグ設定も追随して変更する必要があり、この1筆で成果物名とドキュメントをまとめて調整した。

ソース：https://github.com/tile-ai/tilelang.github.io、https://github.com/tile-ai/tilelang-vscode-edition

4.4 ニュース側：3チャネルとも期間内ヒットなし、外媒の昇騰トレーニング報道は口径不一（11-03）

Google News の中英語組み合わせクエリ、Hacker News、arXiv の3チャネルは本期間内にいずれも TileLang 関連の実質的ヒットがなかった：候補リスト15件のうち12件は GitHub リポジトリ発見によるもの、2件は主題と無関係な Show HN（探偵ゲームとハードウェア設計プロンプトツール）、1件は Wccftech の DeepSeek と昇騰に関する報道である。arXiv 上の all:tilelang の最新ヒットは依然 2026-07-24 の TileSight（tile 中心の GPU 性能分析モデル）であり、TileLang との聯合クエリは期間内に新論文がなかった。唯一記す価値のある業界側の消息には口径の衝突がある：Wccftech は DeepSeek が昇騰 950DT で構成される Atlas 950 SuperPoD 上で V5 をゼロからトレーニングすると述べ、同期の The Information の報道は「早くとも2026年末に 950DT 上で事前トレーニングを開始」であり、Bloomberg と金融時報の早期報道は 950DT の16万枚受注が主に推論向けで、此前に昇騰上でのトレーニングが挫折した後に Nvidia へ回退したと述べている。3種の説は「トレーニングに使うか否か、そしていつ使うか」について一致しておらず、いずれも華為または DeepSeek の公式確認を得ていない。ここに置くのは背景としてのみである：トレーニングが実現するか否かにかかわらず、昇騰側のコンパイラとオペレータへの投資（本期間 #3689、#1975、#1981）はこの経路の土台を敷いているのだ。

ソース：https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods、https://arxiv.org/

5. トレンド観察

第一に、性能回退が一級市民として扱われ始めており、その修正方法は結論よりも方法論の面で記録に値する。#3643 が正確性修正のために導入した ID 保留ルールは、それ自体が必要かつ正しいものである；#3761 はそれを回退せず、保留と再利用を異なる階層で併存させた——スコープを跨いで1つの ID を共有する前に、関与する述語が恒真であることをまず証明し、証明できなければスコープ付きカウントに戻す。著者の「宣言されたバリア資源数のみを変更する」という隔離実験と合わせ、この修正チェーンは「正確性」と「資源宣言」という2つの事柄を切り離した：後者は単独で再現・単独で検収できる変数である。この切り分け能力こそが、コンパイラプロジェクトが「動く」から「保守可能」へ向かう指標にほかならない。

第二に、昇騰のエンジニアリング化は2本の脚で同時に進んでいる。1本はアーキテクチャ：#1975 は2887ファイル、60万行の新規追加で TileLang-Ascend 全体をメインラインのマルチバックエンドフレームワークへ移し、arch="dav-2201" のような識別子でバックエンドを分流させ、A2A3 と A5 を同一リポジトリ内で直交して共存させることを目標とする；もう1本はパイプライン：mlir-ascend の #208 は Codex の初期化、接続、復元の3つのワークフローとタスク引き継ぎを .workspace/ に書き込み、メインリポジトリと昇騰リポジトリもそれぞれ Agent ディレクトリを持つ。「リポジトリを移行する」ことと「インテリジェントエージェントでリポジトリを移行する」ことを同じ日に見ると、昇騰側のボトルネックがすでに「オペレータ不足」から「一貫したエンジニアリング骨格と再利用可能な移行フローの不足」へ転換したことが見えてくる。

  1. オペレータライブラリの自己証明文化が深まっている。TileOPs で今期審査中の数件は、ほぼすべて「証拠」を PR に書き込んでいる。実際のルーティングで測定した Qwen3-30B-A3B EP8 のグループサイズ（48 層のうち偏りの中央寄りにある第 44 層を取得）、35 個のマスク構築による TIR がメインラインとバイト一致、L2 flush の変更には 46 行の実装比較表を添えて旧 flush が 120 MB のダーティ行を残すことを説明、W4A16 のコストモデルは「どの項目がデータで固定できず削除されたか」を明記。TileFoundry は別の道を歩んでいる——同期を証明可能な分析に仕立て、lowering に逆検証させ、分析が拒否したケースを直接列挙する（TMA と wgmma の完了セマンティクスが未宣言、スレッド境界を跨いで等しくないループ）。二つの道を合わせると同じ指向を示す。すなわち「証明していない」ことを認める方が、「動いた」を「正しい」と見なすよりましだということだ。

  2. リスク面はインストール形態に集中しつつある。1 日 7 件の新規 issue のうち 4 件は、pip パッケージを使い、WSL で実行し、コンシューマ向け sm_89 カードを使う同一の外部報告者からのもので、すべて再現スクリプトが添付されていた。#3775 は本番環境での「pip パッケージ + 複数 CUDA コンテキスト + nvrtc バックエンド + CUDA graph キャプチャ」の組み合わせである。これらはコンパイラの深層欠陥ではないが、新規ユーザーが初めて試す際の体験を決める。昇騰リポジトリの日次テスト 3030 項目が全緑でありながら sat=False というセマンティクスの偏差が依然として人手の報告に頼っていることと対照すると、テストカバレッジの形状と実際の使用形状の間にはまだ距離があることが分かる。

最後の観察：今期もニュース側は依然として空白で、15 件の候補のうち 12 件はリポジトリ発見によるものだった。TileLang の対外的な声量は依然として「DeepSeek と昇騰」のような出来事に紐づいており、日常的な進展はコミットの中だけで起きている。このプロジェクトを追う人にとって、これはデイリーレポートの主体が PR と issue であり、報道ではあり得ないことを意味する——これはすでに複数期連続で安定した形態となっている。

付、素材と検証説明

情報源検証表

情報源 検証結果
tile-ai 組織（29 リポジトリ） 調査期間内に 9 リポジトリへプッシュ
メインリポジトリ tilelang メインブランチに 5 件マージ；新規 issue 6 件；審査中更新 23 件；ドキュメントサイト 3 件自動ビルド
TileOPs メインブランチに 1 件マージ；審査中更新 7 件
TileFoundry メインブランチに 1 件マージ；審査中更新 2 件
tilelang-ascend デフォルトブランチ ascendc_pto は調査期間内に新規コミットなし；審査中更新 15 件；日次テスト 3030 項目すべて合格
tilelang-mlir-ascend harness ブランチに 1 件マージ（非メインライン）
tilelang-hygon メインブランチに新規コミットなし；審査中 #19 に更新あり
tilelang-vscode-edition 1 件マージ（ドキュメントと vsix の名称変更）
tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / tilelang-puzzles 調査期間内にプッシュなし
リリースページ（tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon） 調査期間内に新バージョンなし、tilelang の最新は依然として 09-30 の v0.1.15
FlashQLA / TileKernels 調査期間内にプッシュなし
Google News（中英文六組クエリ）/ Hacker News / arXiv 調査期間内に実質的ヒットなし

完全情報源リスト

[23] tilelang-ascend Chunk O フォワード移行（#1866）— https://github.com/tile-ai/tilelang-ascend/pull/1866