TileLang デイリーレポート (2026-10-28)
調査期間:2026-10-27 07:00 ~ 2026-10-28 07:00(北京時間、24時間)。
情報源:GitHub tile-ai 組織のリポジトリ別調査(tilelang / TileOPs / TileOPs-nightly / TileFoundry / tilelang-hygon / tilelang-metax / tilelang-ascend / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / ドキュメントサイトと夜間ベンチマークリポジトリ)+ マージ済みおよび審査中の真値検索(マージ済み 20 件、審査中更新 45 件)+ 新規 issue 全数確認(7 件)+ コミットレベルの diff 確認 + リリースページと PyPI 成果物の確認 + 外部バックエンドリポジトリの確認(Hexagon / CANN オペレータリポジトリ / PPU ラッパー / モデル断片リポジトリ)+ Google News の中国語・英語複数組クエリ(プロキシ経由)+ Hacker News + arXiv。
本期インデックス
- 本日の重点:TileOPs 0.1.0 が正式リリース、206 個のオペレータ、2265 行のベンチマークのうち 2262 行が最速実装より遅くない;24 時間で 20 件をマージ、審査中更新 45 件、新規 issue 7 件;海光 scale バッファのデュアルバンクを 1389 行マージ;サードパーティ Hexagon バックエンドは単日 15 件
-
- コアプロジェクトの進展
- 1.1 メインリポジトリ:ROCm の exp/sin/cos がデフォルトで高速組み込みに降格しなくなり、sin/cos の誤差が 7.8e-04 から 6.6e-08 に低減(#3621)(10-27)
- 1.2 メインリポジトリ:自動ターゲットが現在の HIP デバイスを読み取るよう変更、マルチ GPU マシンで 0 番カードのアーキテクチャに固定されなくなった(#3594)(10-27)
- 1.3 メインリポジトリ:TMA gather/scatter の定数行インデックスを int32 に変換、リテラル記法でコンパイルがクラッシュしなくなった(#3365)(10-27)
- 1.4 メインリポジトリ:NVRTC と CuTeDSL が最初の入力テンソルの所在デバイスに応じて現在のストリームを取得(#3359)(10-27)
- 1.5 メインリポジトリ:自動チューニングのキャッシュキーに compile_flags を組み込み、同一カーネルでコンパイルオプションを変えても古いバイナリを再利用しなくなった(#3338)(10-27)
- 1.6 メインリポジトリ:ThreadSync が address_of を書き込みとして記録、不透明呼び出し後のフェンスを補完(#3346)(10-27)
- 1.7 メインリポジトリ:ランタイム行程数のパイプライン epilogue を直接展開、ページドデコードが再びコンパイル可能に(#3356)(10-27)
- 1.8 メインリポジトリ:スライス充填と負インデックスの静的境界を修正、範囲外書き込みと暗黙の負アドレスを解消(#3352 / #3347)(10-27)
- 1.9 メインリポジトリ:T.ws(-1) 領域が全スレッドで実行され、単独記法でもストアが失われなくなった(#3354)(10-27)
- 1.10 メインリポジトリ:profiler がデバイス未指定時に現在のデバイスを使用、キャッシュフラッシュが 0 番カードに当たらなくなった(#3357)(10-27)
- 1.11 新規欠陥:カーネルキャッシュのヒットがコールドコンパイルを前提とする 6 つのテストを破壊;CuTeDSL の CUDA graph キャプチャが空グラフになる(#3694 / #3693)(10-27)
- コアプロジェクトの進展
-
- マルチバックエンド適配(昇騰 / 海光 / ROCm / Hexagon)
- 2.1 海光:HCU scale バッファをデュアルバンクで分区、blockscaled GEMM が手動 warp 特化下で同一の行マッピングを共有(#29)(10-27)
- 2.2 昇騰メインリポジトリ:PerCoreTask 候補体が内部 if/else を許可、条件付き L1 充填が全体拒否されなくなった(#3686)(10-27)
- 2.3 昇騰メインリポジトリ:CI が AscendC と PTO を 2 つのマトリクスジョブに分割、個別に再実行可能(#3687)(10-27)
- 2.4 昇騰リポジトリ審査中:Cython ラッパーがアナライザを遅延構築、昨日記録した呼び出しごとの再構築問題に直接対応(#1948)(10-27)
- 2.5 昇騰リポジトリ審査中:GLM-5.3 の kpool 圧縮オペレータ、PTO 上のテンソルシフト量(#1910 / #1949)(10-27)
- 2.6 昇騰メインリポジトリ審査中:PTO の T.Parallel ベクトル化、AutoSimVF チャネル、バッチ三次元 ND2NZ、書き込みカウントのマルチバッファ、Bisheng ビルド識別子のキャッシュキーへの組み込み(#3512 / #3320 / #3689 / #3629 / #3691)(10-27)
- 2.7 ROCm 審査中:pip 版 ROCm 10 のアーキテクチャ検出とミラー、FP4 変換と E8M0 型(#3549 / #3665)(10-27)
- 2.8 昇騰リポジトリ日次テスト:2903 項目すべて通過、2 日連続で失敗なし(#1950)(10-28)
- 2.9 沐曦:組織ビューにプッシュタイムスタンプあり、6 ブランチはいずれも調査期間内の新規コミットなし(10-27)
- 2.10 サードパーティ:Hexagon DSP バックエンドが単日 15 件、HVX ベクトル化、VTCM アリーナ、Crouton レイアウト、DMA 組み込みが揃う(10-27)
- マルチバックエンド適配(昇騰 / 海光 / ROCm / Hexagon)
-
- エコシステムと採用側(TileOPs / TileFoundry)
- 3.1 TileOPs 0.1.0 リリース:206 個のオペレータの初の正式版、公開 API がこれで確定(10-27)
- 3.2 TileOPs 0.1.0 ベンチマーク:H200 上で 2265 行のうち 2262 行が最速実装より遅くなく、幾何平均 1.37 倍(10-27)
- 3.3 TileOPs 0.1.0 の既知のギャップと次のステップ:融合オペレータ、モデル全体、計算通信融合(10-27)
- 3.4 TileOPs 4 件マージ:where と masked_fill がチューニング可能に、D128 因果対角が 64 キーごとにスコアリング、スパース長 GDN プリフィルの分割、変更ログをリリース日基準で日付設定(#2782 / #2783 / #2779 / #2781)(10-27)
- 3.5 TileFoundry 2 件マージ:ウィンドウ全体をオペランド自体に応じて発行、動的共有メモリをカーネル実使用量に応じてサイズ決定(#273 / #271)(10-27)
- 3.6 外部エコシステム:CANN コミュニティの TileLang オペレータリポジトリ、PPU 版ゼロ変更ラッパー、DeepSeek V4.1 Flash スニペット例(10-27)
- エコシステムと採用側(TileOPs / TileFoundry)
-
- コミュニティ、チュートリアル、イベント
- 4.1 ドキュメントサイトが単日 6 件の自動更新(10-27)
- 4.2 昇騰リポジトリのドキュメントファクトチェック 4 件、単一コメントが最多
- コミュニティ、チュートリアル、イベント
46 コミット(#1628 / #1655 / #1656 / #1659)(10-27)
- 4.3 審査中:lower trace を独立した開発ツールキットとして抽出(#3583)(10-27)
- 4.4 ニュース側:Google News と HN では24時間以内に今回の調査期間に関する新規報道なし(10-27)
-
- トレンド観察
- 付録、素材と検証説明
今日の重点
TileOPs 0.1.0 が本日の調査期間内にリリースされ、連日続いた「最後の1マスで止まる」状態に終止符が打たれた。PyPI 上の 0.1.0 成果物は 10-27 07:45(北京時間)にアップロードされ、前回デイリーレポートの締め切り時刻より45分遅い;GitHub リリースページは 13:15 に公開され、リポジトリに初めて正式な tag が付いた。リリースノートのハード指標は3グループ:206 個のオペレータすべてに manifest エントリがあり(シグネチャ、テストとベンチマーク負荷、roofline 公式、CI で逐条検証);H200 上の 2265 行の夜間ベンチマークのうち 2262 行が最速の対照実装より遅くなく、幾何平均 1.37 倍、対照には torch.compile、FlashAttention 3、FlashInfer、vLLM、DeepGEMM、FlagGems を含む;206 個のオペレータのうち 195 件の foundry タグ付きマージは TileFoundry がカーネルを生成し同一の CI を通してリポジトリに取り込まれたものである。リリースと同時に公開 API が凍結され、以降 0.x は追加のみで変更なし、さらに API 互換性チェックスクリプトが新設され、前バージョンの書き方で呼び出す PR は直接失敗する。
24時間以内に tile-ai 組織は 20 件をマージ(tilelang 13、TileOPs 4、TileFoundry 2、tilelang-hygon 1)、審査中更新 45 件、新規 issue 7 件(うち 2 件は昇騰リポジトリのデイリーテストボット)。マージの重心は2つの線に分かれる:1つは浮動小数点セマンティクスと実行コンテキストで、ROCm は7つのデフォルト数学関数を高速組み込みから正確実装に戻し(MI300X 上の sin/cos 誤差が 7.8e-04 から 6.6e-08 に低減)、自動ターゲットは固定の 0 番カードではなく現在のデバイスを読むように変更、NVRTC と CuTeDSL は入力テンソルが存在するデバイスに応じてストリームを取得するように変更、自動チューニングのキャッシュキーに compile_flags を追加;もう1つは静的境界とスレッドドメインで、スライス填充は長さではなく領域末端で検証、範囲外の負インデックスはコンパイル期に拒否、T.ws(-1) の領域は全スレッドで実行するように変更。海光は 12 ファイル 1389 行をマージし、scale バッファを独立したデュアル bank に分割し LDS bank と分離した。サードパーティ側では Hexagon DSP バックエンドが単日 15 コミット、HVX ベクトル化、VTCM アリーナ、Crouton レイアウトと DMA 組み込みを一気に Crouton 上のマッピングとリダクションまで展開した。
1. コアプロジェクト進展
1.1 メインリポジトリ:ROCm の exp/sin/cos がデフォルトで高速組み込みに降格しなくなり、sin/cos 誤差が 7.8e-04 から 6.6e-08 に低減(#3621)
10-27 22:39 にマージ、5 ファイル +107/-17、#3618 を修正。HIP 上では、T.exp、T.exp10、T.log、T.log2、T.log10、T.sin、T.cos の7つのデフォルト関数が tl.enable_fast_math の設定に関わらず __expf 一类の高速組み込みに降格していたが、CUDA 側は #875 の取り決めに従い正確実装を使用していた;逆に、明示的な T.__exp などの高速オペレータは HIP 上では全く下沉せず、tilelang.rocm.language 内で直接 AttributeError をスローしていた。本コミットで HIP と CUDA が同一戦略を採用:7つのデフォルトオペレータは HIPMath に委譲、8つのアンダースコア付き高速オペレータには HIPFastMath を指す最低下沉ルールを新設し、統一して FAST_MATH_INTRINSICS リストに移動して ROCm 言語層からエクスポートする。精度比較は #3618 の再現を使用:MI300X 上で exp と exp10 は相対誤差、sin と cos は絶対誤差(x は 1e4 まで)を見る。4列は順に T.exp、T.exp10、T.sin、T.cos、主干は 3.7e-06 / 4.1e-06 / 7.8e-04 / 7.8e-04、本コミットは 8.3e-08 / 8.1e-08 / 6.6e-08 / 6.4e-08。
出典:https://github.com/tile-ai/tilelang/pull/3621
1.2 メインリポジトリ:自動ターゲットが現在の HIP デバイスを読むように変更、マルチ GPU マシンで固定の 0 番カードアーキテクチャを取得しなくなる(#3594)
10-27 22:37 にマージ、2 ファイル +24/-1。determine_target(“auto”) は ROCm 上で登録済みの HIP 検出器に落ちるが、この検出器は無条件に torch.cuda.get_device_properties(0) を読むため、現在の HIP デバイスが 0 番カードでない場合、自動検出は 0 番カードのアーキテクチャを返してしまう。CUDA 側は早くから torch.cuda.current_device() を使っていたが、ROCm は追いついていなかった。修正方法は、まず現在のデバイスインデックスを取得してから属性を照会し、normalize_rocm_arch で機能サフィックスを除去する。追加テストでは current_device() をスタブ化して 1 を返し、同時に検出されたすべてのインデックスを記録し、gfx942 が検出されかつ 1 番カードにしか触れていないことをアサートする。修正前はこのアサートで probed == [0] が得られた。著者は本機に ROCm デバイスがないことを明記しており、テストはスタブ化した torch のみで検出ヘルパーを検証し、実際の HIP ランタイムは実行していない。本日別の #3549 の変更とは異なる関数である(#3549 は tilelang/contrib/rocm.py、本件は tilelang/rocm/target.py)。
出典:https://github.com/tile-ai/tilelang/pull/3594
1.3 メインリポジトリ:TMA gather/scatter の定数行インデックスを int32 に変換、リテラル記述でコンパイルクラッシュしなくなる(#3365)
10-27 09:15 にマージ、2 ファイル +44/-2、#2601 を修正。tma_gather4 と tma_scatter4 の 2 つのフロントエンドヘルパーは、Python 整数形式の行インデックスを受け取るとそのまま透過させ、TIR 内で確定した整数幅を持たないため Blackwell 上でコンパイルクラッシュしていた。他の行式はそのまま維持する。本件は Python 整数を int32 の TIR 定数に変換するのみで、Blackwell ゲート付きのコンパイル回帰を追加し、リテラル行インデックスで gather と scatter をそれぞれコンパイルし、生成された CUDA ソースに対応する命令が確かに現れることを確認する。著者はこのテストがコンパイル可能なことのみを検証し、実行時の数値的正しさは検証しないことを明記している。また提出者の本機の Windows Python 3.13 環境では Home Assistant の pytest プラグインが fcntl に依存しており、pytest の収集段階で失敗するため、本件の TileLang テストはローカルでは実際には実行できていない。
出典:https://github.com/tile-ai/tilelang/pull/3365
1.4 メインリポジトリ:NVRTC と CuTeDSL が最初の入力テンソルの所在デバイスに応じて現在のストリームを取得(#3359)
10-27 08:11 にマージ、2 ファイル +5/-4、#3301 の F149 と #3408 に続く。2 つのアダプタは以前、文字列プレフィックスでターゲット種別を判定し、かつデフォルトで呼び出し元の現在デバイスの現在ストリームを取得していたため、呼び出し元の現在 GPU とテンソルの所在 GPU が一致しない場合に誤ったストリームを選択していた。本件はターゲット種別判定に改め、最初の入力テンソルのデバイスに応じて現在ストリームを照会する。明示的に渡されたストリームハンドルは引き続き最優先であり、テンソルのない NVRTC 呼び出しは元のセマンティクスを維持する。これは昨日の out_idx 出力割り当てストリーム(#3649)に続く、ストリームセマンティクス整理の第 2 弾である。著者は境界を明確にしている。本件はストリーム選択のみを修正し、クロスデバイスのデフォルトストリーム起動コンテキストは未解決項目のままであり、この種の呼び出し元は起動前にテンソルの所在デバイスを現在デバイスに設定する必要があり、パッチには呼び出しごとのデバイスガードはない。検証では実際の Torch ストリームと GPU 起動で同一デバイスとクロスデバイスの非デフォルトストリームをカバーしたが、クロスデバイスのデフォルトストリームコンテキストの正しさは確立していない。
出典:https://github.com/tile-ai/tilelang/pull/3359
1.5 メインリポジトリ:自動チューニングのキャッシュキーに compile_flags を組み込み、同一カーネルでコンパイルオプションを変えても旧バイナリを再利用しなくなる(#3338)
10-27 16:57 にマージ、3 ファイル +28/-1、#3301 の F131 に続く。自動チューナーはコンパイル結果を永続キャッシュするが、キャッシュ識別子に compile_flags が含まれていなかったため、同一カーネルでデバイスコンパイラオプションを変えても旧バイナリにヒットし、測定された性能と実際のコンパイル結果が一致しなかった。本件は CompileArgs に実際に保存させ、元の順序でハッシュに参加させ、さらに flag を JIT 実装から AutoTuner.set_compile_args() を経て候補コンパイル、設定ごとの pass マージ、グループ化されたデバイスコンパイル、キャッシュカーネルロードまで一貫して透過させる。検証は既存の 3 つのキャッシュキーテストと、ネイティブランタイムによる CUDA の定点比較をカバーする。リストと文字列の両方の flag 形式が設定ごとの上書きで生存できる。著者は、従来のマージロジックがコンパイラマクロを落としており、まさに今回塞ぐべき穴であると指摘している。
出典:https://github.com/tile-ai/tilelang/pull/3338
1.6 メインリポジトリ:ThreadSync が address_of を書き込みとして記録し、不透明呼び出し後のフェンスを補完(#3346)
10-27 05:49 にマージ、2 ファイル +62/-1、#3301 の F305 を引き継ぎ、昨日のフェンスセマンティクスの主線を継続する。address_of 自体はアクセス情報を運ばず、ポインタを下流に渡すだけだが、アクセス分類器がその型を読み取りにハードコードしているため、共有バッファが不透明な呼び出しを経て書き込まれても writes 集合に入らず、後続の全ブロック読み取りは競合する書き込みを見つけられず、フェンスも計画されなくなり、スレッドはこうして書き込みを追い越して書き込み前の内容を読んでしまう。実測の CUDA 出力は非常に端的である。T.address_of(B[0]) 形式は extern の後にフェンスがなく、T.access_ptr(B, “w”) 形式にはある。後者はアクセス型を明示的に宣言しているため、この穴はずっと発見されなかった。loop_unswitching.cc はすでに address_of を書き込みとして保守的に扱っており、ThreadSync が唯一の例外である。本コミットは address_of を書き込みとして記録する。方向は保守的である。ポインタの用途が不明なら書き込まれる可能性があり、フェンス計画は安全側で一度誤えばよい。新たに 2 つの用例を追加し、当該カーネルが下沉された後、extern の後に __syncthreads() が続くことを断言する。
出典:https://github.com/tile-ai/tilelang/pull/3346
1.7 メインリポジトリ:実行時行程数のパイプライン epilogue を直接展開、ページングデコードがコンパイル可能に復帰(#3356)
10-27 07:02 にマージ、1 ファイル +5/-1。T.Pipelined ループの行程数が実行期の値である場合、epilogue の展開が 2 ステップ以上で非同期コピーを伴うと、CUDA codegen は TypeError: Downcast from tirx.Sub to ir.IntImm failed を報告する。根本原因は EmitImpl にある。prologue と epilogue の extent が定数の場合にのみ展開し、各反復の待機カウントが定数であることを保証する。一方、epilogue の extent は end - start、すなわち (n + max_stage) - n であり、n が定数の場合は構築期に折り畳まれ、n が実行期の場合は Sub のままとなり、as_const_int は空を返す。そのため epilogue は T.unroll ループとして発射され、その待機は T.ptx_wait_group(1 - k) のようにループ変数を読む形式で書かれ、ptx_wait_group は即値しか受け取らない。修正方法は、判断の前にまず簡約後の extent を見ることであり、定数境界の場合は自身に簡約され、元と同じになる。この欠陥は下流の TileOPs#1852 から報告され、afcebed1 以降に二分され、破壊対象はページングデコード——その行程数は KV 長さに由来する。作者は sm_89、CUDA 13.0 上で両版をコンパイル比較し、新規テストは主干で失敗し、本コミットで通過する。
出典:https://github.com/tile-ai/tilelang/pull/3356
1.8 メインリポジトリ:スライス填充と負インデックスの静的境界を二件修正、越界書き込みと静黙負アドレスを閉鎖(#3352 / #3347)
10-27 05:37 と 05:14 にマージ、合計 4 ファイル +100/-3。前者は Fill コンストラクタの静的境界チェックに対するものである。従来は次元ごとに min >= 0 と extent <= shape のみを検証していたが、extent <= shape はバッファ起点から始まる領域にのみ成立する。起点が内側によるスライスは両項目とも通過するが、越界書き込みに下沉される——T.fill(As[6:12], 7) は長さ 8 のバッファ上でコンパイルを通過し、As[threadIdx.x + 6] = 7 を発射し、6 から 11 に書き込む。修正方法は、領域の末端も目標形状に対して検証することであり、新条件は起点と長さの両方が静的であることを要求し、ゼロに锚定された領域と符号付き起点の領域は影響を受けない。T.clear は同じノードを共有しており、欠陥も同じため併せて閉鎖する。後者は LegalizeNegativeIndex に対するものである。従来は負インデックスにバッファ長を一度加えるだけで済ませ、範囲チェックがなく、長さ 8 のバッファ上で A[-9] は -1、A[-10] は -2 となり、割り当て起点を外れ、いかなる診断もない。一方 NumPy と PyTorch は同じインデックスに対して IndexError を投げる。修正方法は、アナライザが回繞後もなお負であることを証明できる場合にのみ拒否し、インデックスと長さを報告することであり、範囲内のすべての定数とすべての動的インデックスは元の動作を維持する。
出典:https://github.com/tile-ai/tilelang/pull/3352、https://github.com/tile-ai/tilelang/pull/3347
1.9 メインリポジトリ:T.ws(-1) 領域が全スレッド上で実行、取り残しの記法がもはや保存を失わない(#3354)
10-27 04:52 にマージ、2 ファイル +40/-3、#2978 を引き継ぐ。WarpSpecialize のドキュメントでは -1 を「threadIdx.x バインディングをスキップする」と定義しており、つまりその領域は threadIdx の制約を受けない。しかし実装では -1 を通常の warp グループインデックスとして扱い、[-warp_group_size, 0) というどのスレッドにも成立しない区間を構築してしまう。その結果、領域はゼロ個のスレッド上で実行され、区間内のすべてのストアが破棄され、診断も出ない。最小再現は明快だ。256 スレッドでまず全て 0 を書き、次に T.ws(-1) の中で 5 を書くと、結果として 256 スレッドが書いた値は依然として 0 のままである。同じストアを ws でラップしなければ、256 スレッドすべてが 5 を書く。修正方法は、id に -1 が含まれる場合に範囲制約を一切発行せず、実際のグループインデックスは自身のスレッド制約に従わせ、ドキュメント文字列に -1 の例を補うことである——元のドキュメントの例にはちょうどそれが欠けていた。作者は全リポジトリを調査して ws(-1) の用法が現れないことを確認しており、したがって既存コードがその空領域に依存していることはない。
情報源:https://github.com/tile-ai/tilelang/pull/3354
1.10 メインリポジトリ:profiler がデバイス未指定時に現在のデバイスを使用、キャッシュフラッシュが 0 番カードに当たらなくなる(#3357)
10-27 07:26 にマージ、1 ファイル +1/-1。do_bench はデバイスを明示指定せず現在の CUDA デバイスが非ゼロの場合、_cache_device(None) がモジュールレベルのデフォルト cuda:0 に落ちるため、キャッシュフラッシュのための割り当てが実測ワークと異なる GPU に乗る可能性があり、測定された数値と実際の負荷が同じカード上にないことになる。修正方法は、デフォルト文字列をインデックスなしの “cuda” に変更し、PyTorch がバッファ割り当て時に現在のデバイスを解決するようにすることである。明示指定デバイスならびに NPU、MPS のデフォルト値は変更しない。上流 CI(run 37915443170)は Quick Lint、CUDA、ROCm、Metal、Ascend および CuTeDSL の例で通過した。作者は、後続の PR トリガー式実行にはメンテナの承認が必要であり、まだテストを実行しておらず、ローカルにも CUDA ハードウェアや計時ベンチマークがないことを注記している。
情報源:https://github.com/tile-ai/tilelang/pull/3357
1.11 新規欠陥:カーネルキャッシュのヒットが冷間コンパイルを仮定する 6 つのテストを破壊;CuTeDSL の CUDA graph キャプチャが空グラフになる(#3694 / #3693)
両 issue はいずれも lucifer1004 が 10-27 に起票した。前者は、6 つのテストが冷間キャッシュ下では通過するが、同じキャッシュディレクトリで再実行すると失敗することを指摘している。原因は、それらがコンパイルが必ず発生すると仮定していることである。例えば test_tilelang_kernel_gemm_sm75.py の sm75 境界ケースでは 4 つのパラメータで AttributeError: ‘NoneType’ object … が報告される。後者はさらに厳しく、CuTeDSL バックエンドから発行されたカーネル起動が CUDA graph に記録されず、torch.cuda.CUDAGraph のキャプチャが空グラフを取得し、replay() が何も実行しない。一方、同じカーネルを cuda バックエンドでコンパイルするとキャプチャ・リプレイは正常であり、報告には最小再現が添付されている。両者とも「キャッシュと実行コンテキスト」という一族の欠陥の延長であり、1.4、1.5 の 2 件のマージが指し示す同一領域と符合する。
情報源:https://github.com/tile-ai/tilelang/issues/3694、https://github.com/tile-ai/tilelang/issues/3693
2. マルチバックエンド適応(昇騰 / 海光 / ROCm / Hexagon)
2.1 海光:HCU scale バッファをデュアル bank に区分、blockscaled GEMM が手動 warp 特化下で行マッピングを共用(#29)
10-27 08:26 にマージ、12 ファイル +1389/-314、本調査期間のソフトウェア側で最大の一件。変更は scale バッファの行を独立したデュアル bank 区分として計画し、LDS bank と分離するとともに、同一の行マッピングを割り当て、scale コピー、MMAC 沈下の 3 箇所に同時に渡し、3 箇所がそれぞれレイアウトを推し直すことを避ける。詳細には、物理区間化されたライフタイムと m0 アラインメントページングの保持、scale コピーの発行条数の圧縮、アイドル wave の作業を 2 つの scale bank に分散、このレイアウトを手動 warp 特化下の blockscaled GEMM に使用することが含まれる。この一件は昨日マージされたネイティブ B16 ストレージと動的 WASP GEMM(#28)を引き継ぎ、海光は 2 日間で 3 件の主干マージを完了した。
情報源:https://github.com/tile-ai/tilelang-hygon/pull/29
2.2 昇騰メインリポジトリ:PerCoreTask 候補体が内部 if/else を許可、条件付き L1 フィルが全体拒否されなくなる(#3686)
10-27 07:24 にマージ、2 ファイル +85/-24。T.PerCoreTask 内の T.Task は、本体内に if/else が含まれるだけで拒否される。たとえ分岐が内部の作業のみを制御していても——例えば一度のコピー後に条件に応じて L1 フィルを補う場合でも——同様である。d85a70a9 で再現し、具現化段階で「選択候補の制御フローは task スコープの外に置かなければならない」と報告される。本コミットはこの一律の拒否を削除し、task 本体に対する統一的な検証に改める。同期フェンスが候補全体を包み、同時に各分岐は引き続き項目ごとにパイプラインとコアの互換性、ネストされた task、およびサポートされていないクロスコア同期をチェックする。回帰テストは内部 if と if/else を受け入れ、候補本体を保持し、既存の拒否ケースも保持する。さらに完了フラグが外側のガード下で候補全体に追随することをチェックする。task、スケジューリングユニット、遅延の三セット計 113 例がすべて通過し、条件付き L1 フィルを持つカーネルも完全な昇騰ソースコードの沈下を完了できる。デバイス実行は未テスト。
出典:https://github.com/tile-ai/tilelang/pull/3686
2.3 昇騰メインリポジトリ:CI が AscendC と PTO を二つのマトリクスジョブに分割し、個別に再実行可能に(#3687)
10-27 07:05 にマージ、7 ファイル +56/-27。従来、AscendC と PTO は同一の Test for Ascend ジョブ内の二つのステップであり、いずれかのバックエンド固有の失敗がチェック全体を赤にし、再実行は必ずもう一方のセットのサンプルを巻き込んでいた。本コミットはサンプルと単体テストを二つのマトリクスジョブに分割する。Test for AscendC は examples/ascend(target は ascend、4 並列)と testing/ascend -m “not pto”(32 並列)を実行し、Test for PTO は examples/ascend(target は pto)と testing/ascend -m pto(32 並列)を実行する。分割は既存の pto マークに依存し、PTO のみを実行するテストと両バックエンドテスト内の PTO パラメータにマークを补う。二つのマーク式は相互排他かつ合わせて全套をカバーし、重複実行しない。fail-fast は false に設定し、一方のジョブの失敗が他方をキャンセルしない。両方とも「サンプル失敗でもビルド成功時は単体テストを実行する」動作を保持し、それぞれがビルドとテストのサマリーを発行する。共用テスト(同一フロントエンド IR をそれぞれ二つのコード生成器に通す契約チェックを含む)は AscendC 側でのみ一度実行し、その約定を testing/ascend/README.md に記す。
出典:https://github.com/tile-ai/tilelang/pull/3687
2.4 昇騰リポジトリ審査中:Cython ラッパーがアナライザを遅延構築、昨日記録した呼び出しごとの再構築問題に直撃(#1948)
10-27 07:42 更新、作者 zwh1025。昨日のデイリーレポートで記録した新欠陥 #1945 は、昇騰リポジトリの Cython ラッパーが呼び出しごとに新しい TVM アナライザを生成し、実測約 0.28 ミリ秒の固定オーバーヘッドがあると指摘した。本コミットは「TVM アナライザの遅延構築」を題し、このオーバーヘッドに正対する。こうした小さなオーバーヘッドは大規模モデル推論の層ごとの呼び出しで観測可能な遅延に累積するため、昇騰側の性能掃除における低垂の果実に属する。現在も審査中。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1948
2.5 昇騰リポジトリ審査中:GLM-5.3 の kpool 圧縮オペレータ、PTO 上のテンソルシフト量(#1910 / #1949)
いずれも 10-27 に更新。#1910 は昇騰 AI プロセッサに GLM-5.3 の kpool 圧縮オペレータのサンプルを追加する——これは具体モデルの構造部品をサンプルリポジトリに持ち込む手法であり、TileOPs 側の「オペレータはモデルの需要に応じて育つ」という考え方と一致する。#1949 は T.tile.bitwise_lshift が PTO 上でテンソル形式のシフト量をサポートするようにする。従来は定数のみを受け付けていた。PTO は昇騰側で比較的独立したコード生成パスであり、シフト量のテンソル化は、このパスがマスクやビットパッキング系のオペレータにおいて汎用表現へ一歩近づいたことを意味する。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1910、https://github.com/tile-ai/tilelang-ascend/pull/1949
2.6 昇騰メインリポジトリ審査中:PTO の T.Parallel ベクトル化、AutoSimVF チャネル、バッチ三次元 ND2NZ、書き込みカウント式マルチバッファ、Bisheng ビルド識別子のキャッシュキーへの組み込み(#3512 / #3320 / #3689 / #3629 / #3691)
審査中の五件は昇騰側の本調査期間の待ち行列を構成し、いずれも 10-27 に更新がある。#3512 は PTO 上で T.Parallel をベクトル化するパイプライン枠組みと要素ごとのサポートであり、構造的な変更に属する。#3320 は AutoSimtVF チャネルを新設し昇騰パイプラインに登録する。方向性としては SIMT とベクトルの間で自動的に経路を選ぶものである。#3689 はバッチ三次元の GM から L1 への ND2NZ コピーをサポートし、バッチアテンションとタイル GEMM の転送需要に直接対応する。#3629 は書き込みカウント式マルチバッファを追加し DeepGEMM のロードを再配置する。#3691 は Bisheng コンパイラのビルド識別子をバイナリキャッシュキーに組み込む。1.5 の compile_flags のキャッシュキーへの組み込みと同種の考え方であり——昇騰側にもコンパイラ変更後に古いバイナリへ誤ヒットする問題が同様に存在する。
ソース:https://github.com/tile-ai/tilelang/pull/3512、https://github.com/tile-ai/tilelang/pull/3320、https://github.com/tile-ai/tilelang/pull/3689、https://github.com/tile-ai/tilelang/pull/3629、https://github.com/tile-ai/tilelang/pull/3691
2.7 ROCm 審査中:pip 版 ROCm 10 のアーキテクチャ検出とイメージ、FP4 変換と E8M0 型(#3549 / #3665)
#3549 は pip インストール版 ROCm 10 SDK をサポートし、変更にはアーキテクチャ検出と Dockerfile が含まれる。これは 1.2 の auto ターゲット修正と関連するが異なる関数に属する——本件が変更するのは /opt/rocm ベースの get_rocm_arch 検出であり、1.2 が変更するのは tilelang/rocm/target.py 内の検出ヘルパーである。#3665 は HIP テンプレート内の浮動小数点から FP4 への変換を修正し、E8M0 型サポートを補完するもので、ブロック単位スケーリングと低精度量子化の ROCm 上での可用性を指向している。いずれも審査中である。
ソース:https://github.com/tile-ai/tilelang/pull/3549、https://github.com/tile-ai/tilelang/pull/3665
2.8 昇騰リポジトリのデイリーテスト:2903 項目すべて合格、2 日連続で失敗なし(#1950)
昇騰リポジトリのデイリーテストボットによる 10-28 06:49 のレポートでは、合計 2903 項目、合格 2903、合格率 100%、ワークフロー実行番号 40590261193 であった。前日のレポート(#1947)も同樣に 2903 項目すべて合格である。現在も急速に進化を続ける NPU バックエンドにとって、2 日間失敗なしのデイリーテストは主干の可用性が安定していることを示し、これは 2.3 の CI 分割にベースラインを提供する——分割後も 2 つのジョブがそれぞれこの數字を維持すべきである。
ソース:https://github.com/tile-ai/tilelang-ascend/issues/1950
2.9 沐曦:組織ビューにプッシュタイムスタンプあり、6 ブランチいずれも調査期間内の新規コミットなし(10-27)
tilelang-metax リポジトリの pushed_at は 10-26 07:08 を示し、本調査期間の起点より前である。6 ブランチを逐一確認したところ、調査期間内に新規コミットはなく、実質的変更なしとして処理し、項目レベルの内容は記載しない。
ソース:https://github.com/tile-ai/tilelang-metax
2.10 サードパーティ:Hexagon DSP バックエンドが 1 日で 15 件、HVX ベクトル化、VTCM アリーナ、Crouton レイアウト、DMA 組み込みが揃う(10-27)
xwhzz/tilelang-hexagon の hexagon-backend ブランチは 10-27 の 1 日で 15 件のコミットを着地させ、本調査期間のバックエンド側で最も密度の高い外部の動きとなった。当日の主线は 3 段階に分けられる。下層では CodeGenTileLangCBase からコード生成器を再構築し、DMA をコミットグループ付きの IR 組み込みとし、起動グリッドとワーカースレッドプールを TIR に落とし込み、TIR 内で VTCM アリーナを計画した。中層では HVX の単一ベクトル化経路(算術マッピング、ベクトル數学に加え 1 本の要素単位パス)を敷設し、HVX 転置木による複数行リダクション、ソフトウェアパイプライン化された HVX マッピングを実現し、フラグメントと T.reduce_* をサポートした。上層では Crouton レイアウトと領域コピーを導入し、HVX マッピングとリダクションを Crouton タイル上で直接実行させ、セッションごとに DSP 周波数へ投票し、最後にビルドキャッシュと並列デバイステストを追加した。コミットには破壊的マーカー(feat(hexagon)!)が付いており、Crouton レイアウトと領域コピーが既存インターフェースを変更したことを示す。Hexagon は携帯電話と組み込み側の DSP であり、この路線が通れば、TileLang のハードウェアカバレッジは GPU、NPU から DSP へと拡張される。
ソース:https://github.com/xwhzz/tilelang-hexagon
3. エコシステムと採用側(TileOPs / TileFoundry)
3.1 TileOPs 0.1.0 リリース:206 オペレータの初の正式版、公開 API がここに定まる(10-27)
10-27 13:15 に v0.1.0 を公開、リポジトリ初の GitHub release である。PyPI 上の 0.1.0 成果物のアップロードはさらに早く、10-27 07:45(北京時間、すなわち 10-26 23:45 UTC)であり、前回デイリーレポートの締め切り時点より 45 分遅い——これはまさに昨日のデイリーレポートで「初回リリースはまだ最後の枠内」と書いた直後に起きた変化である。リリースノートの位置づけは明確だ。agent 時代のために作り直されたオペレータライブラリであり、各オペレータはまず YAML 仕様(シグネチャ、テストとベンチマーク用の負荷、roofline 式を宣言)を書き、コーディング agent がカーネルを開発・チューニングし、仕様から派生したテストとベンチマークがどの実装がライブラリに入れるかを決める。195 件の foundry タグ付きマージはこのようにして取り込まれた。本バージョンでは公開 API を固定する。オペレータは tileops.
出典:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0、https://pypi.org/project/tileops/
3.2 TileOPs 0.1.0 ベンチマーク:H200 上で 2265 行のうち 2262 行が最快実装より遅くなく、幾何平均 1.37 倍(10-27)
夜間ベンチマークは各 TileOPs カーネルを同一オペレータの他のすべての実装と同一 GPU 上で比較し、単行の比率は最快の比較対象の所要時間を TileOPs の所要時間で割ったもので、1 より大きい場合はより高速であることを示す。全量 2265 行、16 のオペレータファミリで、幾何平均 1.37 倍、2262 行が 0.99 以上に達した。ファミリ別に見ると差異は顕著である。sequence_modeling 4.00 倍、sampling 2.95 倍、convolution 1.86 倍、mamba 1.61 倍、linear_attention 1.54 倍、pool 1.51 倍、quantization 1.45 倍、attention 1.42 倍;体量最大の elementwise 527 行はわずか 1.09 倍、gemm 235 行は 1.28 倍。比較対象には torch、torch.compile、FlashAttention 2 と 3、FlashInfer、FlashMLA、vLLM、DeepGEMM、flash-linear-attention、FlagGems、QuACK、Mamba、Marlin と cuBLASLt が含まれる。0.99 を下回った 3 行はそれぞれ 16M float32 要素上の Where と ClampScalar、および DeepSeek-V3 プリフィル出力投影上の FP8 GEMM であり、差は 1.1% から 1.7% である。この分布は、收益が主にロングテールのオペレータファミリから来ており、主流の GEMM と要素ごと演算ではないことを示す——後者はもともと各実装が最も激しく競争している領域である。
出典:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0
3.3 TileOPs 0.1.0 の既知のギャップと次のステップ:融合オペレータ、モデル全体、計算通信融合(10-27)
リリースノートは 4 つの既知の問題を列挙している。FusedAddLayerNormFwdOp がゼロから遠い行で精度を失い、1e36 に近い行でオーバーフローする(#2773、昨日のデイリーレポートでこの欠陥は既に記録済み);同一ホスト上に異なるアーキテクチャの GPU がある場合、ある GPU 用にコンパイルされたカーネルが別の GPU 上で動作する可能性がある(#2302);Bwd 系オペレータは直接呼び出す必要があり、torch.autograd に接続されない;単一のオペレータインスタンスは同時に 1 つのスレッドと 1 つの CUDA ストリームのみをサポートする。検証範囲は H200、Python 3.12 と 3.13 のみをカバーする;一部のカーネルには SM80、SM86、SM89 パスがあり、これらのパスは対応するデバイス上で動作するがリリース検証範囲には含まれない。ロードマップは 3 ステップである。まずオペレータ境界を跨ぐ複雑な融合(アテンションと前後の投影および正規化、層全体の MoE、GEMM epilogue 内の量子化)を行い、次にモデル全体のプリフィルとデコードを行いサービスエンジンとエンドツーエンドで比較し、最後に計算と通信の融合(テンソル並列 GEMM と all-reduce、エキスパート並列 MoE と all-to-all)を行う。この 3 ステップが実現すれば、TileOPs の對手は単一オペレータライブラリから推論フレームワークに変わる。
出典:https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0
3.4 TileOPs に4件をマージ:where と masked_fill がチューニング可能に、D128 因果対角を64キー単位でスコアリング、疎な長 GDN プレフィルを分割、変更ログの日付をリリース日に設定(#2782 / #2783 / #2779 / #2781)
4件は10-27にマージされ、まさにリリース前の締めくくりとなった。#2782 は request_tune() が WhereFwdOp、MaskedFillTensorFwdOp、MaskedFillScalarFwdOp で ValueError を投げる問題を修正:自動チューニングはランダム生成された整数入力を受け付けないが、これら3つのカーネルは bool 条件やマスクを uint8 として渡していた。ランダムセレクタは各要素がどのオペランドを取るかだけを変え、候補の作業量は変えないため、ランダム整数入力を安全と宣言;3つのテストノードを追加し、主干では3例すべて失敗した。#2783 は D128 密プログラムの因果対角を最適化:因果呼び出しが1つの key セグメントのみで、両方の長さが128行境界上にある場合、最初に消費する warpgroup の対角 key タイルは全行に対して上位64キーをマスクするため、このステップを m64n64 の QK と64キー softmax で実行するよう変更;この選択はコンパイル期フラグ(diag_aligned、カーネルキーに含まれる)である。実行期に64と128幅の QK 間で分岐すると、ptxas が一部の特化で各 WGMMA を直列化する(C7511)ため、フラグに切り替えた後は manifest 内のすべての D128 プレフィル特化で C75xx が出なくなった。#2779 は SM89 上の疎な長系列 GDN プレフィルを状態スキャンと独立出力プログラムに分割し、密で key 主順の D128 入力かつ少なくとも16384トークン、状態スキャン CTA が SM 数の4分の1以下の場合に適用;(1, 16384, 4, 128) bf16 のベンチマークでは所要時間が約0.612ミリ秒から約0.470ミリ秒に短縮され、flash-linear-attention の0.441ミリ秒に対して0.939倍となった。#2781 は CHANGELOG の0.1.0の日付を PyPI に登録された日(2026-10-27)に定め、コントリビューター文書からリリースごとの設定と Actions のクリック手順を削除した。
出典:https://github.com/tile-ai/TileOPs/pull/2782、https://github.com/tile-ai/TileOPs/pull/2783、https://github.com/tile-ai/TileOPs/pull/2779、https://github.com/tile-ai/TileOPs/pull/2781
3.5 TileFoundry に2件をマージ:ウィンドウ全体をオペランド自体として発行、動的共有メモリをカーネル実際使用量でサイズ決定(#273 / #271)
2件とも10-27にマージされた。#273 は sm_89 GEMM の沈み込みを修正:schedule finalize が書き出した内容がコンパイルできない。沈み込みがアキュムレータ全体をゼロ開始・自身と同一スパンのウィンドウ T.ptr_of(acc[0:0+16, 0:0+8]) で包むが、CUDA codegen はフラグメントレジスタタイルのポインタスライスを拒否し「通常のソースレイアウトが必要」と報告;MMA もレジスタフラグメントを lhs_frag_dst[0:0+16, 0:0+32] のような裸の Slice として読み、codegen は同様にハンドラを持たない。修正は沈み込み側で行う:StoragePlan.operand はもともと「ゼロ開始かつオペランド全体にまたがるウィンドウ」をオペランド自体として扱うが、実体化と命令発行は依然として Slice を構築していた。これを一度だけ判定するよう変更——MLIR の TrivialSubViewOpFolder と同じ置換を行う。#271 は動的共有メモリを修正:合成ホストエントリが各 TIR カーネルを dynamic_smem=0 で起動するため、整数 TensorView オフセット経由で共有メモリにアクセスするカーネルに領域が割り当てられず、48 KiB を超える場合は cudaFuncAttributeMaxDynamicSharedMemorySize の設定も必要だがランチャーは一切設定していなかった。作者は #267 の「ターゲット容量どおりに満額割り当てる」提案を採用せず、カーネル実際使用量でサイズ決定する方式に変更。理由は静的共有メモリと動的共有メモリが同じ上限を共有し、各ブロックの申請量が1つの SM に何ブロック載るかを決めるため——sm_89 で99 KiB を満額にすると512バイト使用のカーネルでも1ブロックしか載らなくなる。CUTLASS と TileLang も同じ方式でサイズ決定している。
出典:https://github.com/tile-ai/TileFoundry/pull/273、https://github.com/tile-ai/TileFoundry/pull/271
3.6 外部エコシステム:CANN コミュニティの TileLang オペレータリポジトリ、PPU 版ゼロ変更ラッパー、DeepSeek V4.1 Flash フラグメント例(10-27)
3つの外部リポジトリはいずれも調査期間内に動きがあり、合わせて見ると「昇騰側オペレータ資産」の形成過程が一筋の流れとして浮かび上がる。Atlas-Ascend-Fork/ops-tilelang は CANN コミュニティが昇騰向けに提供する TileLang オペレータリポジトリで、10-20 にリポジトリを開設し、10-27 に「増分テストの自動発見とオペレータマッピングテーブルの削除」をコミットした。README は Ascend 950 以降のアーキテクチャのみをサポートし、CANN 9.3.0 以上を要求し、コミュニティの weekly 版を使用することを明記しており、設計原則はインターフェースとスケジューリングの分離、業務ドメインごとの集約、正確性を性能より優先、benchmark とテストの分離、そしてベースラインをテストファイルとチップクラスタごとにシャーディングして保存することにある。kevin3113/models-tilelang はこのリポジトリの7つのオペレータで DeepSeek V4.1 Flash の「Attention 出力+FFN 前処理」断片(mhc_post_forward、二乗和付き GEMM、リダクション付き RMSNorm、mix 分割、sinkhorn、apply_mix、ffn_norm)を連結しており、HIDDEN=5120、MHC=4、デフォルトの decode 形態は SEQ_LEN=1 で、SuperKernel と npugraph_ex の ACL graph をサポートし、B を 1 から 256 まで実測した SuperKernel の効果は 1.17% から 2.75% である。flytiger-eco/tilelang-wrapper-for-sail は別方向の移行を解決する。CUDA スタイルの TileLang コードを一切変更せずに tilelang-sail(PPU バックエンドディストリビューション版)上で動作させるもので、その手法は純粋に実行時に import プロトコルと site メカニズムをインターセプトするもので、上流ソースコードを変更せず再パッケージングもせず、PPU に対応物がない CUDA 機能はサイレントなデグレードではなく明示的にエラーを報告する。三者はそれぞれオペレータ供給、モデル実装、バックエンド移行の三つの段階に対応している。
出典:https://github.com/Atlas-Ascend-Fork/ops-tilelang、https://github.com/kevin3113/models-tilelang、https://github.com/flytiger-eco/tilelang-wrapper-for-sail
4. コミュニティ、チュートリアル、イベント
4.1 ドキュメントサイトが1日で6件の自動更新(10-27)
tilelang.github.io は調査期間内に6回の Update docs コミット(05:20、06:22、07:19、09:28、22:42、22:43)があり、いずれもボットによってプッシュされ、内容はメインリポジトリのマージ後のドキュメント再生成によるものである。リリース当日のドキュメントサイトの更新頻度は主干のマージリズムと同期しており、正常な連動であって個別に内容を列挙する必要はない。
出典:https://github.com/tile-ai/tilelang.github.io
4.2 昇騰リポジトリのドキュメント・ファクトチェック4件、単一コメント最大46スレッド(#1628 / #1655 / #1656 / #1659)
昇騰リポジトリは 10-27 に「ファクトチェック」を題したドキュメント PR が4件更新された。T.tile.clear(#1628)、T.tile.bitwise_or(#1655)、T.tile.bitwise_not(#1656)、T.tile.bitwise_lshift(#1659)で、コメント数はそれぞれ 10、46、18、20 であり、うち3件にはボットのマークが付いている。1回のチェックが46コメントまで膨らむということは、議論されているのが単なる文言だけではないことを示している。この種の API ごとの動作確認は、実際には昇騰側の意味論的取り決めを記録に留めるためのものである。新しいバックエンドラインにとって、ドキュメントと実装の整合にかかるコストはしばしば過小評価されるが、この一群の PR は目に見える投入である。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1655、https://github.com/tile-ai/tilelang-ascend/pull/1656
4.3 審査中:lower trace を独立した開発ツールキットとして抽出(#3583)
10-27 に更新されたこの審査中 PR は、lower trace をメインリポジトリから tilelang-devkit へ抽出するものである。ダウンストリームプロセスの追跡情報はバックエンド適応のデバッグに非常に有用だが、メインリポジトリに置くとコンパイルパスと依存関係の負担が増える。独立したツールキットとして抽出した後は、メインリポジトリは軽量を保ち、必要に応じて後から接続できる。この種の「デバッグ機能の外部化」の変更は、通常、プロジェクトが複数のダウンストリームバックエンドを持ち始めた時期に現れる。本調査期間内には Hexagon、PPU、昇騰の3ラインが同時に動いており、まさにそうしたタイミングである。
出典:https://github.com/tile-ai/tilelang/pull/3583
4.4 ニュース側:Google News と HN は24時間以内に本調査期間の新規報道なし(10-27)
中英4組のクエリ(TileLang、TileOPs、TileLang オペレータ、tile-ai 智算 kernel)が返した項目はすべて9月末から10月中旬に集中している。1つはDeepSeekとファーウェイが9月29日から30日にAscend版TileLangをオープンソース化したという報道(SCMP、Reuters、Tom’s Hardware、雷峰網、InfoQなど)で、最新の1件は10月22日のSemiAnalysisによるファーウェイとDeepSeekの共生関係に関する分析であり、いずれも調査期間外である。TileOPsの完全一致クエリはゼロヒット。Hacker Newsでは期間内にGameMakerの再実装に関する項目が1件のみあり、ノイズとして除外した。arXiv上のTileLangキーワードの最新論文は2026-07-24。本日の内容の重心はすべてGitHub側にあり、これはリリース日のタイミングと一致している——正式リリースはまだ報道サイクルに入っていない。
5. トレンド観察
第1に、TileOPs 0.1.0は「オペレータライブラリ」のコスト構造を別のアルゴリズムに置き換えた。従来の方法ではオペレータ数がハードウェアとモデル需要に応じて線形に増加し、人手がボトルネックになる。ここでの方法は、各オペレータについてまずYAML仕様を1份書き、シグネチャ、ペイロード、rooflineを明確に宣言し、次にagentがカーネルを生成・チューニングし、最後に仕様から派生したテストとベンチマークで准入を行う。206個のオペレータ、195件のfoundryタグのマージがこのフローの初回全量納品である。これが持続可能かどうかは2つのことにかかっている。仕様が複雑な融合オペレータをカバーできるか(ロードマップの第一歩がまさにこれ)、そしてベンチマークが「ベンチマークに合わせてチューニングされたが実際にはロバストでない」カーネルを阻止できるか——FusedAddLayerNormFwdOpがゼロから遠い行で精度を失うという既知の問題は、まさにこの種のリスクのサンプルである。
第2に、本期間のマージはほぼすべて「サイレントエラー」である。ROCmの数学関数がデフォルトで高速ビルトインを使うことでsin/cosが3桁の差を生む。address_ofが読み取りと見なされることで不透明呼び出し後の共有メモリが書き込み前の内容を読む。T.ws(-1)の領域がゼロスレッドで実行され、保存がすべて破棄されかつエラーが報告されない。スライス填充が長さではなく末端で検証され、直接範囲外書き込みとなる。負のインデックスが回り込んだ後もなお負であり、割り当て開始点より前に落ちる。これら5件の共通点は、コンパイルが通り、実行時にエラーが出ず、結果が誤っていることであり、しかもすべて直近2週間以内に集中的に発見されている。大規模モデルのためにカーネルを書く言語にとって、この種の欠陥のコストはコンパイル期のエラーよりはるかに高く、それらは同じ工学的要請を指している——「エラーは出ないが動作が正しくない」書き方をコンパイル期に拒否すること。本期間の境界検証の2件の修正とws(-1)の修正が行っているのはまさにこれである。
第3に、バックエンドのカバレッジは「ベンダー主導」から「コミュニティ自発」へ移りつつある。公式に保守されているAscend、海光、沐曦、摩尔線程の4ライン之外で、Hexagon DSPバックエンドが1日で15件をマージし線形代数プリミティブをCroutonタイルに展開し、PPU版ではゼロ変更ラッパーが現れ、CANNコミュニティが独自にオペレータリポジトリを構築しそれを使ってDeepSeek V4.1 Flashのモデル断片を組み上げた。これらの動きはいずれもtile-ai組織のリポジトリにはないが、いずれもTileLangの言語層の安定性に依存している。これこそが0.1.0で公開APIを凍結し、API互換性チェックをCIゲートにした意義である——下流が多くなるほど、後方互換性の価値は高まる。
第4に、キャッシュの同一性という問題に2つのバックエンドが同時に着目した。メインリポジトリはcompile_flagsを自動チューニングのキャッシュキーに組み込み(#3338)、AscendはBishengのビルド同一性をバイナリキャッシュキーに組み込んだ(#3691)。同じカーネルでも異なるコンパイルオプションや異なるコンパイラでは異なるバイナリが生成され、キャッシュをソースコードと形状だけでキーにすると誤ヒットし、測定された性能が実際のコンパイル結果と合わなくなる。2つの独立した路線が同じ日に同じことを行ったことは、この種の問題が実用上すでに十分に痛みを伴うことを示している。
付、素材と検証説明
情報源検証表
| 情報源 | 確認結果 |
|---|---|
| GitHub tile-ai 組織(29 リポジトリ) | 期間内に 7 リポジトリへプッシュ:tilelang、TileOPs、TileOPs-nightly、TileOPs.github.io、tilelang-hygon、TileFoundry、tilelang.github.io |
| マージ済み真値(is:merged) | 20 件、1 件ずつマージ時刻が期間内であることを確認(tilelang 13、TileOPs 4、TileFoundry 2、tilelang-hygon 1) |
| 審査中真値(is:open 更新) | 45 件、更新時刻順に全件確認 |
| 新規 issue | 7 件、本文を 1 件ずつ確認、うち 2 件は昇騰リポジトリの毎日テストボットレポート |
| リリースページと PyPI | TileOPs v0.1.0 が 10-27 13:15 に公開、PyPI 成果物のアップロードは 10-27 07:45;tilelang の最新は依然 0.1.15(09-30)、tilefoundry の最新は 0.0.4rc1(10-25) |
| tilelang-ascend リポジトリ | 期間内にコミットなし、ただし PR と毎日テスト issue の活動が密接 |
| tilelang-metax リポジトリ | 組織ビューにプッシュ時刻のタイムスタンプあり、6 ブランチはいずれも期間内コミットなし、実質的変更なしとして処理 |
| tilelang-musa リポジトリ | 前回プッシュ 09-30、期間内にコミットなし |
| 採用側(TileKernels / FlashQLA) | いずれも期間内プッシュなし(10-14) |
| 外部バックエンドリポジトリ | Hexagon バックエンド 15 件のコミット;CANN オペレータリポジトリ、PPU ラッパー、モデル断片リポジトリはいずれも期間内の動きあり、2.10 と 3.6 を参照 |
| Google News 中国語・英語の複数クエリ | 精確クエリはゼロヒット、返却された項目はすべて 9 月末から 10 月中旬に落在し、いずれも期間外 |
| Hacker News | 期間内に GameMaker 再実装の項目が 1 件、TileLang とは無関係、除外 |
| arXiv | キーワード検索で期間内の新規論文なし、最近の 1 本は 2026-07-24 |
完全な情報源リスト
- [1] tilelang ROCm fastmath デフォルト修正(#3621)— https://github.com/tile-ai/tilelang/pull/3621
- [2] tilelang ROCm 自動ターゲットが現在のデバイスを読み取る(#3594)— https://github.com/tile-ai/tilelang/pull/3594
- [3] tilelang TMA gather/scatter 定数行インデックス(#3365)— https://github.com/tile-ai/tilelang/pull/3365
- [4] tilelang NVRTC と CuTeDSL デバイスストリーム(#3359)— https://github.com/tile-ai/tilelang/pull/3359
- [5] tilelang 自動チューニングキャッシュキーに compile_flags を組み込み(#3338)— https://github.com/tile-ai/tilelang/pull/3338
- [6] tilelang ThreadSync で address_of を書き込みとして記録(#3346)— https://github.com/tile-ai/tilelang/pull/3346
- [7] tilelang 実行時行程数パイプライン epilogue 展開(#3356)— https://github.com/tile-ai/tilelang/pull/3356
- [8] tilelang スライス填充境界検証(#3352)— https://github.com/tile-ai/tilelang/pull/3352
- [9] tilelang 範囲外負インデックスの拒否(#3347)— https://github.com/tile-ai/tilelang/pull/3347
- [10] tilelang T.ws(-1) 領域スレッドドメイン修正(#3354)— https://github.com/tile-ai/tilelang/pull/3354
- [11] tilelang profiler 現在のデバイス(#3357)— https://github.com/tile-ai/tilelang/pull/3357
- [12] tilelang カーネルキャッシュヒット破壊テスト(#3694)— https://github.com/tile-ai/tilelang/issues/3694
- [13] tilelang CuTeDSL CUDA graph 空グラフ(#3693)— https://github.com/tile-ai/tilelang/issues/3693
- [14] tilelang-hygon HCU scale バッファデュアルバンク区分(#29)— https://github.com/tile-ai/tilelang-hygon/pull/29
- [15] tilelang PerCoreTask 内部条件(#3686)— https://github.com/tile-ai/tilelang/pull/3686
- [16] tilelang 昇騰 CI で AscendC と PTO を分割(#3687)— https://github.com/tile-ai/tilelang/pull/3687
- [17] tilelang-ascend Cython ラッパー遅延アナライザ(#1948)— https://github.com/tile-ai/tilelang-ascend/pull/1948
- [18] tilelang-ascend GLM-5.3 kpool 圧縮オペレータ(#1910)— https://github.com/tile-ai/tilelang-ascend/pull/1910
- [19] tilelang-ascend PTO テンソルシフト量(#1949)— https://github.com/tile-ai/tilelang-ascend/pull/1949
- [20] tilelang PTO の T.Parallel ベクトル化(#3512)— https://github.com/tile-ai/tilelang/pull/3512
- [21] tilelang AutoSimVF チャネル(#3320)— https://github.com/tile-ai/tilelang/pull/3320
- [22] tilelang バッチ三次元 ND2NZ コピー(#3689)— https://github.com/tile-ai/tilelang/pull/3689
- [23] tilelang 書き込みカウントマルチバッファ(#3629)— https://github.com/tile-ai/tilelang/pull/3629
- [24]
tilelang Bisheng ビルドIDをキャッシュキーに追加(#3691)— https://github.com/tile-ai/tilelang/pull/3691
- [25] tilelang pip版 ROCm 10 サポート(#3549)— https://github.com/tile-ai/tilelang/pull/3549
- [26] tilelang HIP FP4 変換と E8M0(#3665)— https://github.com/tile-ai/tilelang/pull/3665
- [27] tilelang-ascend デイリーテスト 2903 項目全合格(#1950)— https://github.com/tile-ai/tilelang-ascend/issues/1950
- [28] tilelang-metax リポジトリ — https://github.com/tile-ai/tilelang-metax
- [29] tilelang-hexagon バックエンドリポジトリ — https://github.com/xwhzz/tilelang-hexagon
- [30] TileOPs v0.1.0 リリースページ — https://github.com/tile-ai/TileOPs/releases/tag/v0.1.0
- [31] PyPI TileOPs 成果物 — https://pypi.org/project/tileops/
- [32] TileOPs where と masked_fill のチューニング(#2782)— https://github.com/tile-ai/TileOPs/pull/2782
- [33] TileOPs D128 因果対角 64 キー(#2783)— https://github.com/tile-ai/TileOPs/pull/2783
- [34] TileOPs スパース長 GDN プリフィル(#2779)— https://github.com/tile-ai/TileOPs/pull/2779
- [35] TileOPs 変更ログの日付確定(#2781)— https://github.com/tile-ai/TileOPs/pull/2781
- [36] TileFoundry ウィンドウ全体をオペランド単位で発行(#273)— https://github.com/tile-ai/TileFoundry/pull/273
- [37] TileFoundry 動的共有メモリのサイズ確定(#271)— https://github.com/tile-ai/TileFoundry/pull/271
- [38] CANN コミュニティ ops-tilelang オペレータリポジトリ — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [39] DeepSeek V4.1 Flash スニペット例 — https://github.com/kevin3113/models-tilelang
- [40] tilelang-sail PPU 無変更ラッパー — https://github.com/flytiger-eco/tilelang-wrapper-for-sail
- [41] tilelang ドキュメントサイト — https://github.com/tile-ai/tilelang.github.io
- [42] tilelang-ascend ドキュメントファクトチェック(#1655)— https://github.com/tile-ai/tilelang-ascend/pull/1655
- [43] tilelang-ascend ドキュメントファクトチェック(#1656)— https://github.com/tile-ai/tilelang-ascend/pull/1656
- [44] tilelang lower trace の devkit 抽出(#3583)— https://github.com/tile-ai/tilelang/pull/3583
- [45] Google News 検索 — https://news.google.com/