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

情報源：GitHub tile-ai 組織のリポジトリ別調査（29 リポジトリ、期間内プッシュ 9 件）＋マージ済み・審査中ソースの実データ検索（マージ済み 18 件、審査中更新 44 件）＋新規 issue の全件確認（8 件）＋メインブランチのコミット別確認（tilelang 5、TileOPs 7、TileFoundry 1、tilelang-ascend 4、tilelang-hygon 1、tilelang-lsp 1、ドキュメントサイト 5）＋リリースページとタグの確認（5 リポジトリ）＋採用側 2 リポジトリ＋Google News の中国語・英語検索（プロキシ経由）＋Hacker News＋arXiv。

本期インデックス

  • 本日の重点：昇騰 950 の永続化 FlashMLA カーネルの実装（#3678、1997 行追加、42 コミット）；ベクトル化計画における実回帰 1 件の修正（#3734）；#line 指令を1行ずつ再アンカリングするよう変更（#3703）；TileOPs は1日で性能・テスト基盤のコミットを7件マージ；TileFoundry は共有メモリソートを証明可能な解析層にし（#283）、main の lowering が sm_89 で7プログラム中4件誤りと測定；24時間で18件マージ、審査中更新44件、新規 issue 8件
    1. 中核プロジェクトの進展
      • 1.1 メインライン：昇騰 950 の永続化 FlashMLA カーネル、BF16 prefill とコンパクト FP8 decode のデュアルパス（#3678）（11-02）
      • 1.2 メインライン：cast の分離がスキップされた後に残存する 512 ビットベクトル、SM100 で ViT フォワードのコンパイル失敗（#3734）（11-02）
      • 1.3 メインライン：#line 指令を1行ずつ再アンカリングするよう変更、デバッガマッピングはオプションで密モード（#3703）（11-02）
      • 1.4 メインライン：アテンション例がデバイス共有メモリ上限に基づいてタイルを推定、RTX PRO 6000 で2成功7失敗から9件全成功に（#3702）（11-02）
      • 1.5 メインライン：CUDA float32x8 ビット変換を内蔵 bit cast に変更、SASS 命令が19条から18条に（#3748）（11-02）
    1. マルチバックエンド対応
      • 2.1 昇騰：WY Fast フォワードオペレータを新規追加、wy_fast_bwd_split 逆方向セグメント移行（#1827 / #1828）（11-02）
      • 2.2 昇騰：T.tile.fill と T.tile.clear のドキュメントファクトチェックをマージ（#1634 / #1628）（11-02）
      • 2.3 昇騰審査中：cython wrapper がシンボル形状解析をキャッシュ、ホストオーバーヘッドのラインを引き継ぐ（#1967）（11-02）
      • 2.4 昇騰審査中：z3 版同期をメインラインから移入、リポジトリ全体のメインライン同期を推進（#1965 / #1975）（11-02）
      • 2.5 昇騰審査中：リダクションオペレータの2件の欠陥、最初の Parallel 軸ブロードキャストをスカラー読みに変更（#1978 / #1968 / #1979）（11-02）
      • 2.6 海光：gfx948 を gfx946 互換として処理、まず既存命令をテスト可能に（#32）（11-02）
      • 2.7 摩尔線程、沐曦、mlir-ascend は調査期間内にプッシュなし（11-03）
    1. エコシステムと採用側
      • 3.1 TileOPs：1日で7件マージ、性能6件、テスト基盤1件（11-02）
      • 3.2 TileOPs：warp 専用 decode の softmax と QK/PV をオーバーラップ、478回のスキャン呼び出しの幾何平均 0.952 倍（#2817）（11-02）
      • 3.3 TileOPs：MoE がトークンごとに収集して置換を行単位に、H200 で 387.8 TFLOPS（#2815）（11-02）
      • 3.4 TileOPs：16ビット要素ごとに範囲チェック付き除算を削除、8バイトオペランドは1スレッド1ベクトル（#2814 / #2813）（11-02）
      • 3.5 TileOPs：Norm 統計量を行内要素の減算に変更、1e36 付近のオーバーフローを解消（#2811）（11-02）
      • 3.6 TileOPs：tileops_test ターゲットで tests/ops を CUDA なしで実行可能に（#2816）（11-02）
      • 3.7 TileFoundry：共有メモリソートを証明可能な解析層に、83件のフィクスチャ計画で52件拒否31件（#283）（11-02）
      • 3.8 TileFoundry 審査中：CUDA lowering が同期計画を発行し逆検証、main で7プログラム中4件誤り（#285）（11-02）
      • 3.9 採用側：FlashQLA、TileKernels は調査期間内にプッシュなし（11-03）
    1. コミュニティ、チュートリアル、イベント
      • 4.1 昇騰同期の正確性が集中報告面に：1日で8件の新規 issue、2名の報告者がそれぞれ体系的に提出（11-02）
      • 4.2 行跨ぎ・転置ビューの SIMD ロードで DMA 依存が失われる、根本原因は論理領域の再構築に（#3749）（11-02）
      • 4.3 ニュース側：Google News 中英文、Hacker News、arXiv の3チャネルとも調査期間内にヒットなし（11-03）
    1. トレンド観察
  • 付、素材と検証の説明

本日の重点

この日の主线はただ一つ：TileLang が DeepSeek の MLA デコードカーネルを昇騰 950 に移植したことであり、しかもコンパイラがスケジューリング可能な書き方で移植した。#3678 は一度に 1997 行、20 ファイル、42 コミットを追加し、BF16 prefill とコンパクト FP8 decode の二経路をカバーし、さらにオプションの FP8 と FP4 の追加 KV キャッシュ、attention sink、不正インデックス、query ごとの独立長を備える。実装方式は注目に値する：attention ループ、オンライン softmax、逆量子化、スケジューリングはすべて TileLang 内に記述され、T.Stage によってコンパイラにバッファバージョンと同期を管理させ、QK と PV は同一の L0B ストレージリングの二つの異なる layout ビューを使い、L1 の配置は T.annotate_buffer_addresses で直接指定する；prefill は一つのアキュムレータ割り当ての中で四つの互いに重ならない出力スライスに分け、正規化が一片終わるごとに一片を排出する。この書き方を成立させるため、コンパイラ側でも同時に三つの能力を補った——L0C から UB へのコピーで NZ 出力を保持すること、各 AIV ストリームが自身のパイプラインの前に選択されること、明示的にスケジューリングされたステージが物理発行フェーズで同一の場合にソース順を保つこと。対象形状は非常に狭く書かれている：ヘッド数 64、Dqk と Dv はともに 512、topk は 64 の倍数。

第二の線は実際のコンパイラリグレッションの修正である。#3734 が扱うのは SM100 上の ViT フォワード：Cannot convert type float32x16 to CUDA type。経路は長いが各ステップは明確——ベクトル化プランナが 256 ビットの BF16 ストアのために 16 レーンを選び、cast 制約を先送りし、DecoupleTypeCast が計算とメモリアクセスを分離することを期待する；一方 #3579 が直前に「メモリ読み取りのガード」をデカップリングから除外した（これらのガードを複製すると意味が変わる）ため、最終合法化のステップが依然として先送り戦略で幅を選び、残留 FP32 式が 512 ビットベクトルに膨らみ、CUDA バックエンドが表現できなくなった。修正は先送り制約を IR 属性ではなくプランナのパラメータにし、最後の合法化ステップのみ false を渡し、ガードの FP32→BF16 は 8 レーンに減り、デカップリングに成功した BF16 コピーは依然 16 レーンを保つ；アドレス cast もデータベクトル型制約から外し、既存のアクセス解析で合法性を判定する。

第三の線はデバッグ体験である。#3703 は #line 指令を「各 TIR 文ごとに先頭行のみをアンカーする」からオプションの段階的発行に変更した：レベル 0 はオフ、1 は現状維持、2 以上は各行が生成する非空行ごとに再アンカーし、複数行に展開された文や末尾の大括弧も含む。実装は CUDA、HIP、CPU C、C host、Metal、Ascend が共用する C 系出力ストリームに置かれ、vendored の TVM には触れない；ソーススパンのない文と関数は明示的に <tilelang-generated> に切り替え、ユーザー原文はマクロ継続行や文字列リテラルを壊さないよう逐字保持する。新環境変数 TILELANG_EMIT_LINE_DIRECTIVES を Python 側のデフォルト値とし、キャッシュ検索の前に pass 設定へ折り込む。ローカルで 42 項通過、6 項スキップ、CUDA と HIP のケースは作者のマシンでスキップされ、cuda-gdb の実際の動作はまだ直接検証されていない——作者自身がこの点を明記している。同日 LSP リポジトリが一笔追随し、CUDA デバッグを新しい行マッピングに合わせた。

オペレータ側のリズムは止まっていないが、方向は変わった。TileOPs の単日 7 笔のマージのうち 6 笔が性能、1 笔がテスト基盤で、すべて同一作者による；各笔は依然として H200 実測表とチェックリスト回帰ゲートを添付している。最も価値があるのは #2816：テスト時のみ存在する tileops_test ターゲットを新設し、PrivateUse1 上で各オペレータの共有 workload リファレンスを算出し、tests/ops のスモークテストがもはや CUDA 上に立つ必要をなくす。1281 行、59 ファイルの変更で、作者は境界も同時に明確にした——共有リファレンスとの比較はバックエンド非依存の契約と可搬性を証明できるだけで、実際のバックエンドの数値正確性は証明できず、その線は依然 #2611 に掛かっている。

最後は TileFoundry の一本の暗線である。#283 は「共有メモリアクセスをいかに順序付けるべきか」を分析層にした：同期ファミリが lowering の発行順に関数を走査し、各共有メモリアクセスについて命令、発行 mesh、オペランド、バイト範囲を取り、競合アクセスの前に commit/wait とスコープ内フェンスを計画する。83 個のスケジューリングフィクスチャのうち 52 個を計画し、31 個を拒否する（24 個は TMA コピー、6 個は wgmma を使用し、両者の完了セマンティクスはまだ宣言されていない；1 個はスレッド境界をまたいで不等なループにフェンスを挿入する必要がある）。付随する #285 が審査中で、計画を CUDA lowering に変えて逆検証する：sm_89 単卡上で構築可能な 7 プログラムのうち 6 個がリファレンスと完全一致し、synccheck と memcheck はすべてクリーン、一方 main の方式で lower すると同じ 7 個のうち 4 個が結果誤り；#283 の記述で示された直接症状は sm_89 の E4M3 GEMM が最大 659 偏る。

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

1.1 主线：昇騰 950 の永続化 FlashMLA カーネル、BF16 prefill とコンパクト FP8 decode の二経路（#3678）

11-02 09:26 にマージされ、1997 行追加、17 行削除、20 ファイル、42 コミットで、本調査期間で最も大きな一笔となった。カーネルは昇騰 950 向けで、prefill は BF16、decode はコンパクト FP8 を使用し、オプションで FP8 と FP4 の追加 KV キャッシュを持ち、attention sink、不正インデックス、クエリごとの独立長をサポートする。アテンションループ、オンライン softmax、逆量子化、スケジューリングはすべて TileLang で表現し、T.Stage とコンパイラ管理のバッファバージョンおよび同期で駆動する。QK と PV は L0B ストレージリングの 2 つの異なる layout ビューを共有し、L1 の配置は T.annotate_buffer_addresses で指定する。prefill は既存のアキュムレータ割り当て内で 4 つの不相交出力スライスに分割し、独立した同期イベントを備え、正規化が 1 スライス完了するごとに 1 スライスを排出する。decode は行全体の出力コピーを使用する。コンパイラ側では 3 か所のサポートを追加する。L0C から UB へのコピーで NZ 出力を保持し、各 AIV ストリームをそのパイプラインの前に選択し、明示的スケジューリング段階が物理発射フェーズで打ち消された場合にソース順序を維持する。目標はヘッド数 64、Dqk と Dv がいずれも 512、topk が 64 の倍数であり、pair gather は既に記録されたゼロ burst コード生成の回避策を踏襲する。正確性とベンチマークのランナー、バッファ再利用とクエリ境界の回帰、および FlashMLA MIT 宣言付きの README を同梱する。

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

1.2 メインライン：cast 分離がスキップされた後に残存する 512 ビットベクトル、SM100 上で ViT フォワードのコンパイル失敗（#3734）

11-02 09:05 にマージ。e8225d7f5 上で、SM100 向け ViT フォワードがガード付きの FP32 から BF16 への出力ストア箇所で Cannot convert type float32x16 to CUDA type を報告する。プランナは 256 ビット BF16 ストアのために 16 レーンを選び cast 制約を延期し、DecoupleTypeCast が計算とメモリアクセスを分離することを期待する。#3579 がメモリを読むガードに対して分離をスキップさせた後、後続の LegalizeVectorizedLoop は同じ戦略で延期を続け、残存する FP32 式が 512 ビットベクトルとなり CUDA 側で表現できなくなった。修正は defer_cast_constraints をプランナの C++ パラメータ（IR 属性ではない）とし、呼び出し元が段階に応じて選択する。LayoutInference と LowerTileOp は true を維持し、最終合法化は一律 false を渡し、算出された幅を明示的な vectorize_hint として前方伝播し、リライタがデフォルト戦略で再計画するのを避ける。ガードの FP32 から BF16 へは 8 レーンに低下し、正常に分離された BF16 コピーは 16 レーンを維持し、ガードの安全チェックは保持される。アドレス cast とアクセスポインタ引数はデータベクトル型制約から除外され、連続アドレスはスカラベースアドレスとストライドに降格され、既存のアクセス解析で合法性を判定し、int64 インデックスの整列ベクトルアクセスと 4 レーンアトミックのユースケースには影響しない。

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

1.3 メインライン：#line 指令を行ごとの再アンカーに変更、デバッガマッピングはオプションの密集モード（#3703）

11-02 04:06 にマージ、682 行追加、11 ファイル。従来の #line 発行は 1 つの TIR ステートメントが生成する最初の行のみをアンカーし、後続行はコンパイラの論理行番号で進められるため、cuda-gdb は生成されたロード、演算、ストアを無関係な Python 行に帰属させていた。合成ステートメントと末尾の波括弧は進められたマッピングを継承する可能性があり、ソーススパンを持たない関数は前の関数のソースファイルを継承することもあった。本笔は tl.emit_line_directives を段階化する。0 または False で無効化しデフォルトを維持、1 または True でステートメントレベルを維持、2 以上で非空の生成行ごとに再アンカーする（複数行展開のステートメントと末尾の波括弧を含む）。実装は C 系出力ストリームにあり、CUDA、HIP、CPU C、C host、Metal、Ascend をカバーし、vendored TVM は変更しない。同時にステートメントと関数のソーススパンを記録し、ソースのない領域は明示的に <tilelang-generated> へ切り替え、ユーザ原文はマクロ継続行と文字列リテラルを壊さないよう逐語的に保持する。新たに TILELANG_EMIT_LINE_DIRECTIVES を Python 側のデフォルト値として追加し、キャッシュ検索前に pass 設定へ折り込み、明示的設定を優先する（デフォルト値が不正な場合を含む）。ローカルで Metal と LLVM ビルドを用いて検証した。行指令と compile-only の 2 グループ 42 項目が通過、6 項目がスキップ。指令を外した場合、ベクトル化加算と条件選択の 2 つのフィクスチャはレベル 0、1、2 で成果物が完全に一致した。作者は CUDA と HIP のユースケースが本機でスキップされ、cuda-gdb の挙動は直接テストしていないと注記している。

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

1.4 メインライン：アテンション示例がデバイスの共有メモリ上限に応じて tile を推論、RTX PRO 6000 上で 2 通過 7 失敗から 9 全通過へ（#3702）

11-02 06:59 にマージ、345 行追加、13 ファイル。複数のアテンション示例の非チューニング経路が Hopper サイズの tile をハードコードしており、sm_86、sm_89、sm_120 ではブロックあたりの動的共有メモリ上限が約 99 KiB のため、起動が直接失敗していた。エラー原文は、許可される動的共有メモリを 163840 バイトに設定するよう要求するもので、デバイス 0 のブロックあたり最大は 101376 バイトであった。現在、各示例は shared_memory_per_block_optin から tile を推論し、既存の GQA decode 示例（example_gqa_decode.py）の做法を踏襲している。データセンター構成は変更なし。各分岐は前の段の tile 要求を下回る場合にのみ発動するためである。約 99 KiB デバイスでの落とし所：attention sink の四個の示例はそれぞれ 64×64 二段階 128 スレッド（前向）と 64×32 一段階 128 スレッド（逆伝播）、example_gqa_bwd.py は前向 64×64、逆伝播 64×32、example_mla_decode.py は BLOCK_N 32、BLOCK_H 16 を取る。後ろ二者は threads パラメータを新規追加：64 行または 16 行 tile では T.gemm が四個の warp を必要とし、各 warp が 16 行ブロック全体を独占できるようにする。256 スレッドでは直接 lower 失敗となる（それぞれ acc_s と acc_s_cast の layout 衝突、および M=16 N=32 で有効な warp 分割がないと報告）。また、flash-attn 2.x には sm_12x カーネルがなく、四個の varlen 示例は元々参照ステップでエラーを投げていた。現在は flash_attn_varlen_supported を一度プローブし、非対応の場合は純 PyTorch の varlen_attention_ref（GQA、右下揃えの causal、window_size をカバー）に切り替える。検証は RTX PRO 6000 Blackwell（sm_120）上：CUDA 9 項目全通過（元は 2 通過 7 失敗）、CuTeDSL も同様に 9 項目全通過。H100 は 227 KiB の選択可能上限があり、元の tile のままで同様に 9 項目通過。

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

1.5 メインライン：CUDA float32x8 ビット変換が組み込み bit cast に変更、SASS 命令が 19 条から 18 条に減少（#3748）

11-02 09:58 にマージ、20 行追加、3 行削除、#3679 を修正し #3681 を置き換え。変更は CUDA のみ：float32x8 のビット再解釈をビット演算シーケンスから __builtin_bit_cast(uint64_t, make_float2(v, v)) に置き換え、生成される SASS 命令数が 19 条から 18 条に減少。幅度は大きくないが、コンパイラバックエンドにおける典型的な「意味等価、生成物がより短い」類の仕上げに属する。

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

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

2.1 昇騰：WY Fast 前向オペレータ新增、wy_fast_bwd_split 逆伝播セグメント移行（#1827 / #1828）

二者は同じく 11-02 02:33 にマージされ、前後の差は 1 分未満。#1827 は「移行および新增」の方式で WY Fast オペレータを Ascend NPU 側に移行し、#1828 は wy_fast_bwd_split のセグメント逆伝播を移行して補完。WY 表現法は線形アテンション系モデルにおいて再帰行列積を低ランク形式に簡約化する常用手段であり、前向と逆伝播セグメントの組み合わせにより、このオペレータチェーンを昇騰上で訓練可能なレベルまで補完したことになる。両 PR の説明はテンプレートの挨拶文のみで、変更範囲はタイトルとコミットに基づく。本レポートは内部実装の詳細について推測しない。

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

2.2 昇騰：T.tile.fill と T.tile.clear のドキュメントファクトチェックをマージ（#1634 / #1628）

同批でマージされたのはドキュメント側のファクトチェック二者で、それぞれ T.tile.fill と T.tile.clear の API 挙動を対象とする。#1628 の議論にはメンテナンス約定が一句残されている：この二つのインターフェースは比較的近いため、一方に修正がある場合はもう一方にも対応する調整が必要か同期して確認すること。これは継続的に進行中の API ファクトチェックバッチに属し、同期間には T.tile.bitwise_not（#1656）、T.tile.bitwise_and（#1651）、T.tile.bitwise_lshift（#1659）、T.tile.bitwise_or（#1655）も審査中である。

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

2.3 昇騰審査中：cython wrapper でシンボル形状解析をキャッシュ、ホスト側オーバーヘッドの系統を承接（#1967）

11-02 更新。#1967 は cython wrapper 内でシンボル形状の解析結果をキャッシュし、呼び出しごとの dispatch オーバーヘッドをさらに抑えることを目標としており、欠陥 #1898 に対応する。この系統は前号でマージされた #1948 と同一目標に対する2つの着地点である：#1948 は呼び出しごとに無条件で TVM Analyzer を構築する固定オーバーヘッドを遅延構築に変更し（910B3 上で約 450 マイクロ秒から約 106 マイクロ秒に短縮）、#1967 はシンボル経路における重複解析そのもののコストを扱う。同族には審査中の #3757 もあり、メインライン側では動的 ABI と共有シンボル検証を整合させる：表示名ではなく Var オブジェクトをキーとし、同名異変数が同一シンボルと見なされるのを回避し、シングルトン次元上の stride はフォールバックとしてのみ扱い、値の決定を先行させない。

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

2.4 昇騰審査中：z3 版をメインラインから同期移入、リポジトリ全体のメインライン移行も同時進行（#1965 / #1975）

#1965 は 11-02 12:18 に更新され、メインラインの z3 版 insert sync を Ascend リポジトリに移行した。#1975 は同日 03:28 に更新され、引き続き TileLang-Ascend 全体をメインラインのマルチバックエンドアーキテクチャへ移行するものである。両者は同一の収束経路における2つの層であり、一方は同期アルゴリズムの実装を差し替え、他方はリポジトリの骨格を差し替える。これに配套するメインライン側の動作として #3752（z3 rlimit 設定）が審査中であり、また #3735 は thread sync 内でモジュラーインデックスの非交差性を証明することを提案している——前号レポートで ThreadSync のロールバックが残した証明予算の問題は、「求解リソースを設定する」と「証明が必要な形状を減らす」の両端から同時に処理されつつある。

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

2.5 昇騰審査中：リダクションオペレータの2つの欠陥、先頭 Parallel 軸のブロードキャストをスカラ読み出しに変更（#1978 / #1968 / #1979）

3件とも 11-02 に更新された。#1978 は AscendVidReduction が行分割ループ内でブロック全体 [1,N] の UB に対して tile オペレータサイズを誤って半減させる問題を修正し（#1800 に対応）；#1968 は threads=2 時に sigmoid と round のカウントを半減させ（#1897 に対応）；#1979 は先頭 Parallel 軸上のブロードキャストオペランドをスカラ読み出しに変更し、VEC の UB アライメントの罠を回避する（#1885 に対応）。3者はいずれも同種の問題に属する：昇騰側のベクトル化経路は形状とアライメントに対する仮定が CUDA 側より厳しく、境界条件が破られるとコンパイル時または実行時のハードエラーとなる。

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

2.6 海光：gfx948 を gfx946 互換として処理、まず既存命令をテスト可能に（#32）

11-02 10:46 にマージ。変更は gfx948 を gfx946 互換として扱い、既存の MLS atoms、MMAC、barrier、コンパイルオプションを再利用するもので、目的は gather load の接続前に既存 builtins を新型号上で動作させることである。これは典型的な2段階アプローチである：まず新アーキテクチャ番号を既存経路に接続してテスト可能なベースラインを得て、その後で本当に新しい命令を必要とする部分を補う。当該リポジトリでは前回のウィンドウで通常のコピーに対する LDS swizzle 修正がマージされている。

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

2.7 摩尔線程、沐曦、mlir-ascend はウィンドウ内でプッシュなし（11-03）

tilelang-musa の最終プッシュは 09-30 で停止し、tilelang-metax は 10-26 で停止、tilelang-mlir-ascend は 10-31 08:21（本ウィンドウ起点より前）で停止している。そのため本期のバックエンド動向は昇騰と海光の両側に集中しており、うち昇騰はマージ18件中の4件、審査中更新44件中の多数を占めている。

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

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

3.1 TileOPs：単日7件マージ、6件が性能、1件がテスト基盤（11-02）

7 件すべてが 11-02 05:10 から 05:22 の間にマージされ、作者はいずれも lcy-seso で、attention、MoE、要素ごと演算、正規化、テストターゲット、CI 契約の 6 方向をカバーしている。前号の「1 日 13 件」と比べて件数は半減したが、構造が変わった。前号は shape ごとのチューニングの仕上げだったのに対し、今号ではテスト基盤（#2816）と CI 契約（#2818）という 2 種類の非性能変更が現れ、オペレータライブラリがマルチバックエンドと保守性に向けた構造的な準備を始めたことを示している。各件は引き続き H200 の実測表を添付し、manifest に測定済み行の回帰がないことを宣言している。

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

3.2 TileOPs：warp 専用 decode の softmax と QK/PV を重ね、478 回のスキャン呼び出しで幾何平均 0.952 倍（#2817）

11-02 05:21 にマージ。変更は warp 専用 GQA decode の softmax と QK、PV 行列積を重ねるもので、確率はレジスタに留め、K と V は独立した barrier を使い、1 つの 3 スロットリングバッファを共有する。選択条件は D が 128、64 個の key の tile、KV ヘッドごとに 1 行の tile、softcap なし、かつ「1 CTA 1 SM」のグリッドが占有率とストリーム長の境界を満たせること。満たさない呼び出しは引き続き直列プログラムを通る。9 つのプランナ境界はいずれもホスト側で固定デバイスファクトを用いて検証され、実際の attention テンソルの割り当てや 132 SM のカードを必要としない。実測環境は H200、CUDA 13.2、PyTorch 2.13.0+cu132、TileLang 0.1.12 で、CUPTI のデバイスビジー時間を使用：形状 (16, 8, 4096, 64, 8, 128) で float16 は 0.074610 ミリ秒対 fa3 の 0.077490（1.039 倍、3.654 TB/s）、bfloat16 は 0.074260 対 0.077220（1.040 倍、3.671 TB/s）。影響を受ける 478 回のスキャン呼び出しで記録された幾何平均レイテンシは main の 0.952 倍だった。

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

3.3 TileOPs：MoE を token ごとに収集して事前配置に変更、H200 で 387.8 TFLOPS（#2815）

11-02 05:19 にマージ。tight レイアウトの MoE 入力は token ごとの収集に変わり、同一の hidden 行がその複数のルート間で再利用される。token が GPU を満たさない場合はルートを異なるブロックに分割する。同時に tight スキャンで使われていなかったエキスパート順のソースインデックス割り当てと書き込みを削除した。スケジューリングには新しい規則を追加：未融合、K 主順の tight グループで平均行数が 128 より大きく 192 以下なら、H200 上の ping-pong 64×256 tile を選ぶ。グループサイズを明示した呼び出しは引き続きコストモデルを通る。実測 (4096, 256, 8, 7168, 2048) で bfloat16 は 7.4417 ミリ秒対 flashinfer-cutlass の 7.6982（1.034 倍、387.844 TFLOPS）、float16 は 7.5040 対 7.9869（1.064 倍、384.624 TFLOPS）。記録済みのルートおよび偶数グループのスキャンは両 dtype とバイアスモードをカバーし、manifest にノイズ超過の回帰はない。

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

3.4 TileOPs：16 ビット要素ごと演算で範囲チェック付き除算を除去、8 バイトオペランドは 1 スレッド 1 ベクトル（#2814 / #2813）

#2814 は 11-02 05:18 にマージ。bfloat16 の Sigmoid と 16 ビット SiLUAndMul に対して 4 分の 1 スケールの逆数を使い、範囲チェック付き除算を回避し、この補助関数を SiLU と共用する。同時に指数演算で不要な非正規数ガードを省き、融合ゲート出力をスレッドをまたいで隣接する 16 バイトベクトルごとに格納し、gate が負の 88 に近いときに無限値を乗じる未丸め SiLU 活性化を保持する。実測で bfloat16 Sigmoid は 16777216 要素で 0.017504 ミリ秒対 torch-compile の 0.018656（1.066 倍、3.834 TB/s）、268435456 要素で 0.248545 対 0.262529（1.056 倍、4.320 TB/s）。SiLUAndMul は (2048, 14336) で 0.041856 対 0.045440（1.086 倍）。#2813 は 05:17 にマージ。8 バイトオペランドを長い起動下で 1 スレッド 1 ベクトルとして処理するもので、変更前に TileFoundry の roofline 分析でこの種のビット単位 OR オペレータがすべてメモリ帯域制限であることを確認している。4.8 TB/s の HBM 帯域に基づく各 workload の下限は、hidden-state-prefill の int64 が 41.9 マイクロ秒、int32 が 21.0 マイクロ秒、llama-7b-ffn の int64 が 56.4、int32 が 28.2、cnn-feat-broadcast の int64 が 42.8、linear-bias の int64 が 14.0、int32 が 7.0 である。

ソース：https://github.com/tile-ai/TileOPs/pull/2814、https://github.com/tile-ai/TileOPs/pull/2813

3.5 TileOPs：Norm 統計量を行内要素の減算に変更し、1e36 付近のオーバーフローを解消（#2811）

11-02 05:10 にマージ。根本原因は非常に直接的で、すべての平均と分散のカーネルが生の行に対して fp32 の総和を取っており、和が行のオフセット量（約 N 掛ける x ）とともに増大し、1e36 付近の行は直接オーバーフローして出力がすべて NaN になる。オーバーフローしなくても、平均はオフセット量の桁で丸められ、x − mean がこの丸めを各要素に持ち込み、その結果、定数行で算出される分散は丸めの二乗となり、正負の weight ではなく bias に正規化されてしまう。修正方法は、行から構成される各 fp32 量に不変条件を一つ加えることである。すなわち、それは行の拡散度によってのみ制約され、オフセット量によっては制約されない。実装上は「行内のある要素」を基準としてから減算し、warp 主導のカーネルでは warp の先頭要素を、シャードカーネルでは行の先頭要素を用いる。これは前回のレポートで #2812 が列挙した四つの Norm オペレータの欠陥と同じ種類の問題に対する正面からの修正である。

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

3.6 TileOPs：tileops_test ターゲットで tests/ops を CUDA から切り離して実行（#2816）

11-02 05:20 にマージ。本調査期間で単一変更として最大のもので、1281 行追加、524 行削除、59 ファイル、#2803 をクローズ。テスト時にのみ存在する tileops_test ターゲットを新設し、PrivateUse1 上で各オペレータの共有 workload 参照を計算し、tests/ops のスモークテストはこれを経由して実行される。サポートされない参照呼び出しは明示的に unvalidated として列挙してスキップし、CUDA 上でのみ成立する除外項目も同じターゲット下で検証する。オペレータから workload へのレジストリは workloads/calls を通じて共有され、ベンチマーク専用のバインディング、コピー会計と入力上書きは benchmarks/_cases に残る。さらにテストデバイス生成器、畳み込みサポートを追加し、テスト機能マーカーを修正し、参照実装が GQA 逆方向の scale と空行のセマンティクスに従うようにした。作者は限界を目立つ場所に書いている。このターゲットがテストするのはバックエンド非依存の契約とテストの可搬性であり、共有参照との比較は実際のバックエンドの数値的正しさを証明するものではなく、#2611 は依然としてオープンのままである。利用可能な参照がない呼び出しは引き続き unvalidated の形で可視化される。同じバッチの #2818 は dispatch 契約を workload データ生成から切り離し、契約ケース向けには形状のみをインスタンス化し、使われないテンソル内容の生成も検証も行わず、五つのホスト構造テストスイートを compile-contract-gate 内で一度だけ実行するようにした。

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

3.7 TileFoundry：共有メモリの順序付けを証明可能な分析層にし、83 個のフィクスチャで 52 計画・31 拒否（#283）

11-02 02:09 にマージ。問題は現実のものである。finalize が共有メモリアクセス間に commit、wait、barrier を発行しないため、共有メモリ経由で一時保存するスケジューラは競合し、カード上の sm_89 の E4M3 GEMM は torch に対して最大 659 ずれる。#267 で定めた方向は「HIR 上でアクセス順序を並べる分析」に「ターゲットに従って順序を発行する lowering」を加え、最小の正しい形式のみを発行し、証明できなければ拒否するというもので、本コミットが納品するのは分析部分である。新設の同期ファミリは lowering の発行順に各関数をたどる。各共有メモリアクセスについて、analysis/emission.py から命令、発行 mesh とオペランドを取得し、IterationScope からスコープを取得し、メモリファミリが記録した配置からバイト数を取得する。競合アクセスの前に、commit 一回と wait 一回（未完了の cp.async 書き込み向け）および「両方を同時に保持する最内スコープ」を覆うフェンス一回を計画し、先に実行される子ノードのアクティブ位置で番号付けする。オペレータは新レジストリ visitor_registry/completion.py を通じて効果がどのように完了するかを宣言し、何も宣言しないものは UNKNOWN として扱う。このファミリは UNKNOWN 後のアクセス、戻り時にまだ飛行中の UNKNOWN アクセス、スレッド境界をまたいで待たないループ内へのフェンス挿入、および予約呼び出しで呼び出し先が共有メモリを保持する状況を拒否する。デフォルトではこのファミリを要求する呼び出し元がないため、finalize を行う 75 個のスケジューリングフィクスチャのバイトは main と完全に同一である。83 個のフィクスチャのうち 52 を計画し 31 を拒否した（24 個は TMA コピーを使用、6 個は wgmma を使用し完了セマンティクスが未宣言、1 個はスレッド関連境界のループ内へのフェンス挿入が必要）。コーパスケースはこれらの拒否を個別にアサートしている。

ソース：https://github.com/tile-ai/TileFoundry/pull/283

3.8 TileFoundry 審査中：CUDA lowering が同期計画を発行し逆検証、main 上で 7 プログラム中 4 つが誤り（#285）

11-02 12:40 更新。これは #283 の実行側である。CUDA lowering は同期計画に従ってステップを発行し、実際の発行が計画と一致するかを一つずつ検査し、一致しなければ LoweringError を投げる。コーパス上の結果：デフォルトフローでは 14 個のプログラムが同期付きで通過し、31 個を拒否した。この 14 個の TIR は main との差分が新增ステートメントのみで、別の 30 個は main と完全に同一である。コア visitor を通した後、main が finalize する 75 個のプログラムはバイト単位で一致した。単一の sm_89 上でビルド可能かつ同期付きの 7 個のプログラムを実行した：6 個は参照と完全に一致し、synccheck と memcheck の下で全てクリーンであった。同じ 7 個を main の方式で lower すると、4 個で結果が誤った。作者は同時に main に既存する二つの問題も報告している：GEMM と loop_outputs で racecheck が同一 cp.async の複数レーンによる同一アドレスへの書き込みを報告し（HIR の copy_async は分割ターゲットを受け付けないため、各スレッドがタイル全体をコピーする）、loop_outputs はさらに一つの出力を誤った要素へ書き込む。

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

3.9 採用側：FlashQLA、TileKernels は調査期間内にプッシュなし（11-03）

QwenLM/FlashQLA の最終プッシュは 10-30 03:02 で停止し、deepseek-ai/TileKernels は 10-14 05:49 で停止しており、両リポジトリとも本調査期間内にコミットはない。TileKernels は 10-14 以降、既に複数期連続でプッシュがなく、採用側の可視活動は今期すべて TileOPs と TileFoundry という自社リポジトリの二つに集中している。

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

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

4.1 昇騰の同期正確性が集中報告面に：単日 8 件の新規 issue、二人の報告者がそれぞれ体系的に提出（11-02）

調査期間内にメインリポジトリで新規 issue が 8 件追加され、番号は #3740 から #3749 に集中し、すべて昇騰 AscendC と PTO の二つのバックエンドのコード生成と同期挿入を指している。うち 7 件は同一報告者により提出され、環境は TileLang 0.1.15、PyTorch 2.10.0+cpu、torch_npu 2.10.0 で、主題は同一プログラムが二つのバックエンドで挙動が分岐する点である：生成コードが起動するコアで不一致（AscendC は mix(1,2)、PTO は vector を出力）、enable_auto_schedule=True のとき AscendC が同期を挿入しない、自動スケジューリングをオフにすると生成成功するがオンにすると失敗しエラーメッセージが不明瞭、同一コードが AscendC では生成できるが PTO では失敗する。この一群の報告の共通点は、いずれも再現可能な prim_func と行単位で対照した生成コードを添付しており、直接回帰テストケースに変換できる形態であることだ。さらに審査中の #3735（thread sync でモジュール式インデックスの非交差を証明）と #3580（述語と昇格インデックスの同期）を加えると、昇騰側の同期正確性は現在最も集中した一線となっている。

出典：https://github.com/tile-ai/tilelang/issues/3747、https://github.com/tile-ai/tilelang/issues/3745

4.2 行跨ぎと転置ビューの SIMD ロードが DMA 依存を失う、根本原因は論理領域の再構築（#3749）

同批で技術的含有量が最も高い一件で、報告者の環境は main ブランチ e8225d7f5、Ascend950PR、CANN 9.2.0。問題は二箇所の連続 SIMD ロードが、より早い GM から UB へのコピーへの依存を失うことである：一つは行優先ロードが行の途中から開始し次の行へ跨ぐ場合、もう一つは転置ビューにアクセスし物理連続次元が最後の論理次元でない場合。根本原因の経路は詳細に記されている：生成された tl.access_ptr は物理ロード範囲を覆うが、MemoryAccessDetector::LogicalPointerRegion が再構築した論理領域はその一部を漏らし、依存分析は入力 DMA とロードが無関係と見なし、AutoSchedule が入力コピーをロードの後ろへ移動し、InsertSync も MTE2 から V への set/wait を発行しない。最小再現では work.shape が (2, 128) で、まず T.copy(A, work[1, 0:64])、次に T.SimdVF() 内で T.simd.vld(work[0, 64]) により 128 個の連続半精度要素を読む：生成された読み取り記述子は tl.access_ptr(work[0, 64], 128, 1) で、物理的に [64, 192) を覆い、うち [128, 192) はまさに直前に DMA が書き込んだ部分であり、再構築された領域は work[0:1, 0:128] すなわち物理 [0, 128) で、書き込み区間を完全に外している；行数は ceil(128/128) で 1 行と算出され、開始列 64 が算入されていない。この件は本調査期間の主力が補強している FlashMLA、PTO ベクトル化と同じ土台であり、一旦命中すればサイレントな誤結果となる。

ソース：https://github.com/tile-ai/tilelang/issues/3749

4.3 ニュース側：Google News 中英語、Hacker News、arXiv の三チャネルはいずれもウィンドウ内でヒットなし（11-03）

6組の Google News クエリ（英語 TileLang、TileLang GPU kernel、tile-ai TileLang、中国語 TileLang オペレータ、TileLang 昇騰、沐曦または海光または摩尔線程）が返した最新の関連報道は 10-21 の SemiAnalysis 華為と DeepSeek の記事で、残りは 09-29 から 10-01 の DeepSeek オープンソース昇騰コンポーネントの一連に集中しており、すべて本ウィンドウより前である。arXiv 上の all:"TileLang" の最新ヒットは 2026-07-24 の TileSight で、TileLang と Ascend の複合クエリは結果なし。Hacker News の候補エントリのうち TileLang 関連はリポジトリプッシュのエントリのみ。候補リスト 14 件のうち 12 件は GitHub リポジトリ発見によるもので、2 件は無関係な Show HN と LSP リポジトリであり、本期の情報面は完全にリポジトリ側で構成されている。

ソース：https://news.google.com/、https://arxiv.org/

5. トレンド観察

この日最も記録に値する判断は、昇騰がもはや「適応目標」ではなく「初回リリース目標」になったということだ。FlashMLA のような DeepSeek 自社デコードパス上の核心カーネルが、昇騰 950 を直接ターゲットとしてメインラインに書かれており、しかも T.Stage、annotate_buffer_addresses といった明示的スケジューリング向けの低レベル機能を使っており、別リポジトリを開いて移植するのではない。コンパイラ側の 3 つのサポート（L0C から UB への NZ 保持、AIV ストリームのパイプライン前選択、ソース順での明示的ステージ発行）と合わせると、昇騰パスは「動く」から「スケジューリング意図どおりに動く」へと移行しつつあることがわかる。

第二の手がかりは、正確性の重心が下がっていることだ。前号のメインラインでは TIRX 移行のために ThreadSync 正確性改善一式をリバートしたが、今号ではすぐに 2 つの補完が現れた。メインライン #3752 は z3 でリソース上限を解き、#3735 はモジュール化インデックスの非交差性で証明が必要な形状を減らすことを提案。同時に #3749 は「依存性分析が物理ロード範囲を見えない」欠陥を MemoryAccessDetector::LogicalPointerRegion の層まで掘り下げた。これらを TileFoundry の暗線と合わせて見ると——分析層は 83 のフィクスチャのうち 52 の共有メモリ順序を証明でき、lowering を main の方式で行うと 7 プログラム中 4 つが誤計算になる——同じシグナルが 3 回現れている：TileLang の現在の主要リスクは機能不足ではなく、「生成コードが見た目は正しいが実際は正しくない」ことであり、チームが選んだアプローチは正確性を却下可能な証明にすることであり、パッチで個別に塞ぐことではない。

第三は、テストと契約が独立した層になり始めたことだ。TileOPs の tileops_test はスモークテストを CUDA から切り離し、compile-contract-gate は dispatch 契約と workload データ生成を分離し、両者は同じ 1 時間内にマージされた。作者は同時に「共有参照は実バックエンドの正しさを証明できない」と明記している。この「何をテストできるか」と「何を証明したか」を分けて表現する習慣は、性能パッチを数本書き足すよりも、このプロジェクトがプロトタイプ期を抜けつつあることをよく示している。

最後に一つの対照に注意：オペレータ側は単日 7 件の性能マージ、6 件が H200 実測表付きである一方、ニュース側の三チャネルは 24 時間ゼロヒット。TileLang の対外プレゼンスは依然「DeepSeek オープンソース昇騰コンポーネント」のようなイベントに集中しており、日常の進展はほぼリポジトリ内でのみ起きている。このプロジェクトを追う者にとって、これはデイリーレポートの主体が報道ではなくコミットと PR にならざるを得ないことを意味する。

付、素材と検証説明

情報源検証表

情報源 検証結果
tile-ai 組織（29 リポジトリ） ウィンドウ内プッシュ 9 リポジトリ
メインリポジトリ tilelang メインブランチ 5 件マージ；新規 issue 8 件；ドキュメントサイト 5 件自動ビルド
TileOPs メインブランチ 7 件マージ；サイト 1 件
TileFoundry メインブランチ 1 件マージ；1 件レビュー中新增
tilelang-ascend メインブランチ 4 件マージ；6 件レビュー中更新
tilelang-hygon 1 件マージ
tilelang-lsp 1 件マージ（CUDA デバッグ行マッピングに整合）
TileOPs-nightly 1 回自動プッシュ
tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack ウィンドウ内プッシュなし
リリースページ（tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon） ウィンドウ内新バージョンなし
FlashQLA / TileKernels ウィンドウ内プッシュなし
Google News（中英語 6 組クエリ）/ Hacker News / arXiv ウィンドウ内実質ヒットなし

完全な情報源リスト