レポート期間:2026年11月2日(月)から11月6日(金)まで、計5営業日 素材情報源:本誌の毎日TileLangコミュニティ動態デイリーレポート5期(11-02、11-03、11-04、11-05、11-06)。5期の収集ウィンドウは端から端まで連続しており、合計で10-30 07:00から11-06 07:00までの168時間をカバーする(月曜日の分は72時間ウィンドウで前週末をカバー、火曜日から金曜日までは24時間)。前回のウィークリーレポート(10-26から10-30)とは重複しない。なお、tile-ai組織の各リポジトリに対する168時間ウィンドウの補足調査、各リポジトリのリリースページとタグ、TileLang×RISC-Vの4情報源巡回再確認を別途実施した(詳細は付録を参照)


1. 今週の主要ニュース

  • 昇騰950の永続化FlashMLAカーネルを直接メインラインに書き込み、別リポジトリへの移植ではなくコンパイラがスケジューリング可能な記述方式を採用(11-02マージ、#3678):1件の変更で1997行追加、17行削除、20ファイル、42コミットに及び、BF16 prefillとコンパクトFP8 decodeの2経路をカバーし、オプションのFP8およびFP4追加KVキャッシュ、attention sink、不正インデックス、queryごとの独立長をサポートする。アテンションループ、オンラインsoftmax、逆量子化、スケジューリングはすべてTileLangで表現し、T.Stageによってコンパイラにバッファバージョンと同期を管理させる。QKとPVはL0Bストレージリングの同一領域を異なるlayoutビューとして共有し、L1の配置はT.annotate_buffer_addressesで直接指定する。prefillは1つのアキュムレータ割り当て内で4つの互いに素な出力スライスに分割し、正規化が1スライス完了するごとに1スライスを排出する。この記述方式を成立させるため、コンパイラ側では同時期に3つの機能を補った:L0CからUBへのコピーでNZ出力を保持する、各AIVストリームを自身のパイプラインより前に選択する、明示的スケジューリングされたステージが物理発行フェーズで同一の場合にソースコード順序を維持する。対象形状は非常に狭く定義されている:ヘッド数64、DqkとDvはいずれも512、topkは64の倍数。情報源:https://github.com/tile-ai/tilelang/pull/3678

  • ThreadSyncを週に3回改修、まず移行のために道を譲り、その後性能と証明を補う(11-01、11-04、11-05):月曜日の#3733は3週間前にマージされた正確なスレッドカウント(#3291)を丸ごとロールバックした。理由は率直に書かれている——その変更は高コストなカウントクエリをフルブロックフェンスチェックの前に移動させ、すべてのスレッドがフェンスに到達する場合でも証明を実行する必要があり、TVMの求解予算を使い果たして部分カウントを返し、正当なremapおよびsequence補助kernelが「Cannot prove a thread rectangle」でコンパイル失敗する。商ガード、正確な参加スレッド境界、非矩形スレッド集合という3種類の正当性改善も併せて撤回され、#3732は移行後の再最適化のためのトラッキングチケットとして残された。木曜日の#3761はロールバックに派生した副作用を処理する:#3643の名前付きバリア重複番号を修正するため、同一グループのスレッドはスコープを跨いでもIDを再利用できなくなり、TileKernelsのmhc_pre_big_fuseカーネルはブロックあたり96スレッドのため、2段階のcp.asyncパイプライン上で余分に1つのIDを取得し、コンパイル生成物が宣言するバリアリソースが4から5に増え、B300上で隠れ次元1280のバリア占有上限が16ブロック/SMから12ブロック/SMに低下した。修正方法はThreadSyncが順序付きスコープリストを受け入れるようにし、スコープを跨いでIDを共有する前に新しいアナライザで参加述語が恒成立することを証明するというもの。受け入れ基準はリソース宣言が4に戻り、占有が16ブロック/SMに戻ること。#3643の修正はロールバックされていない。金曜日に審査中の#3787は証明器の盲点を引き続き補う:192スレッド下でtx < 160 and tx // 32 % 2 == 0のような「境界項とwarpインデックスの剰余項をandで結合した」形態が誤検出される。根本原因はFloorDivとFloorModのエンコーディングが「上界+剰余」の組み合わせで2つの除算の関係を確立できないことであり、修正方法はthreadIdx.xをwarp_index * warp_size + laneに書き換えること。情報源:https://github.com/tile-ai/tilelang/pull/3733、https://github.com/tile-ai/tilelang/pull/3761、https://github.com/tile-ai/tilelang/pull/3787

  • TIRx移行が「口頭の理由」から1週間に3つの確固たる証拠へ、ベースの世代交代が公開準備段階に(11-03~11-06):火曜日の#3736でtilelang/language/ast/、tilelang/language/parser/、tilelang/language/tir/entry.pyが一括削除された——初期のローカルTVMScriptパーサー分岐の残骸であり、1行追加、3384行削除、10ファイル、言語層は上流のtvm.tirx.script.parserに切り替わり、ASTと文のオーバーライド登録はtilelang.language.overridesで行われる。木曜日の#3752が依存するtile-ai/tvm#91と金曜日の#3752により、Z3のrlimitがハードコードされた100,000からkernelごとに設定可能になり、作者の言葉によれば「これによりZ3証明器に関連する多くの問題が間接的に修正された」。金曜日の#3803では完全な移行草案(draft)が提出された:vendoredのTVMを独立したApache TIRxに置き換え、フロントエンド、IR変換、コード生成、JIT適応、パッケージングをすべてネイティブオブジェクトモデルに移行し、apache/tvm-tirxの7e49a148とapache/tvm-ffiの3e906bc6をpinし、依存ツリーはクリーンな上流を維持しパッチを注入しない。3つの線は同じことを指している:メインラインのこの段階の第一優先はベースの交換であり、その他の最適化はすべてそれに道を譲る。出典:https://github.com/tile-ai/tilelang/pull/3736、https://github.com/tile-ai/tilelang/pull/3752、https://github.com/tile-ai/tilelang/pull/3803

  • 「静かな誤コード生成」が今週の主要な欠陥類型となり、しかも4つのリリースバージョンにまたがる(11-03~11-05):最も重いのは#3808——条件関連のスカラー化の後、共有メモリepilogueで__syncthreads()が1つ欠け、0.1.12、0.1.13、0.1.14、0.1.15の4バージョンすべてで静かに誤ったCUDAが生成され、コンパイル時に一切エラーが報告されず、報告者は完全に実行可能な再現スクリプトを提供した。同族には#3800(warp specializationをオフにすると、T.Pipelinedがtcgen05のbarrier待機の前に次のstageをリロードし、MMAが書き換えられたoperandを読む)、#3721(sm_100a上でTMEMに対してT.gemmを実行し、threads/32がNを整除できない場合PTXにtcgen05.mmaが1つもなく、積が静かに破棄される)、#3775(本番環境のTP2 sglangにnvrtcバックエンドを追加するとCUDA graphキャプチャが破損し、バイト一致の設定を3回再現すると3つの異なる最初の失敗点が出る)、#1991と#1992(昇騰910_9362上でthreads=2の場合AscendVidReductionが注入されたtmp workspaceも無条件に半減し、sin/cosが間接アドレッシングtile上で3978/4096要素を誤計算、すなわち97.1%)がある。この族はすべてエラーを報告せず、通常のテストをすべて通過し、的を絞った再現でのみ捕捉できる。出典:https://github.com/tile-ai/tilelang/issues/3808、https://github.com/tile-ai/tilelang/issues/3800、https://github.com/tile-ai/tilelang-ascend/issues/1991

  • TileFoundry、同期を「分析で証明できる」から「発射は証明と一致しなければならない」へ進める(11-02 から 11-06):#283 は「共有メモリアクセスをどのように順序付けるべきか」を証明可能な分析層にした——lowering の発射順序に従って各関数をたどり、衝突アクセスの前に commit、wait、スコープ内フェンスを計画し、83 個のスケジューリングフィクスチャのうち 52 個を計画、31 個を拒否した(24 個は TMA コピー、6 個は wgmma を使用し、両者の完了セマンティクスはまだ宣言されていない)。デフォルトでは呼び出し元からの要求がないため、finalize される 75 個のフィクスチャのバイトは main と完全に同一である。木曜にマージされた #285 はその実行側である:CUDA lowering は同期計画に従って発射し、項目ごとに逆検証を行い、不一致があれば LoweringError を投げる。単一の sm_89 上でビルド可能な 7 個のプログラムのうち 6 個は参照と完全に一致し、synccheck と memcheck はすべてクリーンであったが、同じ 7 個を main の方法で lower すると 4 個の結果が誤りとなり、直接の症状は sm_89 の E4M3 GEMM の最大偏差 659 である。金曜にマージされた #286 は WGMMA の非同期完了セマンティクスを補い、水曜にマージされた #284 は FSDP と HSDP のトレーニングセマンティクスを HIR 評価器に組み込み、4 枚の sm90 上で実際の 2×2 PyTorch HSDP と整合させた。著者は #286 に明確な限定を残している:行列命令エミッタはまだ接続されておらず、このコミットは実行の正しさを証明するものではなく、性能も与えない。出典:https://github.com/tile-ai/TileFoundry/pull/283、https://github.com/tile-ai/TileFoundry/pull/285、https://github.com/tile-ai/TileFoundry/pull/286

  • TileOPs は 5 つの期間で 15、7、1、7、9 件をマージし、競争点は「オペレータ数」から完全に「セレクタと準備コスト」へ移った(全週):月曜期間の 15 件のうち 13 件は 11-01 の 1 日に集中し、ほぼすべてが H200 実測表と「manifest に測定済み行がなければフォールバック」という宣言を伴っていた。木曜の 7 件、金曜の 9 件でリズムは連休後で最も密な状態に戻った。3 つのハードな数字がこの週の方向性を要約している:GQA 稠密フォワードは 1 つの実測ルールで decode または prefill プログラムを選択し、最速候補に対する選択リグレットは 1.818 から 1.017 に低下した(#2808)。warp 専用 decode は softmax と QK、PV を重ね合わせ、影響を受ける 478 回のスキャン呼び出しの幾何平均レイテンシは main の 0.952 倍であった(#2817)。ブロードキャストオペランドのインプレース読み取りにより、lerp、where、clamp、masked_fill の 5 種類のオペレータのブロードキャスト行の所要時間はメインラインの 0.13 から 0.44 倍に低下した(#2832)。同期間には roofline 计价の書き換え(物理的に不可能な 3 つの帯域幅行を 7.48、7.50、5.30 TB/s から 0.62、0.62、2.65 TB/s に修正、#2830)と 3 段階タイルファクタの個別フィッティング(選択リグレットは 0.89% から 0.06% に低下、#2836)もある。出典:https://github.com/tile-ai/TileOPs/pull/2808、https://github.com/tile-ai/TileOPs/pull/2832、https://github.com/tile-ai/TileOPs/pull/2830

  • 国産バックエンドは「自研ブランチ」から「公開レビュー」へ向かい、海光が初めて正式な PR を開く(11-05):海光リポジトリのメインブランチは 11-02 の #32(gfx948 を gfx946 互換として処理)で止まっており、今週の動きはまず feat/hcu-scaled-cast ブランチ上に 4 件積み上げられた(E8M0 スケーリングの FP8/BF8/FP4 変換、64 bank の LDS 読み取りフェーズ、混合 LDS wrap の最大 2 経路の証明、物理読み取り幅に基づく混合チェック)。金曜にはさらに LDS フェーズ修正を 3 件落とし、09:08 と 19:52 にそれぞれ #33(scaled cast)と #34(reduction の XOR shuffle に DPP と lane swap を使用)の 2 つの PR を開いた——これは国産バックエンドの中で「自研ブランチから公開レビュー」という経路を完走した最初のものである。mlir-ascend は同期間に BishengIR ツールチェーンアダプタを contrib に移し(#209)、さらに #203、#211、#212 の 3 つの小さい PR で tl.sub/mul/div/max/min、tl.shl、tl.shr を一つずつ TLIR 方言に移した。摩尔線程は 11-05 に README ドキュメントのコミットを 1 件落とし、09-30 以来の完全な沈黙を終えた。出典:https://github.com/tile-ai/tilelang-hygon/pull/33、https://github.com/tile-ai/tilelang-hygon/pull/34、https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

  • ハードウェアサポートが両方向へ同時拡大:Blackwell の新 SKU と国産アクセラレータ(全週):一方は新しいカード。B300 で報告された sm_103a は 2 つのパスでそれぞれ引っかかる——tcgen05.mma.kind::i8 は sm_100a と sm_110a にのみ存在し、カーネルは lowering を経て ptxas で Feature '.kind::i8' not supported を報告(#3770)。CuTeDSL の TMA ヘルパー関数はアーキテクチャのホワイトリストを維持しており、B300 は DSL に入る前に拒否される(#3772、sm_103a と gr100 の sm_107a を追加後、B300 では 8 通過 1 スキップ、gr100 では 13 通過 1 スキップ)。SM120 上の FP8 と fp32 累加の GEMM は常に SM89 の mma.sync アトミックを使用しており、ue8m0 単位スケーリング付きの block-scaled MMA に置き換えた後、RTX 5090 上で 8192 立方が 463 から 727 TFLOPS に上昇(審査中 #3790)。もう一方は国産カード:海光は 64 bank の LDS フェーズでマイクロアーキテクチャレベルのチューニングを行い、昇騰は Cube と Vector のコア比を固定 1:2 から設定可能に変更し(#3783)、AutoSchedule の周期下限を厳格化して Z3 求解時間を 10.04% 削減(#3776)。出典:https://github.com/tile-ai/tilelang/pull/3772、https://github.com/tile-ai/tilelang/pull/3790、https://github.com/tile-ai/tilelang/pull/3783


2. コアプロジェクト進捗(メインリポジトリ tilelang)

バージョンとリリース

  • メインリポジトリの最新タグとリリースは依然として v0.1.15(09-30 リリース)であり、全週にわたって新規タグも新規リリースページ項目もなく、今週金曜で満 37 日に達した。
  • 組織内の 2 つの新バージョンはいずれも先週のまま:TileOPs は v0.1.0(10-27)、TileFoundry は 0.0.4rc1(10-25)で、今週はいずれも正式版に移行していない。
  • 適応ラインのタグは全週動きなし:tilelang-ascend v0.1.2.000(09-09)、tilelang-metax v0.1.15(10-16)、tilelang-musa v0.1.15+musa.1(09-30)、tilelang-mlir-ascend は v0.1.2.020 系列で停止。
  • マージ量とペース:月曜の調査期間は 72 時間(10-30 07:00 から 11-02 07:00)で、マージ 35 件;火曜から金曜の 4 つの 24 時間調査期間はそれぞれ 18、8、12、18 件で、5 期間合計 91 件。審査中更新は順に 82、44、48、42、57 件で、合計 273 件;新規 issue は順に 11、8、7、4、5 件で、合計 35 件。メインブランチのコミット分布:tilelang 7、5、5、3、5 件、TileOPs 15、7、1、7、9 件、TileFoundry 3、1、1、1、1 件、tilelang-ascend 5、4、0、1、2 件、tilelang-mlir-ascend 1、0、1、2、0 件、tilelang-hygon メインブランチ 2、1、0、0、0 件(別途機能ブランチのコミットあり)、tilelang-lsp と VS Code 拡張は各 1 件、tilelang-musa 1 件。
  • ペースとしては今週は「週初に清算、週中に底打ち、週末に回復」の形を示した:水曜はわずか 8 件のマージで今週の最低点となり、金曜には 18 件に戻った;審査中キューは週初の 82 件から 42 件まで低下した後、金曜に 57 件へと回復した。これは週初に滞留したレビューが今週大量に消化されたことを示している。

コンパイラ、言語層と実行パス

  • 昇騰 950 FlashMLA カーネルと付随するコンパイラ機能(#3678、11-02 マージ、20 ファイル +1997/-17):要聞の第1条を参照。注目すべきは、このコミットが「明示的スケジューリング」に必要な三つの低レベル機能を同時にメインラインへ追加しており、昇騰ブランチ内でパッチを当てているわけではない点である——L0C から UB へのコピーで NZ 出力を保持し、各 AIV ストリームをパイプラインの前に選択し、物理発射フェーズが同一の場合はソースコード順序を維持する。
  • cast デカップリングのベクトル幅リグレッションを連続修正する2コミット(#3734 11-02 マージ、#3759 11-05 マージ):#3734 は SM100 上の ViT フォワードにおける Cannot convert type float32x16 to CUDA type を修正——ベクトル化プランナが 256 ビット BF16 ストアのために 16 レーンを選択し cast 制約を遅延させ、#3579 がデカップリングでメモリ読み取りのガードをスキップさせた後も、最終リーガライズでは遅延戦略が引き続き使われ、残留した FP32 式が 512 ビットベクトルに膨張していた。修正方法は defer_cast_constraints を IR 属性ではなくプランナの C++ パラメータにすることで、ガードの FP32 から BF16 への変換は 8 レーンに低減され、デカップリングに成功した BF16 コピーは引き続き 16 レーンを維持する。#3759 はその残りを処理する:条件内の int64 cast が以前はデータベクトル幅を 16 レーンから 4 レーンに圧縮していたが、その条件が 16 レーン上で完全に一致する場合でも同様であった。修正方法は条件一貫性解析を cast 幅の免除に再利用することで、if 文、tirx.if_then_else、スカラ述語の Select 条件を個別に扱い、節内の cast は引き続き制限される。
  • ThreadSync の三段階的進化(#3733、#3761、#3787):要聞の第2条を参照。#3733 は明確な暫定措置であり、原文には「TIRX 移行のために道を譲る」と明記され、付随する TVM fork 最適化提案 tile-ai/tvm#89 はすでにクローズされている。#3761 の 401 行の変更で最も記録に値するのは帰属実験である:著者はコンパイル成果物で宣言されたバリアリソース数を単独で変更し、実行可能命令には一切触れずに同じ遅延を再現し、「性能リグレッション」を感覚から特定可能な単一変数へと変えた。
  • 言語層とフロントエンドの削減:#3736 は 3384 行の遺留 TVMScript パーサフォークを削除し(要聞の第3条を参照)、そのリスクを具体的に記述している——誰かが直接 tilelang.language.parser.operation をインポートすると、tirx.PrimExpr と tirx.IterVar 上に互いに衝突するオペレータを登録してしまう;#3760 はネストした T.macro に有界なプロセス内変換キャッシュを導入し、コンテキストに依存しない変換結果のみを保存し、マクロインスタンスは現在の globals とクロージャ値を再バインドすることで、「カーネルはキャッシュにヒットしたがマクロは依然として重複解析・コンパイルされる」問題を解決する;#3802 は T.annotate_host_mapped_buffers を追加し、注釈されたパラメータのインデックスを PrimFunc の tl.host_mapped_params 属性に書き込み、実行時には cuMemHostGetFlags で登録状態を検証し、cuMemHostGetDevicePointer でデバイス可視アドレスを解決し、通常のページ可能メモリは引き続き拒否される。
  • デバッグと可観測性:#3703 は #line 指令を「TIR 文ごとに先頭行のみをアンカーする」から段階的発行(0 はオフ、1 は文レベル、2 以上は行ごとに再アンカー)に変更し、実装は CUDA、HIP、CPU C、C host、Metal、Ascend が共用する C 系出力ストリームに置かれ、新たに TILELANG_EMIT_LINE_DIRECTIVES を Python 側のデフォルト値として追加した;著者は CUDA と HIP のユースケースがローカルでスキップされ、cuda-gdb の実際の動作は直接検証されていないと注記している。同日 LSP リポジトリがフォローアップの1コミットを行い、CUDA デバッグを新しい行マッピングに整合させた;木曜日には VS Code 拡張リポジトリがさらにフォローアップの1コミットを行い、ドキュメントを更新し成果物を vsix に改名した。
  • 自動チューニングの基準:#3546 は set_autotune_inputs の供給クロージャを「呼び出し点から決して渡されないパラメータ」から get_current_device() に変更した。以前は CUDA のないマシンでフォールバック分岐が各キャプチャテンソルを CPU に移動し、チューニングの各ラウンドでランチャーのデバイス不一致により失敗していた;デュアルカード Ascend950 上では3件の昇騰自動チューニングリグレッションが修正前は全滅、修正後は全通過した。同族の #37

84 AscendC CI 上で、クロスデバイスコピープローブの前後にデバイス同期を1回ずつ追加し、非同期エラーが try ... except RuntimeError のスコープ内に戻るようにした——これまではそれが一歩遅れて浮上し、無関係な torch.arange() が4バイトのmemcpyサイズ不一致を報告したように見えていた。

  • Z3 求解リソースをハードコードから調整可能に(#3752、11-06 マージ、tile-ai/tvm#91 に依存):rlimit のデフォルトは引き続き 100,000 で、設定しない場合の動作は変わらない。これは木曜日に昇騰側で AutoSchedule の周期下限を厳格化したこと(#3776、Z3 求解時間が 10.04% 減少)と同じ方向の両端である。一方は探索予算を変更し、もう一方は探索開始点を変更している。

正確性と欠陥修正

  • ThreadSync とバリア:#3733 の一括リバート(要聞参照)と #3761 のクロスドメイン再利用(要聞参照)のほか、審査中の #3765 は「単一スレッドインデックスの境界ガード付き部分バリア」をスレッド次元が warp 数の整数倍である場合に明示的に受け入れ、#3767 は複数インデックスの fragment アクセスにおける StructuralEqual 内部アサーションを 2 つのインデックス式を持つ ValueError に置き換え、#3786 は T.copy の目的領域の範囲外アクセスをコンパイル時エラーに変更し、#3788 は ProveFragmentContains の高速パスを修正した——大きな fragment を完全コピーする際に真を直接返し、スレッド帰属チェックをスキップするため、「小さな fragment がスレッド 0 から 127 を所有し、大きな fragment が 0 から 95 のみを所有する」場合が包含成立と誤判定されていた。
  • 静かに誤ったコードを生成する一族:#3808、#3800、#3721、#3775、#1991/#1992(要聞 4 番目参照)のほか、今週はさらに #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 インストール」という組み合わせの集中的な障害報告面を形成している。
  • 依存性分析と同期挿入の根本原因:#3749 は今週最も技術的に高度な欠陥報告である——連続する 2 箇所の SIMD ロードが、より早い GM から UB へのコピーに対する依存性を失った。根本原因は MemoryAccessDetector::LogicalPointerRegion が再構築した論理領域が物理ロード範囲の一部を欠落させていることにある:最小再現では work.shape が (2, 128)、生成された読み取り記述子は物理 [64, 192) をカバーするが、再構築された領域は [0, 128) であるため、依存性分析は入力 DMA とロードが無関係と判断し、AutoSchedule は入力コピーをロードの後ろに移動し、InsertSync も MTE2 から V への set/wait を発行しない。同一期間にメインリポジトリでは昇騰 AscendC と PTO の 2 バックエンドのコード生成と同期挿入を指す issue が新たに 8 件追加され(#3740 から #3749)、うち 7 件は同一報告者によるもので、いずれも再現可能な prim_func と行単位で対照した生成コードが添付されている。
  • 境界条件が破られると内部エラーを報告:#3738 は「padding 後の共有 tile より小さいグローバル領域からの T.copy(例えば K=100 を BK=112 にコピー)が TMA の一括コピーと判定されてハード失敗する」問題を修正;#3766 は cta_group::2 の TMA バイトカウントを修正——CTA ペアのバイトがすべて偶数 CTA のバリアに落ちるが、cluster バリア付きの tma_copy は従来 rank 0 に対してのみ cluster_size × tile_bytes を期待しており、(2,1,1) と (1,2,1) では偶然正しく、(4,1,1) と (2,2,1) では誤っていた。
  • オペレータとサンプルの境界:#3709 は CPU スカラー GEMM が目的領域のみをクリアするようにし、#3712 は LegalizeNegativeIndex が到達不能な分岐をスキップするようにし;#3714 は fp32 から tf32 への変換に最近接偶数丸め(cvt.rn.tf32.f32、非飽和は SM90 以上、飽和は SM100 以上が要求)を追加し;#3702 は複数のアテンションサンプルが shared_memory_per_block_optin から tile を推論するようにし、RTX PRO 6000(sm_120)上で CUDA 9 項目が「2 成功 7 失敗」から 9 全成功に変わり、H100 は 227 KiB の選択可能上限を持ちながら元の tile のままである。
  • 昇騰側の 2 つの意味論的偏差:#3762 は T.simd.vcvt(src, "uint8", sat=False) が int32 ベクトルに対して sat=False を静かに無視し常に飽和することを報告。ドキュメントに書かれた回り込み意味論(-1 が 255、-128 が 128、300 が 44 に)は実測では切り捨て(-1 が 0、-128 が 0、300 が 255 に)となり、int8 量子化の終盤における f32 から int8 への変換経路に影響する;#3789 は PTO バックエンドが整数 T.atomic_add に対して符号付きの pto.si32/si64 を発行するが、PTO のアトミックインターフェースは符号なしキャリアのみを受け入れることを報告し、当日中に #3792 と #3794 の 2 方向の修正案が出た。

ランタイム、CI とツールチェーン

  • CI の2種類の安定性問題はどちらも対応済み:#3784 は AscendC CI の自動チューニング供給テストケースを修正(前述)、#2833 は TileOPs の RUNNER_VERSION を 2.336.0 から 2.338.2 に引き上げ(GitHub が旧 runner を廃止した後 gpu-smoke がすべてキュー待ちになったが、entrypoint.sh は --disableupdate 付きのため自己アップグレードできなかった)、#2827 は TileOPs の off-CUDA スイートを夜間ジョブに移した(従来 preflight ではこの2ステップが ubuntu-latest 上で13〜18分かかっていたが、移行後 preflight は収集チェックのみを残し、テストレポートの破棄と cuda_only の必失敗チェックも外れ、約3分になった)。
  • 並行マージに post-merge ゲートが欠けている実証:#2824 が21:06にマージされた際、grouped GEMM の roofline 行を sizes の宣言必須かつそこから G を導出する形に変更したが、#2830 は21:08にマージされ、依然として G のみで sizes がないままだった。2件は約2.5分差で、それぞれの CI はグリーンだったが、マージ後の main の GPU Smoke が即座に RowError: row misses ['sizes'] を報告し、当日 #2837 で補修した。これは今週最も記録に値するプロセス上の記録である。単一の CI 通過では「2件の意味的相交」をカバーできない。
  • ドキュメントとローカライゼーション:tilelang.github.io は今週の4つの収集ウィンドウでそれぞれ 3、5、3、3 件の Update docs 自動ビルドがあり、内容は #line の階層マッピング、ThreadSync スコープ一覧、host-mapped 注釈、TF32 rounding、Z3 設定などのインターフェース変更に対応する。TileOPs.github.io は1回プッシュがあり、TileOPs-nightly はブランチプッシュのみでメインブランチに新規コミットはない。
  • 昇騰側のツールチェーン詳細:#1956 は AI Core 例外ダンプを CANN 9.2.0 に互換化(acldumpTensorInfo の必須フィールドを補完、空アドレスとゼロサイズテンソルをスキップ、登録インターフェースを廃止済みの aclrtSetExceptionInfoCallback から aclrtExceptionInfoCallbackRegister に変更、バージョン要件を 9.3 から 9.2 に引き下げ);#1932 は A2/A3 のパフォーマンスチューニング技能を Ascend950 ワークフロー構造に揃え、チップ種別識別とチップ情報検索の2技能を新規追加した。

レビューキューの観察

  • 最大の一笔は依然として昇騰リポジトリ全体の移行(#1975):11-04 00:34 更新時点で追加 607317 行、削除 97595 行、2887 ファイルに達し、TileLang-Ascend をメインラインのマルチバックエンドアーキテクチャに移行、既存の昇騰コンポーネントは ascend/a2a3 配下に移し、arch="dav-2201" とメインラインの A5 パスで分流する。作者は説明の中で TODO を実に率直に書いている。3ディレクトリ下のファイルと元ファイルとの差異はまだ収束待ち、静的チェックとフォーマット CI は復旧待ち、ドキュメントと Agent 関連は移行待ち。
  • TIRx 移行草案(#3803)は全週で最も重みのある審査中項目:これは今後の TVM 側修正がどのようなリズムでメインラインに入るかを決め、昇騰、海光、mlir-ascend のグルーレイヤーを追随変更するかどうかも決める。
  • TileOPs の審査中キューはセレクタとレイアウトに集中:#2836 は coop2_swap の3段階タイルに個別フィッティング係数(1.36)を与え、#2838 は「K を連続次元とみなし、TN レイアウトの K=1 の正当な呼び出しを否決してしまう」問題を修正し、#2839 は FP8 グループ GEMM で TMA を使って A のスケーリングをロードし、K ≤ 1536 の非融合タイトグループに 64×256 ping-pong タイルを割り当て、#2840 は SM89 の prefill sinks を補完し sink の有無をコンパイルインスタンス同一性に組み込む。
  • 審査中数の推移:82、44、48、42、57 笔。週初めの高水準は先週の積压によるもので、週中に42笔まで下がった後金曜に回升し、マージ量の週末回升と同期しており、通常の清算—起票リズムである。

3. マルチバックエンド適配(昇騰 / 海光 / 摩尔線程 / 沐曦)

昇騰(Ascend):「適配目標」から「初回リリース目標」へ

  • 初出側:FlashMLA カーネル(#3678)は昇騰 950 をターゲットにメインラインへ直接書き込まれ、T.Stage、T.annotate_buffer_addresses といった明示的スケジューリング向けの低レベル機能を用いており、別リポジトリを開いて移植するものではない。ターゲット形状はヘッド数 64、Dqk と Dv は 512、topk は 64 の倍数。付随するバッチ転送機能は 11-03 に補完された:#3689 により [B, M, N] 領域を一度の T.copy で記述でき、lowering がバッチ ND2NZ ディスクリプタを生成し、ソースのバッチと行ストライドを保持し、転置バッチは DN2NZ で処理され、GM 境界を超えるソースバッチはスキップされる。検証は Ascend950DT と CANN 9.2 上で行われ、50 の AscendC サンプルテストと 1733 の昇騰単体テストが通過し、うち 28 の BF16/FP32 デバイステストケースがバッチと M/N の端数ブロック、バッチ全体の越境、非連続ソースストライド、転送先ストレージがソースより大きいケース、矩形転置をカバーしている。
  • 同期の正確性:#1981 は copy_pipe_to_l1 の同期注釈を修正——A2/A3 の PTO クロスコア V→C 通信は GM を中継し、三段の転送はそれぞれ AIV の MTE3、AIC の MTE2、AIC の MTE1 に属するが、operation_config.h では PIPE_MTE3 と注釈されており、自動同期は AIC 自コアの MTE3 を待つがそのキューは空でイベントが即座にセットされ、MTE1 の L1 読み出しが MTE2 の書き込み完了前に許可される可能性がある。再現証拠は二つ:生成ソース内の PIPE_MTE3 を PIPE_MTE2 に変更すると 10/10 すべて正確に復帰し、ローカルでは最新メインライン上で 10/10 回実行エラー、単回最大 16384/16384 要素の不一致。修正法は名前を直接置換するのではなく分割する:copy_pipe_to_l1 は PIPE_MTE2 に注釈変更し、新たに copy_pipe_to_l1_bind を PIPE_MTE1 と注釈(A5 直結パス)し、A2/A3 を修正しつつ A5 でエラーにならないようにする。
  • メインラインへの収束:#1971 の RFC は同期実装の置換ロードマップを示す——メインラインの同期コアと独立した Z3 証明サービスを移植し、旧 TIR アダプテーション層を通じて既存のコンパイルパイプラインに接続し、TVM、スケジューリング、メモリプランニング、AscendC/PTO codegen、実行 ABI を維持し、受け入れ通過後にデフォルト実装へ切り替えて旧アルゴリズムを削除する。付随する #1965 はメインラインの z3 版 insert sync を Ascend リポジトリへ移行し、#1975 はリポジトリ級スケルトンの移行である。三つは同一パスの三層である。同時期に #3752 は Z3 の rlimit を kernel ごとに設定可能にし、#3735 は thread sync 内でモジュール化インデックスの非交差性を証明することを提案しており、まさに「求解リソースの設定」と「証明が必要な形状の削減」の両端から先週のロールバックが残した証明予算の問題に対処している。
  • ホストオーバーヘッドの系列:#1948 は CythonKernelWrapper.forward() 内で毎回無条件に構築される TVM Analyzer() を遅延構築・必要時再利用に変更し、910B3 上で vec add 単一ブロックが 442〜467 マイクロ秒から 105.1 マイクロ秒に、32 ブロックが 441〜493 から 108.0 に、kpool 圧縮が 447〜450 から 111.3 に低下し、素の ctypes 呼び出しは約 12.5 マイクロ秒でほぼ不変。引き継いだ #1967 は cython wrapper 内でシンボル形状の解析結果をキャッシュし、シンボルパスのコストを引き続き削減する。同系列には審査中の #3757 もある(プリント名ではなく Var オブジェクトをキーとし、同名異変数が同一シンボルと見なされるのを防ぐ)。
  • オペレータ供給:#1827 は WY Fast 順伝播を新規追加し、#1828 は wy_fast_bwd_split のセグメント化逆伝播を移行・補完(両者はいずれも 11-02 02:33 前後に 1 分未満の差でマージされた)。これにより線形注意力のこのオペレータチェーンを昇騰上で学習可能なまでに補完した。#1866 は Chunk O 順伝播(chunk ゲーテッド線形注意力)を移行し、開発者モードに 4 つの pass 設定を加え、単一の @tilelang.jit カーネルに固定 20 コアの grid-stride ループを組み合わせ、オンチップ融合、ゼロ GM workspace、ゼロホスト ATen オペレータを実現した。#1879 は MHA 順伝播サンプルを新規追加(BSHD レイアウト、Expert モードの三分派 MMA パイプライン、wrapper に T.* コードを一切含まない)。cann-bench 20 ケースの精度はすべて通過し、torch-npu ベースラインとの比較で平均加速比 0.97 倍、総合スコア 73.13(ローカル 910B2 基準)——0.97 倍はまだ純益がないことを意味するが、骨格は示された。#1966 の審査中の GLM-5.3 kpool fp8_mqa_logits サンプルはデュアルカーネルに AIV/Cube の分業を加えたもので、ヘッド数 32、次元 128、ページサイズ 64。
  • エンジニアリング基準:#1954 は T.tile.bitwise_rshift の int64/uint64 におけるコンパイル時セグメンテーション違反を、実際の dtype とサポートリストを伴う ValueError に変更した(int16/uint16/int32/uint32 のみ許可)。「64 ビットを新規追加しない

右シフト対応」;#3783 は Cube と Vector のコア比を明示的に設定可能またはデバイスごとに導出可能にし(vector_core_num / cube_core_num、取得できない場合は 2 にフォールバック)、選択されたグループ数を sid ドメイン、タスク割り当て、同期、コード生成まで一貫して引き継ぐ;#3776 は AutoSchedule の起動間隔に単一タスク境界、リソースグループ境界、自己依存境界の三つの必要条件を追加し、三者の最大値で二分探索を初期化、昇騰 950PR に CANN 9.3.0 週次更新版を組み合わせた全サンプルがコンパイル通過し、スケジューリング結果は不変、Z3 求解フェーズの累計所要時間が 10.04% 削減。

  • リグレッション状態:月曜日の調査期間における日次テストは 2970 項目すべて通過、11-01 に workflow 失敗が 1 回(#1976);水曜と木曜の日次テスト記録は 3030 項目すべて通過、通過率 100%(#1980、#1988)。テストケース規模は一週間で 2970 から 3030 に増加しており、カバレッジが依然拡大中であることを示すが、#1991/#1992 と #3762 の二つの問題は人手による報告で初めて発見されたものであり、「通過率」を正確性の代替指標と見なすことはできないことを示している。

海光(Hygon):自社開発ブランチから公共レビューへ

  • メインブランチ:#31 は 10-30 にマージ、動的 MLS 境界推論と MatrixStore の stride 処理を改善、MMAC と blockscaled GEMM の clear_accum の下沉を修正、持続的 WASP パスで M パーティションコンシューマと独立した A/B フェンスに切り替え、slot と phase のローテーションを保持;#32 は 11-02 10:46 にマージ、gfx948 を gfx946 互換として扱い、既存の MLS atoms、MMAC、barrier、コンパイルオプションを再利用、目的は gather load を導入する前に既存の builtins を新型号で動作させること。これは典型的な二段階アプローチであり、まず新アーキテクチャ番号を既存パスに組み込んでテスト可能なベースラインを取得し、次に本当に新しい命令が必要な部分を補完する。
  • 機能ブランチ:feat/hcu-scaled-cast は 11-04 の 8 時間以内に 4 コミットを記録、それぞれ E8M0 スケーリングの FP8/BF8/FP4 変換(スカラーとパックの 2 種類の lowering、新規 223 行の HCU scaled_cast.h と lower_scaled_cast.cc、メイン変換ファイルの変更 650 行)、bank-word ジオメトリに基づく plain-copy の swizzle 導出と狭幅 b16 tile 向けの明示的逆変換の提供(12 ファイル +1536/-112、新規テストファイル 5 件)、混合 LDS wrap の「最大 2 経路」証明、混合 wrap チェックを物理読み取り幅に変更;11-05 にさらに 3 コミット(デフォルト swizzle を 64 bank に適応、MN スライス消費側で swizzle 計画を復元、K パーティションのレイアウト認識読み取り)。
  • レビュー状態:11-05 09:08 に #33(scaled cast、すなわちブランチ本体)を起票、19:52 に #34(reduction の XOR shuffle を DPP と lane swap で最適化、加えて「レイアウト推論は既知制約を先に伝播してから自由候補を割り当てる」修正)を起票。機能ブランチに 4 コミットを蓄積し PR 未起票の状態から 2 つの正式 PR へと至ったのは、国産バックエンドで初めてこのパスを完走した事例であり、テーマは共有メモリの bank 衝突とフェーズアラインメントに集中しており、当該バックエンドが「コンパイル可能」から「マイクロアーキテクチャに基づくチューニング」段階へ移行したことを示している。
  • 他に審査中の #19(タイトルは Distributed のみ、説明は空、11-03 11:06 更新)があり、本レポートではその実装詳細について推測しない。

摩尔線程(MUSA)と沐曦(MetaX):一方はドキュメント保守を再開、一方は依然アップストリーム同期のみ

  • 摩尔線程:11-05 19:46 にメインブランチへ docs: update README installation and operator development references の 1 コミットを記録、本リポジトリでは 09-30 以来の初コミット;同時期に 6 件の遺留 PR(依存アップグレード、TileOPs サブモジュール、MUSA 統合の証拠基準ドキュメント)のタイムスタンプが一括更新されたが、内容は依然 04 月から 08 月の古い提案である。判断は保守的にすべき:ドキュメントコミットはリポジトリに保守者がいることを示すが、バックエンド開発の再開を証明するには不十分であり、真のシグナルは後続にオペレータや lowering のコミットが現れるかどうかである。
  • 沐曦:デイリーレポートでは 5 つの調査期間すべてで「プッシュなし」と記録されていたが、168 時間の調査期間の再確認でメインブランチに今週確かにコミットがあることを確認、ただし逐一見るとすべてアップストリーム tilelang の同期(#3760、#3752、#3759、#3802、#3714、#3784、#3772、#3534、#3761、#3689、#3546、#3736、#3758、#3748、#3678、#3734、#3702、#3703 など)であり、作者はすべてアップストリームのメンテナで、沐曦側の独自コミットは見られない。したがって正確な表現は:メインブランチはアップストリーム同期中、独自開発は依然として静穏期にあり、すでに連続 3 番目の調査期間となる。dev ブランチの最新は依然 10-16 で停止している。

mlir-ascend と TLIR:セマンティクスを一層ずつ方言へ落とし込む

  • 構造整理:#206 は 10-31 にマージされ、CI 内で内部キャッシュサービスを指していた内容を削除した。再現性のクリーンアップに相当する;#209 は 11-04 16:35 にマージされ、BishengIR ツールチェーンの適応コードを contrib に移し、「上流ツールチェーンの進化に伴い頻繁に変更が必要となるグルーレイヤー」をコアの lowering から分離し、今後の BishengIR アップグレード時の変更範囲を contrib 内に収束させた。
  • Agent ワークフロー:#208 は harness ブランチにマージされ(メインラインではなく、引用時に「メインラインにマージ」と書くべきではない)、2715 行、66 ファイルを新規追加し、オペレータ harness に最初の Codex ワークフローを追加した:プロジェクト初期化時に反復質問を用い、既存のオペレータリポジトリを接続し、永続化コンテキストから開発を復元し、チューニングまたは分析を行う;3 つのスキルは .workspace/ 配下の編集可能な Markdown 記録を共有し、タスク引き継ぎでは「現在の候補」と「検証済みの最良候補」を区別し、失敗およびロールバックした実験を保持する。著者は今回実際の NPU 実測は行っていないと明言しており、NPU 実行適応と npm 配布は今後の作業として挙げられている。
  • 方言オペレータ:#203 は TLIR に tl.sub/mul/div/max/min を追加し、TL から HIVM への lowering を補完した;#211 は tl.shl を追加し、#212 は tl.shr を追加し、3 件は 11-05 の 01:42、19:44、20:58 に順次進められた。このラインはペースは遅いが方向は単一である:TileLang のセマンティクスを一層ずつ MLIR 方言に落とし込んでいく。

4. エコシステムと採用側

TileOPs:0.1.0 以降の第 3 週、チューニングからセレクタエンジニアリングへ

  • マージペース:月曜日のウィンドウは15件(うち13件は11-01)、火曜日7件、水曜日1件、木曜日7件、金曜日9件、週合計39件。火曜日のバッチにはテスト基盤(#2816 の tileops_test ターゲット、1281行、59ファイルで、tests/ops のスモークテストが CUDA 上に必須でなくなる)と CI 契約(#2818)という非性能変更が2種類現れ、オペレータライブラリがマルチバックエンドと保守性に向けた構造的準備を始めたことを示している。木曜日と金曜日は純粋な性能改善に戻った。
  • アテンション:#2808 は GQADenseFwdOp と GQADenseWithLSEFwdOp のプログラム選択を1つのルール AttentionCall.dense_decode_program にまとめた。入力は query と key の長さ、バッチ、ヘッドグループ、head dim、RoPE、SM 数で、2~64 個の query 位置は warp 専用 decode を経由する。利用可能な最速プログラムに対する選択後悔(幾何平均と最大値)は 1.818 と 25.2 から 1.017 と 1.81 に低下し、メインライン上の「score scale に log2(e) を掛けた後 FP32 で消失し NaN になる」問題と「D256 の複数位置呼び出しが SM90 で共有メモリ不足により起動失敗する」問題という2つの実バグを修正した。#2817 は warp 専用 decode の softmax を QK、PV とオーバーラップさせ(確率はレジスタに保持、K と V は独立 barrier、共有3スロットリングバッファ)、H200 上の (16, 8, 4096, 64, 8, 128) で float16 は 0.074610 ミリ秒対 fa3 の 0.077490(1.039倍、3.654 TB/s)、bfloat16 は 0.074260 対 0.077220(1.040倍、3.671 TB/s)。#2807 はページド MHA decode を「1ブロックの共有メモリに収まる」各 head dim まで拡張し、#2797 は GQAPagedFwdOp から旧ページド decode kernel を削除した。
  • MoE と GEMM:#2815 は tight レイアウトの MoE 入力をトークンごとに収集し、同じ hidden 行をその複数ルート間で再利用させる。(4096, 256, 8, 7168, 2048) 上で bfloat16 は 7.4417 ミリ秒対 flashinfer-cutlass の 7.6982(1.034倍、387.844 TFLOPS)、float16 は 7.5040 対 7.9869(1.064倍、384.624 TFLOPS)。#2805 は SM の少なくとも4分の3を埋める部分一波 W4A8 prefill グリッドを stream-K ではなく1 CTA 1 tile に変更し、llama-8b-up-b33 は 23.52 マイクロ秒から 17.98(vLLM 20.42)に短縮、リスト外 M 33~128 の44組の N/K スキャンでスケジュール形状変更は5%~36%高速化、中央値18%、遅くなったものは皆無だった。#2834 は小 M の NT GEMM のために一波を埋められる幅のネイティブ幅タイルを作り、decode-down(128×7168×2048)はメインライン上で132個の SM のうち112個しか使っていなかったが、変更後は最速比較項目を3%リードし、gate-up-b32 はメインラインより1.4%高速化した。
  • 要素ごと演算とリダクション:#2814 は bfloat16 の Sigmoid と16ビット SiLUAndMul を4分の1スケールの逆数に変更し、範囲チェック付き除算を回避した。#2813 は8バイトオペランドを長起動時に1スレッド1ベクトルで処理する(まず TileFoundry の roofline 分析でこの種のビット単位 OR オペレータがすべてメモリ律速であることを確認)。#2832 は lerp、where、clamp、masked_fill の5種類のオペレータが stride 0 のブロードキャストテンソルを先に出力形状に展開しなくなり、ブロードキャスト行の所要時間はメインラインの0.13~0.44倍に低下した。#2829 は「本体に最大2つの安価な命令、読み書き総量が L2 の4分の3以下」の単項要素ごと演算を256スレッドの一波に縮小し、ついでに ReLU が NaN 入力に対して0を返す古い問題を修正した。#2828 は行ブロック単位の ArgMax/ArgMin で「一波駐留」を前提に各スレッドが8つの16バイトベクトルを持つようにし、(2048, 4096) の bf16 ArgMin は 8.704 マイクロ秒から 7.393(FlagGems 8.543)に短縮した。#2835 はリダクション kernel の軸置換と出力形状をインスタンスごとに1回だけ計算するようにし、6組の FP32/BF16 VarMeanFwdOp 構成の完全なホットコール所要時間が4.33%~6.17%低下、生成される CUDA とホストソースコードは完全に一致——節約されたのは呼び出し準備コストであり、GPU 時間ではない。
  • 精度という種類の欠陥が正面から修正された:#2811 は行から構成される各 fp32 量に不変条件「行の広がり度のみに制約され、オフセット量には制約されない」を追加し、実装上は行内のある要素を基準として減算する(warp 主導のカーネルは warp 先頭要素、シャードカーネルは行の先頭要素)ことで、1e36 付近のオーバーフローを解消した。#2821 は分散統合を空のアキュムレータではなく最初の実部分量から初期化するようにし、定数入力の分散が NaN にならなくなった(tests/ops/test_reduce.py は155から158に増加)。これ

「空アキュムレータ + 大オフセット」のような数値的な落とし穴は、保持軸の前後で同時にリダクションする形状でのみ露呈し、通常の軸全体リダクションでは検出できない。

  • 計測口径が製品特性になる:#2830 は roofline を「今回の呼び出しで実際にどれだけ読んだか」で書き直した——GQA の稠密、varlen、ページング前向きは「query が見ることのできる key 行」に対してのみ K/V を計上し、key が実際に到達する截断位置に対して RoPE 表を計上する。GroupedGEMM は新しい grouped_gemm_roofline により batch_sizes に従ってある行のグループを価格付けする。修正後、物理的にあり得ない三行は 7.48、7.50、5.30 TB/s から 0.62、0.62、2.65 TB/s になった。#2822 は bench_kernel の L2 フラッシュを「int8 バッファに対する zero_()」から「プライベートな int32 バッファを読む、サイズは報告 L2 の二倍」に変更した。旧 flush は約 120 MB のダーティ行を残し、その書き戻しが次の計時呼び出しに落ち込むためである(2048×11008 f32 の Sum が 27.81 から 21.38 マイクロ秒に低下)。
  • セレクタとコストモデル:#2825 は W4A16 の _select_config を _legal_configs 上で _config_cost の argmin を取る形に変更し、tile 優先・最宽 K の同点ウィンドウと joint_switch_margin の三つの例外規則を削除し、stage_fill 項を新設した。#2826 は稠密セレクタの校正契約をドキュメントストリングに書き込んだ(スキャン、フィット、チェックの三步、ソフトアニーリングとホールドアウト検証を含む)。#2836 は coop2_swap の三階段タイルに個別のフィット因子(1.36)を与え、regret は幾何平均で 0.89% から 0.06% に、最大で 5.04% から 0.99% に低下した。
  • リリース形態とテスト口径:#2816 の tileops_test はスモークテストを CUDA から切り離し、著者は同時に境界を明確にした——共有参照との比較はバックエンド非依存の契約と移植性を証明できるのみで、実バックエンドの数値的正しさは証明できず、その線は依然として #2611 に掛かっている。#2827 は off-CUDA スイートを夜間に移動した。#2833 は runner をアップグレードした。

TileFoundry:同期分析、訓練セマンティクスと命名ガバナンス

  • 同期:#283 の分析層と #285 の lowering 実行端(要聞第五項参照)に加え、#286 が 11-06 13:40 にマージされ、各オペランド効果に対して完了条件を宣言し、WGMMA の fence、commit、wait をそれを必要とする HIR ノードに紐付け、lowering が命令の発射点でこれらのステップを消費し、後続スコープが集団待機を収められない場合は発射スコープ内で完了させる。同時にレジスタ Reshard の別名を修正し、要素がコピーを保持する各スレッド上で同じレジスタに留まるようにした。著者は二つのリスク声明を保持している:行列命令エミッタはまだ接続されておらず、このコミットは実行の正しさを証明せず性能も与えない。TMA の完了セマンティクスは別途扱う。11-06 20:12 に別途 #288 を開き、「完了済みの部分リダクション」の分析における計價を扱う。
  • 訓練セマンティクス:#284 が 11-05 16:36 にマージされた。以前の評価器は reshard をグローバルテンソルとして扱い、異なる参加者の寄与の差異を見落としていた。本コミットは参加者ごとにローカル値を求め、GPU 座標、呼び出しとループ、gather/scatter および軸沿いの Partial(sum) リダクションを行い、Reshard の別名証明を帰属に敏感にした。検収:コアスイート 775 項、バックエンドスイート 778 項(二卡 sm90 上で 13 項スキップ)、十組の HIR 配置が CPU 上で通過し、二種の駐留戦略が四卡 sm90 上で実際の 2×2 PyTorch HSDP と二更新ステップを実行し、ローカル勾配、パラメータ、モメンタムがすべて一致した。
  • スケジューリングとリリース:#280 と #282 は「各呼び出しで何を発射するか」「どの mesh が発射するか」「ウィンドウをどう切るか」の決定を analysis/emission.py 一箇所に集約した。検収はバイト単位の対照である——83 個のスケジューリングフィクスチャのうち 75 個の finalize 出力がメインラインと完全に一致し、残り 8 個はメインライン上でも同様に失敗する。サイト倉は 10-30 に 0.3 版スケジューリング記事《Scheduling, Designed for Agents》をリリースし、「agent 向けスケジューリング」をコードとリリースノートから公開叙述へ移した。
  • 命名ガバナンス:#287 はドキュメント向け issue であり、コードベースに「一つの概念に五つの名前」があると指摘する——mesh の一メンバーが placement 133 箇所、program 110 箇所、instance 92 箇所、participant 44 箇所と呼ばれ、さらに unit という一語が空型、ハードウェア計算ユニット、mesh メンバーを同時に指し、Python と CUDA ランタイムで local_layout のシグネチャも一致していない。この種の issue は通常すぐにコードを変えることはないが、コンパイラプロジェクトが「書ける」から「長期間人に読める」へ進むための必经ステップである。

TileSight:性能モデルが論文からリポジトリへ

  • tile-ai/TileSight(24 スター)は、07-24 の arXiv 論文「TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters」のオープンソース実装である。タイル形状・データ移動・ソフトウェアパイプライン・ハードウェアプロファイルから出発し、CPU 上でオペレータのコスト、キャッシュ流量、実行時間を見積もる。11-05 にリポジトリ初の PR(#1)が出現した。H20 は Hopper/GH100 であり、requires_arch=("wgmma",) 付きのプランは、アーキテクチャプロファイルが WGMMA サポートを宣言していないために従来は UnsupportedError を投げていた。本コミットは support_wgmma=True を設定し、かつ明示的に support_utcmma=False を維持するため、こうしたプランは受け入れられ、一方で tcgen05 と Tensor Memory の要求は引き続き拒否される。検証は CPU 上で 847 項目すべて通過であり、作者は明示的に「GPU 性能の検証ではない」と宣言している。
  • この路線の意義は、TileOPs の roofline 見積もり、TileFoundry の解析レイヤーと合わせて三角形を形成することにある。一つは解析モデルを提供し、一つは行単位の見積もりを提供し、一つは証明可能な同期プランを提供する。三者はいずれも今週、実質的な動きがあった。

採用側:TileKernels と FlashQLA

  • 二つの採用側リポジトリは週を通じてコミットがなかった。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、隠れ次元 1280 上でバリア占有予算の 25% を失っていた。これは、採用側のカーネルがすでにコンパイラ側の回帰プローブになっていることを示す。TileKernels がコードをプッシュしないこの数週間、そのカーネルは依然としてメインラインの変更を通じて受動的に「テストされ続けて」いる。エコシステムを追跡する者にとって、これは注目に値する逆方向のシグナルである。

コミュニティブランチとサードパーティリポジトリ

  • RISC-V 研究チェーンの形成:コミュニティ研究者 zishu0131-png が 10-31 に同日三個のリポジトリを整理した。tilelang-riscv は総合エントリ(研究目標、二段のコンパイルチェーンとバージョン組み合わせを説明)、numba2tilelang は Numba CUDA Python を静的解析して TileLang ソースコードに変換、tilelang2rvv はサポートされる実在の tvm.tirx.PrimFunc を RVV intrinsic C に変換し、さらに Clang で RV64GCV ELF にコンパイルして QEMU virt 上でシングル hart で実行する。作者は境界を明確に記している。TileLang 公式登録のバックエンドではなく、threadIdx に依存するアドレッシング、hart 間同期、未サポートの IR は一律拒否され、block は順次実行される。研究環境は TileLang 0.1.15、Clang 23.1.3、QEMU 11.1.0 である。168 時間の調査期間における四情報源の巡回では、今週新たな変化はなかった。
  • PTO-ISA/tilelangGym の独立リポジトリ化:この感度・回帰テストセットは PTO-ISA/tilelang の examples/ascend/t_parallel_st_suite から分割された。理由は、テストセットが TileLang と PTOAS のバージョンを独立して固定できるようにするためである。現在の三つの主要パスは、昇騰 SoftLib SimtVF(target=ascend、cython、主要回帰ゲート)、VectorizeParallelToPTO と VL の vload/vstore、そして benches 配下の @pto.jit で書かれた手書き SIMD DSL 対照実装である。各テストケースディレクトリには config.toml と VERDICT.md が付属する。
  • サードパーティの小規模リポジトリ:sablin39/tilelang-cuda-skills(149 スター、TileLang の記述と CUDA ツールチェーンでのデバッグのためのスキルセットを提供)、xwhzz/tilelang-hexagon(2 スター、TileLang を Qualcomm Hexagon DSP に持ち込む)、Index4NekoMeteor/Tilelang-AttnRes(1 スター、TileLang で Kimi K3 のアテンション残差構造を実装)、Tianchen-ai/tilelang-weft-bench(0 スター、CPU 側オペレータベンチマークの統一エンジン)、flytiger-eco/tilelang-wrapper-for-sail(0 スター、説明なし)。最初の三者は方向性が興味深い。一つはツールチェーン側の開発体験、二つは新しいターゲットアーキテクチャと先端モデル構造への移植である。後二者はごく初期段階であり、まず観察リストに記録し、技術的進展として書くべきではない。

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

  • 不具合報告はインストール形態に集中:水曜1日で7件の新規issueのうち、4件は同じ外部報告者によるもので、pipパッケージを使い、WSLで実行し、コンシューマー向けsm_89カード(RTX 4060 Laptop)を使用しており、すべて再現スクリプトと完全な出力を添付していた。技術的に最も高度な#3775は、本番環境における「pipパッケージ + デュアルH200 NVLink + 複数CUDAコンテキスト + nvrtcバックエンド + CUDA graphキャプチャ」の組み合わせで、報告者は3回の実行の対照条件をすべて書き出していた——3回バイト単位で一致する設定で再現したところ、最初の失敗点が3つ異なり、/proc/environとクリーンな対照との違いはTILELANG_EXECUTION_BACKEND=nvrtcという1つの環境変数だけであった。これらはコンパイラの深層的な欠陥ではないが、新規ユーザーが初めて試す際の体験を決定づけるものである。
  • 昇騰側に「報告者すなわちレビュアー」の閉ループが出現:#1991/#1992の報告者は#1897のレビュアーであり、#1897の修正方法はcount側のセマンティクス(HasModifiedBufferがworkspaceスロットをスキップ)のみをカバーしており、workspace側は未処理のままである——同じレイアウトロジックの別の側面が2週間後に指摘された。これは、昇騰バックエンドのレイアウト不変条件がまだ検査可能な書面契約として形成されていないことを示している。
  • ドキュメントとローカライゼーションの链路は正常:tilelang.github.ioは週を通じて複数のUpdate docsが自動同期され、TileOPs.github.ioは1回プッシュされた。LSPリポジトリはCUDAデバッグを新しい行マッピングに整列させ(71e762e0)、VS Code拡張リポジトリはドキュメントを更新し成果物をvsixに改名した。これは11-02の「CUDAデバッグをTileLang行マッピングに整列」に続くもので——前号の主線でマージされた行単位のアンカリングが生成コードの行番号ソースを変更したため、拡張側のデバッグ設定もそれに合わせて変更する必要がある。
  • チュートリアルと知識資産:今週は新しいチュートリアルリポジトリは登場せず、コミュニティの投入は既存リポジトリのメンテナンスと上述した2つのRISC-V方向の小規模リポジトリに集中した。PTO-ISA/tilelangGymの独立リポジトリ化は、テスト資産が初めて単独でバージョン管理されたことを意味する。
  • ニュース側は週を通じて実質的ヒットゼロ:5つの収集期間のGoogle News中英文クエリ(精密語、GPU kernel、中国語オペレータとコンパイラ、昇騰と沐曦/海光/摩尔線程をカバー)において、最新の関連報道は依然として10-21のSemiAnalysis「The Huawei-DeepSeek Symbiosis」であり、残りは09-30の一波「DeepSeek オープンソース昇騰版TileLang、DeepGEMM、DeepEP」の報道に集中しており、すべて本周期より前のものである。Hacker Newsでtilelangを24時間検索した結果は0件。arXivの最新ヒットは11-03の「Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving」(TileLangを引用した作業、期間外)と07-24のTileSight本体である。したがって報告本体はすべてコード側からのものであり、これはすでに複数期連続の安定した形態である。
  • 産業側の口径の衝突は特筆に値する:水曜の期間において、WccftechはDeepSeekが昇騰950DTで構成されるAtlas 950 SuperPoD上でゼロからV5を訓練すると報じたが、同時期のThe Informationの報道は「早くても2026年末に950DTでの事前訓練を開始する」であり、Bloombergと金融時報の初期報道は950DTの16万枚の受注が主に推論用であり、以前に昇騰での訓練が挫折した後にNvidiaへ回帰したと伝えている。3つの説は「訓練に使用するか否か、そしていつ使用するか」について一致しておらず、いずれもHuaweiまたはDeepSeekの公式確認を得ていない。本レポートは背景記録としてのみ扱い、これに基づいて昇騰側の投入の実装ペースを推測するものではない。

6. トレンド観察

第一に、TIRx移行はすでに主線の事実上の最優先事項であり、それが「減法」の形で現れ始めている。 一週間のうちに3つの証拠がある:月曜に移行のためにThreadSync正確性改善の一式を回退し(#3733、同時にTVM fork側の最適化提案#89をクローズ)、火曜に3384行の遺留パーサー分岐を削除し(#3736)、金曜に完全な移行草案をコミットした(#3803、2つの上流コミットをpinし、パッチを注入しないことを明示)。その間にZ3リソースをハードコードから設定可能への変更(#3752とtvm #91)も挟まっている。これらの動作に共通するのは、複雑さをクリティカルパスから撤退させることである:一時的に保守的な実装に戻しても、vendored依存の内部構造を言語層に生やし続けることはしない。代償として、主線は移行期間中に「既知の非最適」パスを一批抱えることになり、收益としては基盤の世代交代が完了した後にすべての下流バックエンドが安定したIR境界に対して適配できることである。

第2に、正確性リスクは形態を変えた:「コンパイルが通らない」から「コンパイルは通り、実行時にエラーも出ず、結果が誤る」へ、しかもバージョンを跨いで潜伏する。 今週は少なくとも7件がこの族に属する:#3808 は 0.1.12 から 0.1.15 までの4バージョンにわたり __syncthreads() が1つ欠けており、#3800 は Blackwell の tensor memory パス上で MMA が上書きされた operand を読むようにし、#3721 は sm_100a 上で GEMM の積全体を黙って破棄し、。#3775 は本番環境で CUDA graph のキャプチャを破壊し、。#3749 は依存関係分析から物理ロード範囲を見えなくし、。#1991/#1992 は昇騰の threads=2 下で sin/cos が 97.1% の要素を誤計算し、。#3774 と #3773 はコンシューマー向け카ードで誤った結果を出す。この族はユーザーへの被害が最大である。なぜなら「コンパイラは正しいか」という前提を、検証を要する仮説に変えてしまうからだ。今週の対応方向は、正確性を棄却可能な証明にすることだ——TileFoundry の #283/#285 はまず証明してから発射し、発射後に逆検証する。#3786 と #3787 は黙って誤ったコードを生成するのを、コンパイル時エラーまたは明示的な拒否に変える——単体テストの数を増やし続けることではない。

第3に、昇騰の役割は「適応ターゲット」から「初回ターゲット」へ変わり、エンジニアリングが2本の脚で同時に進んでいる。 1本はアーキテクチャだ。FlashMLA のような DeepSeek 自身のデコードパス上のコアカーネルが、昇騰 950 を直接ターゲットとしてメインラインに書かれており、使っているのは T.Stage、annotate_buffer_addresses といった明示的スケジューリングのために用意された低レイヤ機能だ。コンパイラ側の3箇所の配套変更と合わせ、昇騰パスは「動く」から「スケジューリング意図どおりに動く」へ転換しつつある。もう1本は骨格だ。#1975 は 2887 ファイル、60万行の追加で TileLang-Ascend 全体をメインラインのマルチバックエンドフレームワークに移し、arch="dav-2201" でバックエンドを振り分ける。mlir-ascend の #208 は Codex の初期化、接続、復旧の3つのワークフローとタスクハンドオーバーを .workspace/ に書き込む。「移行リポジトリ」と「エージェントによる移行リポジトリ」を同じ週に見ると、昇騰側のボトルネックが「オペレータ不足」から「一貫したエンジニアリング骨格と再利用可能な移行プロセスの不足」へ移ったことが分かる。

第4に、オペレータライブラリの競争点は完全にセレクタ品質と準備コストへ移り、性能差の残された空間は「どう計算するか」にはない。 金曜の TileOPs の9件のマージのうち、実際に GPU kernel を変えたのはブロードキャスト就地読みとネイティブ幅タイルの2件だけだ。reduction レイアウトキャッシュが節約するのは呼び出し準備(4.33% から 6.17%、生成ソースコードは不変)、3段階タイルファクタが変えるのは1つの定数(regret が 0.89% から 0.06% へ)だ。さらに roofline 価格計算の書き換えと L2 フラッシュを読み取り専用化を加えると、明確な路線が見える。まず各形状の合法な設定を列挙し、統一されたコストモデルで順位付けし、キャリブレーション過程を再現可能な契約として書く(#2826 は「スキャン、フィッティング、チェック」の3ステップをドキュメント文字列に書いている)。この道を正しく進めば、性能は個別カーネルの手作業チューニングに依存しなくなる。誤れば「ある形状は特に速く、他の形状は運任せ」に退化する。

第5に、国産バックエンドは「一度に大きなバックエンドをマージ」から「レビュー可能な最小単位を公共リポジトリへ push」へ転換している。 海光は自研ブランチから #33 と #34 の2つの正式 PR へ進み、摩尔線程は1本の README で2ヶ月の静寂を終わらせ、mlir-ascend は3つの小さな PR で tl.sub/mul、tl.shl、tl.shr を1つずつ TLIR 方言へ移し、昇騰はコア比率、例外ダンプ、性能スキルといったエンジニアリング詳細を項目ごとに補っていく。3本の線の共通点は、もはや大きなマージを待たないことだ。これは上流にとって良いことである。レビュアーは配置と lowering の設計に早い段階で意見を出せる。3万行の一括マージの時に初めて見るのではない。下流にも代償がある——メインラインはこうした期間、「既知のいくつかが動かない」状態にあり、より強力なリリース前回帰規律が必要になる。

最後のリスク面:ベーススイッチはすべてのグルーレイヤのコストを一度に露呈させる。 もし #3803 が進めば、3rdparty/tvm の vendored 構造は2つの pin された上流コミットに置き換えられ、その時点で TVM 内部 API に依存するすべてのバックエンド(昇騰、海光、mlir-ascend のグルーレイヤ)が追随して変更しなければならない。草案には「依存パッチを注入しない」と明記されており、これは約束でもあり制約でもある——今週 mlir-ascend が BishengIR アダプタを contrib へ移し、昇騰が性能スキルをチップ世代別に分岐させたことは、いずれもこうしたベーススイッチのために用意された交換可能面と見なせる。この線が健全かどうかは、移行草案のマージ後2週間以内に、これらのバックエンドがコア lowering に手を入れずに適応を完了できるかで判断する。


付録:素材と検証説明

情報源チェック表

情報源 検証結果
GitHub tile-ai 組織のリポジトリ別(29 リポジトリ) 五つの収集ウィンドウで順に 10、9、9、8、10 のリポジトリにプッシュがあり、168 時間ウィンドウの再検証では 14 のリポジトリにプッシュがあった
マージ実績検索(is:pr is:merged) 五つのウィンドウで 35、18、8、12、18 件、合計 91 件。最初のウィンドウは 72 時間(10-30 07:00 から 11-02 07:00)
審査中実績検索(is:pr is:open updated) 五つのウィンドウで 82、44、48、42、57 件、合計 273 件
新規 issue 照合 五つのウィンドウで 11、8、7、4、5 件、合計 35 件
メインブランチのコミット単位検証 tilelang 7、5、5、3、5 件;TileOPs 15、7、1、7、9 件;TileFoundry 3、1、1、1、1 件;tilelang-ascend 5、4、0、1、2 件;tilelang-mlir-ascend 1、0、1、2、0 件;tilelang-hygon メインブランチ 2、1、0、0、0 件;tilelang-lsp と VS Code 拡張機能が各 1 件;tilelang-musa 1 件
リリースページとタグ tilelang 最新 v0.1.15(09-30)、週全体で新バージョンなし;TileOPs 最新 v0.1.0(10-27);TileFoundry 最新 0.0.4rc1(10-25);tilelang-ascend v0.1.2.000(09-09);tilelang-metax v0.1.15(10-16)
採用元 TileKernels は 10-14 で停止;FlashQLA は 10-30 で停止、最新リリースは依然 09-30 の v0.1.3
TileLang×RISC-V 四情報源巡視 168 時間ウィンドウで変化なし
サードパーティとエコシステムリポジトリ 研究チェーン三リポジトリ、tilelangGym、cuda-skills、hexagon、AttnRes、weft-bench、wrapper-for-sail、TileSight のいずれもプッシュまたは新規作成あり
Google News(中英四から六組) 五つのウィンドウいずれもウィンドウ内に実質的なヒットなし、最新の関連報道は 10-21 の SemiAnalysis の記事
Hacker News tilelang で 24 時間検索は 0 件、ヒットしたのは主題と無関係な Show HN
arXiv 最新ヒットは 11-03 の KV 圧縮論文(TileLang を引用、ウィンドウ外)と 07-24 の TileSight 本体

主な引用

com/tile-ai/tilelang/pull/3783>

説明

  • 本レポートの事実、バージョン番号、日付はすべて調査期間内の5回分のデイリーレポート原文およびその一次情報源から取得しており、推測による補足は行っていない。5つの収集期間は連続しており(10-30 07:00 から 11-06 07:00 まで合計168時間)、前回のウィークリーレポートとの重複はないため、5期間のマージ数、レビュー中件数、新規 issue 数はそのまま合算した。
  • マージ数、レビュー中件数、issue 数は GitHub 検索の実測出力によるものであり、メインブランチのコミット数はリポジトリごとの atom 照合の結果である。両者は集計基準が異なる(1回のマージが複数のコミットに対応する場合があり、非デフォルトブランチのコミットはマージ数に含まれない)ため、付録では2種類を別々に記載し、合算はしていない。
  • TileOPs の「15件」は月曜日の72時間枠内に現れたもので、うち13件が 11-01 に集中しているため、引用する際に「月曜単日15件」と書くのは適切ではない。
  • #1991 と #1992 はタイトルと本文が一致しているため、「同題2件」として記録し、1件に統合していない。海光の機能ブランチの7件と摩尔線程のドキュメントコミットはブランチレベルの動向として表現しており、メインラインにマージ済みであることを意味しない。
  • 沐曦は5期間すべてで「プッシュなし」と記録されていたが、168時間の再確認でメインブランチにコミットがあることが判明した。ただし1件ずつ照合したところ、すべて上流 tilelang の同期(作者は上流メンテナ)であったため、「メインブランチは上流同期中、独自開発は依然として静穏期」と表現した。TileSight と TileFoundry の WGMMA 変更はいずれも作者自身が能力メタデータまたは同期計画に限定すると述べており、レポートはその限定に従っている。