TileLang デイリーレポート (2026-10-23)
調査期間:2026-10-22 07:00 ~ 2026-10-23 07:00(北京時間、24 時間)。
情報源:GitHub tile-ai 組織のリポジトリ別確認(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon ブランチ / tilelang-sunrise / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / 3 つのドキュメントサイト)+ マージ済み・審査中の真値検索(is:merged と is:open の 2 グループ、リポジトリ別にグループ化して統計)+ コミット単位の diff とリリースページの tag 確認 + 採用側リポジトリの確認(TileKernels / FlashQLA)+ コミュニティの新規リポジトリ検索 + Google News の中国語・英語複数グループ検索(プロキシ経由)+ Hacker News + arXiv + 産業メディア報道。
本期インデックス
- 本日の重点:35件のマージのうち、メインリポジトリはアナライザのライフサイクルに重点を置き、DeepSeek V3.2 の top-k セレクタの範囲外書き込みが CUDA_ERROR_ILLEGAL_ADDRESS に特定された。昇騰 NZ レイアウトが一部の UB→L1 コピーを実現(18ファイル)。CANN コミュニティが新たに ops-tilelang オペレータリポジトリを開設
-
- コアプロジェクトの進展
- 1.1 メインリポジトリ:ベクトル化書き換え前にアナライザをクローン(#3612)(10-22)
- 1.2 メインリポジトリ:並列ループ分割後に旧アナライザのバインドが無効化(#3226)(10-22)
- 1.3 メインリポジトリ:deepseek_v32 の top-k セレクタ候補バッファの範囲外(#3280)(10-22)
- 1.4 メインリポジトリ:CUDA 16ビット逆三角関数とバイナリ数学の降級を補完(#3172)(10-22)
- 1.5 メインリポジトリ:sm8x スパースメタデータの非全単射形状を拒否するよう変更(#3125)(10-22)
- 1.6 メインリポジトリ:eager フロントエンドが Python ソース位置を保持(#3558)(10-22)
- 1.7 審査中:pass 機能チェック、符号なし mulhi と assume 述語(10-22)
- コアプロジェクトの進展
-
- マルチバックエンド対応(昇騰 / 海光 / 沐曦 / 摩尔線程)
- 2.1 昇騰:NZ バッファが一部の UB→L1 コピーをサポート(#3530)(10-22)
- 2.2 昇騰 PTO:MAD 方向述語の指数増加によりコンパイルが停滞(#3581)(10-22)
- 2.3 昇騰 PTO:SIMT ベクトルのローカル変数をレーンごとに割り当て(#3523)(10-22)
- 2.4 昇騰リポジトリ:cann-bench に Dilation2D を追加、同日 depthwise_conv_2d をロールバック(#1916 / #1919)(10-22)
- 2.5 昇騰リポジトリ審査中:sparse_mla の順伝播と逆伝播(#1929 / #1928)(10-22)
- 2.6 海光:ブランチをマージして enable_register_pipeline を削除(10-22)
- 2.7 沐曦 / 摩尔線程 / MLIR-Ascend:調査期間内にプッシュなし(10-22)
- 2.8 産業側:SCMP が SemiAnalysis レポートを転載、TileLang は6リポジトリ中唯一利用可能(10-22)
- マルチバックエンド対応(昇騰 / 海光 / 沐曦 / 摩尔線程)
-
- エコシステムと採用者(TileOPs / TileFoundry / オペレータライブラリ)
- 3.1 TileOPs:FP8 ページド warp 特化 kernel の共有メモリエイリアスにより NaN(#2708)(10-22)
- 3.2 TileOPs:dense prefill を永続化走査に変更、D64 三コンシューマ(#2706)(10-22)
- 3.3 TileOPs:SM90 dense GQA decode が warp 特化を採用(#2694)(10-22)
- 3.4 TileOPs:因果スライディングウィンドウが1つの KV head の2つの key tile でクエリ head 群全体にサービス(#2696)(10-22)
- 3.5 TileOPs:ページド decode の自動チューニングが以前は1つの block table を共有(#2703)(10-22)
- 3.6 TileOPs:masked softmax の FMA オーバーフローと varlen RoPE の奇数半幅(#2699 / #2700)(10-22)
- 3.7 TileOPs:FP8 grouped GEMM の階層的 scale ステージングと gate/up 乒乓 warp group(#2704 / #2693)(10-22)
- 3.8 TileOPs:述語ブロードキャストとブールオペランドを16バイトベクトル読み込みに変更(#2702 / #2697)(10-22)
- 3.9 TileOPs:SM89 の3経路を同日推進(#2705 / #2691 / #2690)(10-22)
- 3.10 TileOPs:リダクション折り畳みの単一発行と空 DSA 選択の検証(#2689 / #2698)(10-22)
- 3.11 TileOPs:ベンチマーク参照を float32 上での計算に変更(#2692)(10-22)
- 3.12 TileFoundry:sm_89 アーキテクチャと RTX 6000 Ada デバイスのインストール(#259)(10-22)
- 3.13 TileFoundry:gmem の価格計算をフットプリントと実行 mesh に基づく方式に変更し、被呼び出し側の次元をバインド(#258)(10-22)
- 3.14 TileFoundry:デバイスが kernel の起動を拒否した場合は必ずエラーを送出(#256)(10-22)
- 3.15 ディストリビューション版:Sunrise が 0.1.15+sunrise.1.2.0 のリリース PR を作成(#9、審査中)(10-22)
- 3.16 採用者:TileKernels と FlashQLA は調査期間内にプッシュなし(10-22)
- エコシステムと採用者(TileOPs / TileFoundry / オペレータライブラリ)
-
- コミュニティ、チュートリアル、イベント
- 4.1 CANN コミュニティが ops-tilelang オペレータリポジトリを開設(10-22)
- 4.2 tilelang/agent:オペレータ最適化を agent の SKILL ワークフローとして記述(10-22)
- 4.3 コミュニティ新リポジトリ:Sail ラッパー、リモート自動チューニング agent、チュートリアルリポジトリ(10-22)
- 4.4 ドキュメントとバージョン:ドキュメントサイトを3回更新、VERSION は依然 0.1.
- コミュニティ、チュートリアル、イベント
15(10-22)
-
- トレンド観察
今日の重点
本調査期間中に tile-ai 組織は 35 件をマージした(TileOPs 20、tilelang 9、TileFoundry 3、tilelang-ascend 2、TileFoundry ドキュメントサイト 1)。さらに 71 件の審査中 PR が調査期間内に更新された。メインリポジトリの 9 件は、先週の「形状と幾何検証」からアナライザのライフサイクルへと重心が移っている。ベクトル化書き換え前にアナライザをクローンする(#3612、単一ファイル 2 行)、並列ループ分割後に旧バインディングが無効になる(#3226、6 ファイル)、eager フロントエンドが Python ソース位置を保持する(#3558)——この 3 件はいずれも kernel の計算を誤らせるものではなく、書き換え後にコンパイル期が陳腐化した状態を読むようにさせるものだ。数値側で最も記録に値するのは #3280 だ。deepseek_v32 の top-k セレクタは radix 候補を固定の [2, 4096] 共有バッファに一時保存するが、第一ラウンドで選出されるしきい値バケットのサイズを制約していない。低位仮数が同一の FP32 スコアや全等な入力では境界外書き込みが発生し、CUDA_ERROR_ILLEGAL_ADDRESS でクラッシュする。昇騰側の #3530 は当日最大の規模のコミット(18 ファイル +993/-256)で、圧縮 NZ パディングをストレージ正規化まで保持し、一部の UB→L1 コピーに対して C0 グループ行間隔を明示的に与える。#3581 は PTO 上でコンパイルを 10 分以上停滞させる述語式の指数増大を修正する。エコシステム側では 2 本の新ラインがある。CANN コミュニティは 10-20 に ops-tilelang オペレータリポジトリを設立し、10-22 に初回修正をプッシュした。Sunrise ディストリビューションは 0.1.15+sunrise.1.2.0 のリリース PR を打ち出し、TANG バックエンド wheel をディストリビューションに取り込んだ。
1. コアプロジェクトの進展
1.1 メインリポジトリ:ベクトル化書き換え前のアナライザクローン(#3612)
10-22 23:43 にマージ。ベクトル化の該当ソース 1 箇所のみを変更し、+2/-1。ループのベクトル化は IR を書き換える際に渡された算術アナライザを再利用するが、アナライザは書き換え前の定数と区間制約をキャッシュしている。ベクトル化がループ境界やインデックス式を変更すると、後続の pass が得るのは依然として古い結論となる。修正方法は、ベクトル化書き換えに入る前にアナライザを複製し、書き換えは複製に対してのみ作用させることだ。このコミットには回帰テストが添付されていないが、その価値は別の入口から #3226 と同源の問題を塞ぐ点にある——同日のそのコミットは並列ループ分割後の無効なバインディングを扱っており、2 件を合わせて見ると、分析コンテキストのライフサイクルが独立した欠陥クラスとして整理されていることがわかる。
出典:https://github.com/tile-ai/tilelang/pull/3612
1.2 メインリポジトリ:並列ループ分割後の旧アナライザバインディング無効化(#3226)
6 ファイル +55/-19。loop_partition、loop_vectorize、lower_tile_op の 3 つの pass とそのヘッダファイルに関わる。並列ループが分割されるとループ本体の反復変数が再バインディングされるが、pass は依然として分割前の変数でアナライザを参照し、旧変数の区間と整除情報を取得するため、簡約と境界判定が歪む。修正は分析コンテキストを pass ごとに分離し、ループ分析スナップショットの初期化経路を簡素化し、最小化された並列ループ分析の回帰テストケースを付属させた。コミッタには Codex の連署が含まれる。
出典:https://github.com/tile-ai/tilelang/pull/3226
1.3 メインリポジトリ:deepseek_v32 top-k セレクタ候補バッファの境界外(#3280)
2 ファイル +91/-10、#1351 を修正。tl_topk_impl は radix 候補を固定の [2, 4096] 共有バッファに格納するが、第一ラウンドの radix で選出されるしきい値バケットが 4096 要素を超えないことを制約するものは何もない——FP32 スコアの低位仮数が同一、または入力が全等の場合、アトミック追加がバッファ末尾を越えて書き込み、CUDA_ERROR_ILLEGAL_ADDRESS でクラッシュする。改法は tl_topk ラッパー層で (batch, 2, seq_len) のグローバル workspace を確保し、2 つの radix バッファをいずれも seq_len で受け止めるようにする。公開されている tl_topk() のシグネチャは不変。79 行の再現テストケースを付属させ、#1351 の入力構成に従っている。
出典:https://github.com/tile-ai/tilelang/pull/3280
1.4 メインリポジトリ:CUDA 16 ビット逆三角関数とバイナリ数学の降級補完(#3172)
3 ファイル +463/-8。asin、acos、asinh、acosh、atanh、log1p、atan2、hypot、copysign、ldexp、nextafter には CUDA 専用の下降ルールが欠けており、汎用ルールは float32/float64 のみを扱うため、float16 と bfloat16 の呼び出しは codegen 段階で誰も引き受けていない。修正はこれらを CUDAMath に登録し、half_t / bfloat16_t のオーバーロードを補う。10 個のオペレータは float32 で評価してから変換し直し、nextafter は 16 ビット符号化による整数比較ステップに変更し、fast math 下で bfloat16 の非正規数を保全する。float32/float64 の関数名は不変で、common.h 側に 182 行のテンプレート実装を追加した。
出典:https://github.com/tile-ai/tilelang/pull/3172
1.5 メインリポジトリ:sm8x スパースメタデータの非全単射形状を拒否に変更(#3125)
単一ファイル tilelang/layout/gemm_sp.py +6/-0、#2606 を修正。make_cutlass_metadata_layout_sm8x の interleave は、M がグループサイズの整数倍かつ K が偶数の場合にのみ全単射となる。満たさない場合、従来は誤ったメタデータを暗黙的に生成し、スパース GEMM の結果が通知なく誤っていた。現在は両条件とも明示的に拒否するよう変更。当初 PR に付随して追加されたスパースメタデータ形状テストはレビューで削除を求められ、メインブランチには入っていない。
出典:https://github.com/tile-ai/tilelang/pull/3125
1.6 メインリポジトリ:eager フロントエンドが Python ソース位置を保持(#3558)
4 ファイル。eager フロントエンドが Python の for / with 文を生成呼び出しに降ろした後、エラーとデバッグのマッピングが生成コードの行番号に戻ってしまう問題。修正により、生成されたループ呼び出しを for 文に、コンテキストラッパーを with 文にマッピングし、イテラブル引数のアンパック位置も補完する。#3125 と同様、レビューにより付随するソース位置回帰テストの削除が求められた。
出典:https://github.com/tile-ai/tilelang/pull/3558
1.7 審査中:pass 機能チェック、符号なし mulhi と assume 述語(10-22)
メインリポジトリの審査待ちキューは調査期間内に 50 件更新。注目すべきは [Transform] Locate VerifyBufferInit reads outside tile ops(#3620)、[Codegen] Add unsigned mulhi intrinsic(#3617)、[Transform] Keep assumption checks with local dependencies in scope(#3611)と [Transform] Preserve tl.assume predicates that simplify to true(#3610)、[CPU] Support clz in C and C host code generation(#3609)、昇騰側では [BugFix][Ascend] Reject odd FP4 view geometry in GM->L1 copies(#3608)と [Ascend] Support exact placement of versioned on-chip buffers(#3607)が同日更新。ほかに #3613 は「レビュー不要」と明記されたネイティブ argmin/argmax リダクション草案。
出典:https://github.com/tile-ai/tilelang/pulls
2. マルチバックエンド適配(昇騰 / 海光 / 沐曦 / 摩尔線程)
2.1 昇騰:NZ バッファが部分 UB→L1 コピーをサポート(#3530)
当日最大の変更:18 ファイル +993/-256。圧縮 NZ レイアウトのパディングは従来、論理レイアウト段階で平坦化されていたため、部分 UB→L1 コピー時にソースとデスティネーションそれぞれの C0 グループ行ストライドを表現できなかった。修正により、コンパクト NZ パディングをストレージ正規化段階まで保持し、容量に応じて一度だけ正確にマテリアライズする。コピー領域の降下時に 32 バイトブロック単位のソース/デスティネーション C0 グループ行ストライドを明示的に与え、行ストライドを scatter テンプレート、Ascend/PTO の DMA codegen とアクセスフットプリントまで通貫させる。同時に幾何・アラインメント・境界チェックを追加し、物理ビューのオーナーを保持し、呼び出し側を論理割り当て形状+明示的行ストライドへ移行。src/ascend/op/nd2nz.h は 83 行追加。
出典:https://github.com/tile-ai/tilelang/pull/3530
2.2 昇騰 PTO:MAD 方向述語の指数増大によるコンパイル停滞(#3581)
2 ファイル +48/-5。繰り返し可変ガードを持つ PTO カーネルは、行列積を含まなくても分析に 10 分以上かかることがある——これは NPU 実行前のコンパイル時停止である。根本原因は PTOMmadDirectionAnalyzer がガードによる可変メモリ読み取り時に進入状態を保守的に保持することにあり、両分岐が MAD 方向を保つ場合、Merge は同じ述語 S を受け取り、実装は常に S || S を構築して簡約する。算術簡約器は再現ケースでこの重複を消せず、連続分岐により述語式が倍々で膨張する。修正では両側の述語が同一の場合に再構築をスキップし、20 行の的を絞ったテストを追加。
出典:https://github.com/tile-ai/tilelang/pull/3581
2.3 昇騰 PTO:SIMT ベクトルローカル変数をレーンごとに割り当て(#3523)
4 ファイル +273/-29。SIMT ベクトルのローカル変数がスカラー PTODSL バッファとして割り当てられ、レーン数が計算されていないにもかかわらず Python の添字でアクセスされていた。AllocatedBufferValue は項目ごとの代入をサポートしておらず、ループ索引の書き込みは PTODSL AST 書き換え段階でより早く失敗していた。修正ではベクトルスロットごとに十分なスカラー要素を割り当て、スロット添字をスカラーオフセットに変換し、既存の連続 load/store パスを再利用することで、AscendC のベクトルローカルセマンティクスを維持し、ループ展開も変更せず、新たな下降パスも追加しない。float32x2 / float32x4 の 2 組の codegen テストを追加、計 209 行。
出典:https://github.com/tile-ai/tilelang/pull/3523
2.4 昇騰リポジトリ:cann-bench に Dilation2D を追加、同日に depthwise_conv_2d を回退(#1916 / #1919)
#1916 は Dilation2D フォワードをマージ(cann-bench level3 タスク):NHWC、fp16/fp32、任意の kernelSize/stride/rate と SAME/VALID に対応し、範囲外 tap の寄与は -inf + filter として golden とビット単位で一致し、NaN 入力の伝播、+inf filter tap により範囲外寄与が NaN になることを確認。実装は純粋な Vector kernel に Python wrapper を加えたもので、昇騰コピーエンジンの 32B 行ストライド整列制約に従ってディスパッチし、grid はジオメトリに応じて適応的に決め、コア数を固定値で書き込まない。同日 #1919 は 10-21 にマージされたばかりの depthwise_conv_2d(#1915)を全体として回退し、コミットメッセージには理由が記載されておらず、サンプルリポジトリは「1つ入れて1つ戻す」状態にある。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1916
2.5 昇騰リポジトリ審査中:sparse_mla フォワードと逆伝播(#1929 / #1928)
#1929 は sparse_mla_fwd(DSA 疎 MLA フォワード:varlen packed ランダム行 gather、オンライン softmax、QK/PV デュアル GEMM 融合)を新規追加し、examples/dsa_sparse_finetune/ に配置。デュアル kernel(fast 特化と portable 汎用)を host ディスパッチで組み合わせる。fast パスの適用範囲は H=128 かつ G=1、dim/tail は 512/64、topk は 256 の整数倍、NI は 1 から 6 の間と固定され、範囲外は汎用パスを通る。#1928 は sparse_mla_bwd(dQ GEMM チェーン、P/dP ベクトルチェーン、dKV ランダム行スキャッタ)を新規追加し、examples_experiment/dsa/ に配置してオペレータテスト一覧に登録、フォワード結果 k1 は本オペレータに折り込まれた。昇騰リポジトリの審査中は合計 11 件。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1929
2.6 海光:マージブランチで enable_register_pipeline を削除(10-22)
gcy_merge_main は 10-22 11:11 に 6 ファイル +10/-24 のクリーンアップをプッシュし、enable_register_pipeline スイッチを全体として削除した。HCU レジスタソフトパイプラインの注入とパイプライン計画の 2 つのパス、pipeline_utils.h、および言語層の loop 拡張が関与し、サンプルと性能スクリプトも呼び出しを同期して変更した。スイッチの削除は通常、レジスタパイプラインパスがすでにデフォルトで有効になっており、二重メンテナンスが割に合わなくなったことを意味する。ただし main ブランチは依然として 10-20 の preserve constant K origins in hoisted MLS windows(#25)のままで、この一連の変更はまだマージブランチ上で落地を待っている。
出典:https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main
2.7 沐曦 / 摩尔線程 / MLIR-Ascend:調査期間内にプッシュなし(10-22)
tilelang-metax の最終プッシュは 10-16 で停止、tilelang-musa は 09-30 で停止、tilelang-mlir-ascend は 10-19 で停止し、3 つの適応ラインはいずれも調査期間内にコミットがない。昇騰、海光の 2 ラインが本調査期間のバックエンド変更のすべてを担った。
出典:https://github.com/tile-ai/tilelang-metax
2.8 産業側:SCMP が SemiAnalysis レポートを転載、TileLang は 6 リポジトリ中唯一利用可能(10-22)
南華早報 10-22 に『Jensen’s nightmare』を掲載し、SemiAnalysis が木曜日に発表したファーウェイ-DeepSeek レポートを伝えた。DeepSeek の最新ソフトウェアリリースに加え、昇騰 950PR に関するさらなる実験から、米国との差が縮まっていることが示され、DeepSeek が 2027 年または 2028 年までに「主に」ファーウェイの「すでにかなり成熟した」技術スタックに依存すると判断している。このレポートの原文は前号ですでに収録したものであり、「昇騰のサポートを得た 6 つのリポジトリのうち TileLang のみが使用可能」という判断は変わらず、また昇騰 950 バックエンドは独立した適配層にとどまらず TileLang のメインリポジトリに統合されたと述べている。同一情報源による二次的な伝播であり、新たな技術的進展を構成するものではないが、この判断が主流メディアの論調に入りつつあることを示している。
3. エコシステムと採用側(TileOPs / TileFoundry / オペレータライブラリ)
3.1 TileOPs:FP8 ページド warp 特化 kernel の共有メモリアライアスにより NaN(#2708)
2 ファイル +80/-30。GQAPagedFP8WSKernel が「7 つの単一 token リクエストが 300 token リクエストより前に並び、4099 行の FP8 cache、48 行のページ」という組み合わせで NaN を返す。根本原因は、共有メモリプランナが warp ロールごとのバッファを異なる時間帯にアクティブと読み取ったため、split handover を依然 TMA で埋められている key または value リングアドレス上に配置し、handover の running maxima もリングフラグビットに押し込めてしまったことにある。修正ではこれらのバッファを共有メモリ内で明示的に分離し、24 行の回帰テストケース(qwen3-fp8-ragged 基準)を追加した。
出典:https://github.com/tile-ai/TileOPs/pull/2708
3.2 TileOPs:dense prefill を永続的走査に変更、D64 三消費者(#2706)
6 ファイル +590/-5、新規 GQADensePersistentWSKernel(対外名 gqa_dense_persistent、gqa_dense より優先):query tile 数が SM 数を超える場合、1 つの CTA が 1 つの SM を占めて交互のラウンドで自身の tile を処理し、現在の tile の epilogue 中に次の tile の query をロードする。D64 では各 tile は 192 行の query を 3 つの 64 行消費者 warpgroup に分散し、K/V の再読み込みが 3 分の 1 に減少;D128 は 128 行を維持する。
出典:https://github.com/tile-ai/TileOPs/pull/2706
3.3 TileOPs:SM90 dense GQA decode を warp 特化に変更(#2694)
新規 GQADecodeWSKernel:TMA プロデューサー warp が 2 スロットの K/V リングを埋め、1 つの WGMMA warpgroup が head グループのオンライン softmax を維持する;split 数は常駐 CTA スロットを一度に埋めるように定め、各 split が少なくとも 4 つの KV tile を持つことを保証し、単一 split 時は CTA 内で正規化して直接出力を書き、combine kernel を経由しない;KV head は CTA 間で最も速く変化する。
出典:https://github.com/tile-ai/TileOPs/pull/2694
3.4 TileOPs:因果スライディングウィンドウで 1 つの KV head の 2 つの key tile をクエリ head グループ全体にサービス(#2696)
新規 GQADenseSlidingWindowBandKernel:GQADenseFwdOp は従来、各 query head に対して同じ 2 つの key/value tile を再ロードし、1 つの CTA は 1 つ(query tile, head)のみを処理していた;新 kernel では 1 つの CTA が(query tile, KV head)に対して band の 2 つの key/value tile を一度だけロードし、順にそのグループの query head に対してリダクションを行い、同時に次の head の query tile を飛行させる。対象はスライディングウィンドウ行数が 128 key を超えないシナリオ。
出典:https://github.com/tile-ai/TileOPs/pull/2696
3.5 TileOPs:ページド decode の自動チューニングが以前は同一の block table を共用していた(#2703)
4 ファイル +24/-13。ページ分割 decode の autotune 供給は、各リクエストに同じ block table を渡すため、実質的に 1 リクエスト分のページしかテストしておらず、しかもそのページは常時 L2 に滞留している。候補ランキングは、サービス対象バッチが実際に支払ったトラフィックを反映している。実測では 1500 MHz の H200 上で、MHA decode チューナーは opt-1.3b-b16 において block_N 128(コールド 80.0 us)を選び、block_N 64(72.5 us)を放棄した結果、flashinfer の 0.994 まで低下した。修正により、チューニングはリクエストごとに異なる cache ページを読み取るようになった。
出典:https://github.com/tile-ai/TileOPs/pull/2703
3.6 TileOPs:masked softmax の FMA オーバーフローと varlen RoPE の奇数半幅(#2699 / #2700)
#2699(4 ファイル +54/-6)は dense GQA スライディングウィンドウ前向きの NaN 行を修正する:因果ウィンドウ下の 128x64 チューニングタイルでは、各 query tile が先頭行のウィンドウから開始するため、最後の数行は連続して 2 つの全マスク key tile に遭遇する。最初のブロックの後、2 つの running maxima はともに -1e38 のクランプ上にあり、fast math が prev * scale - curr * scale を FMA に融合すると、丸め誤差(約 2 の 99 乗)がオーバーフローして NaN になる。修正方法は、スケーリング再スケールの前に running maxima を先に減算することである。#2700 は varlen query RoPE の奇数半回転幅(rotary_dim 10、14 など)におけるループ境界を修正する:並列度 rows 掛ける chunks が回転分岐のスレッド数で割り切れない場合、TileLang はループ上限を失い、warp 特化パスの後続 query 行が誤って計算され、汎用パス(dim 256)はそのままハングする。
出典:https://github.com/tile-ai/TileOPs/pull/2699
3.7 TileOPs:FP8 grouped GEMM の段階的ステージング scale と gate/up ピンポン warp group(#2704 / #2693)
#2704 は設定キー stage_scales を新設する:プロデューサーが cp.async で各段階の A/B scale を共有メモリにコピーし、着地時にその段階の barrier に到達させる;stage_scales=0 は引き続き旧 kernel を生成する。背景は、masked FP8 grouped GEMM が数学 warpgroup 内で各ステップ K の A scale をグローバルから読み取るため、リングが深くなると L1 が圧縮されこれらの load が停止し、ds-v3-e8-up-masked が DeepGEMM の 0.979 しか出ないことである。#2693 は GemmTemplate にピンポン形態を追加する:2 つの数学 warpgroup がそれぞれ完全な 64 行 tile を実行し、CTA の tile を交互に取得し、一方のグループの epilogue が他方のグループの mainloop の下に落ち、ローテーションバリアで順序付ける。
出典:https://github.com/tile-ai/TileOPs/pull/2704
3.8 TileOPs:述語ブロードキャストとブールオペランドを 16 バイトベクトル読み取りに変更(#2702 / #2697)
#2702 は行ブロードキャスト述語(Eq/Ne/Gt/Lt/Ge/Le と LogicalAnd/Or)において、1024 列ごとのブロック全体の stride-1 オペランドを 16 バイトベクトルで load16_evict_last 経由で読み取るようにする。一部のブロックと他の staged body は fragment コピーを保持する;linear-bias-float32(1024x4096 対 1x4096)の 6 種類の比較は 0.99 から 1.06 に向上し、論理オペレータは 1.04 となった。#2697 は比較と論理の二値 kernel において全サイズオペランドを 16 バイト evict-first ベクトル読み取りに変更する。float32 行は以前 torch.compile の 0.99 から 1.00 であった。
出典:https://github.com/tile-ai/TileOPs/pull/2702
3.9 TileOPs:SM89 の 3 パスを同日に推進(#2705 / #2691 / #2690)
#2705 は SM89 の SSD chunk scan のデフォルトパイプライン段数を 2 段に下げ、scan kernel のバッファ需要を削減する。6 形状の FP16/BF16 実測は 3 ラウンド、各ラウンド 200 回の CUPTI サンプリングの中央値を取る。#2691 は既存プランナーがマルチパーティションを選択する場合に SM89 の paged DSA key 区間を分割し、既存の FP32 部分出力と log-sum-exp の merge を再利用する;単一パーティション時は BF16 出力と単一 kernel を維持し、コンパイルターゲットを merge キャッシュ識別子に含め、ターゲットごとに独立した JIT builder を構築する。#2690 は SM89 上で query head グループ 32/64、head dim 64/128 の NSA top-k 選択に 32 列 tile を使用するよう変更し、64 行 tile と同じスコアリングソート実装を保持する。8 つの FP16/BF16 ユースケースの GPU busy 中央値は 2.1% から 15.0% 低下した。
ソース:https://github.com/tile-ai/TileOPs/pull/2691
3.10 TileOPs:リダクション折り畳みの単一発射と空 DSA 選択検証(#2689 / #2698)
#2689 は ReduceFoldKernel を新規追加し、各保持インデックスの lead 段の連続サフィックスを一度の発射で折り畳む(従来は軸プレフィックスとサフィックスのリダクションに二回の発射が必要だった)。また ReduceEdgeKernel と VectorNormEdgeKernel に対し、行を 512 スレッドブロックでインフライト保持できる呼び出しをすべて委譲するようにした。背景は、Sum が 3d-multidim-float16(0.977)と mlp-intermediate-float32(0.990)で torch-compile に劣後していることにある。#2698 はページング DSA 検証を修正する。query 行に有効な選択 key が一切ない場合、参照実装はゼロ出力と負の無限大 LSE を返すが、参照 LSE を減算すると NaN 誤差界となり、正しい結果を誤りと判定してしまう。修正では空行に正確なゼロ LSE 界を与え、「すべての行が空」の場合に検証器が非ゼロ出力と有限 LSE を拒否しなければならない制御用例を追加した。
ソース:https://github.com/tile-ai/TileOPs/pull/2689
3.11 TileOPs:ベンチマーク参照を float32 で計算するよう変更(#2692)
ref_program は従来、正確性参照であると同時に torch ベースライン計測の対象でもあり、多くの参照が 16 ビット入力 dtype のまま直接 torch を呼んでいたため、結果はデバイス側 torch の各ステップの丸め方に依存していた。CUDA はこの差異を覆い隠すが、CPU torch と torch_musa は覆い隠さない。GroupNorm と InstanceNorm は MTT S5000 上で失敗する一方、同一 kernel は float64 参照に対しては合格していた。修正では ref_program を float32 以上で計算し、torch ベースラインは入力 dtype で計時するようにした。これは国産 GPU バックエンドによって露呈したが、まず参照実装層で修正すべき稀な問題である。
ソース:https://github.com/tile-ai/TileOPs/pull/2692
3.12 TileFoundry:sm_89 アーキテクチャと RTX 6000 Ada デバイスをインストール(#259)
11 ファイル +392/-22。nvidia_sm89.toml と nvidia_rtx_6000_ada.toml の二つのターゲット記述(それぞれ 114 行と 124 行)を新規追加し、docs/spec/target.md にも 42 行を同期追加した。このコミットは TileFoundry のデバイスモデリングを Hopper から Ada に拡張するもので、同日 #256 で露呈した「sm_90 ターゲットでコンパイルされたテストが sm_89 単卡上で値のみ誤る」問題と表裏一体である。併せて sm_89 の e4m3 warp mma 宣言も追加した(#261、審査中)。
ソース:https://github.com/tile-ai/TileFoundry/pull/259
3.13 TileFoundry:gmem 計價をフットプリントと実行 mesh に基づく方式へ変更し、被呼び出し側の次元をバインド(#258)
当日最も広範な変更:74 ファイル +1128/-553、#239、#252、#253 を修正。従来は 132 個の CTA が 40 個しか探査しない場合、計價は「各ウィンドウの各位置を一回移動」で計算するため、under-fill を全く検出できなかった。被呼び出し側のみが宣言できる次元はバインドできず、analyze の出力ターゲットと関数ごとのメモリ行の意味も不明瞭だった。修正では check と analyze の次元バインドを instantiate に集約し、被呼び出し側またはサブモジュールが宣言する次元をバインド可能にした。コミットはコミュニティ貢献者が起案し、メンテナが引き継いで書き直した。
ソース:https://github.com/tile-ai/TileFoundry/pull/258
3.14 TileFoundry:デバイスが kernel 起動を拒否した場合は必ずエラーを投げる(#256)
7 ファイル +69/-11。従来はランチャーが kernel<<<grid, block, smem, nullptr>>>(...) を発行した後、起動状態を一切読まなかったため、デバイスが起動を拒否しても呼び出しはそのまま返り、出力は旧内容のままとなり、エラーが呼び出し側に届かなかった。sm_89 単卡上では、バンドルされた sm_90 ターゲットでコンパイルされたテストは値が正しくないという形でのみ現れ、compute-sanitizer によって初めて cudaLaunchKernel が cudaErrorNoKernelImageForDevice を返すことが示された。修正では CUDA ランチャーが cudaGetLastError() を読んだ後にエラーを投げるようにし、CPU 側 module も同期して変更、仕様ドキュメントとテストも併せて更新した。
ソース:https://github.com/tile-ai/TileFoundry/pull/256
3.15 ディストリビューション版:Sunrise が 0.1.15+sunrise.1.2.0 のリリース PR を作成(#9、審査中)
10-22 16:48 candidate/sunrise-0.1.15-1.2.0-20261019 ブランチを作成して PR を開き、Sunrise ディストリビューション版を 0.1.15+sunrise.1.2.0 に更新し、TileLang 0.1.15 およびそれに伴いレビュー済みの TileOPs、TileKernels、TVM、tvm-ffi の変更を取り込んだ。ソースのインポートとマージ、ディストリビューション適配は別々にコミットし、最終的に TANG バックエンド向けの Sunrise wheel を適配設定してスライス GEMM のリグレッションを修正した(2 つのアキュムレータフラグのいずれも公開テンプレートで検証し、stride、origin、offset のアサーションは保持)。自己申告の検証:ローカル契約テスト 55 件通過、TANG パイプラインの codegen テスト 54 件通過、2 つの lint ジョブと信頼境界チェックが通過。ただし S2 wheel のビルドでパッケージが生成された後、インストール後の import スモークチェックが tangGetDeviceCount の箇所で失敗した。ブランチ上の 3 コミットはそれぞれ 0.1.15 ディストリビューション適配、CI を TANG バックエンドに絞り込む変更、スライス GEMM アキュムレータフラグのテストである。
出典:https://github.com/tile-ai/tilelang-sunrise/pull/9
3.16 採用側:TileKernels と FlashQLA は調査期間内にプッシュなし(10-22)
deepseek-ai/TileKernels の最終プッシュは 10-14 で停止、QwenLM/FlashQLA も同様に 10-14 で停止しており、両採用側ともに本調査期間内にコミットはない。
出典:https://github.com/deepseek-ai/TileKernels
4. コミュニティ、チュートリアルとイベント
4.1 CANN コミュニティが ops-tilelang オペレータリポジトリを開設(10-22)
Atlas-Ascend-Fork/ops-tilelang は 10-20 にリポジトリを開設し、CANN コミュニティが昇騰 NPU 向けに提供する TileLang 高性能オペレータリポジトリと位置づけられている。Kernel 実装、自動化テストケース、標準ワークロードと性能ベースラインを統一的に管理し、再利用可能・検証可能・持続的に進化可能なオペレータ資産の形成を目標としている。リポジトリ内には src、tests、scripts、docs の 4 層と、OAT.xml、SECURITY.md などの開放原子コミュニティ規範ファイルがすでにあり、10-22 には「正当な上界シードとケースハッシュの組み合わせでオーバーフローし CPU テストが失敗する」問題を修正するコミットをプッシュした。それ以前のコミットには「Ascend A5 の参照を Ascend 950 に置換」と「benchmark が失敗または中断したセッションでも性能ベースラインを上書きしてしまう」問題がある。単一 DSL を唯一の実装言語とするコミュニティレベルのオペレータリポジトリはこれまで見られず、この動向は週次で追跡する価値がある。
出典:https://github.com/Atlas-Ascend-Fork/ops-tilelang
4.2 tilelang/agent:オペレータ最適化を agent の SKILL ワークフローとして記述(10-22)
tilelang/agent(08-28 開設、3 スター)は 10-22 にプッシュされ、README 自体が実行可能な prompt テンプレートになっている。@tilelang-operator-dev/SKILL.md でドメイン知識を注入し、入力はオペレータ名、GPU/CPU 参照実装、shape と「tilelang npu 納品、中国語納品」の制約のみである。コメントには、agent は A5 環境で実行する必要があり、ロード後は 10 ステップのクローズドループに従うことが明記されている。Skill には選型決定木、8-CSV ルール、L0 から L4 までの戦略、11 の落とし穴が含まれる。4.1 のオペレータリポジトリ、3.15 の命令レベルスケジューリングと本リポジトリを合わせて見ると、昇騰側では「人がルールを書き、agent がチューニングし、リポジトリにベースラインを蓄積する」という分業が形成されつつある。
出典:https://github.com/tilelang/agent
4.3 コミュニティ新規リポジトリ:Sail ラッパー、リモート自動チューニング agent、チュートリアルリポジトリ(10-22)
flytiger-eco/tilelang-wrapper-for-sail は 10-21 に開設され 10-22 にプッシュ(Python、説明なし)。cc-c122/TileLang-Remote-Autotuning-Agent(Python、10-22 プッシュ)と easy-tilelang/tilelang-devkit、easy-tilelang/tilelang-tutorials が同日にプッシュされた。数が多く単一リポジトリのスターは低く、チュートリアルとツール層の自然な増加であり、まだ技術的方向性を判断できるものではない。
出典:https://github.com/flytiger-eco/tilelang-wrapper-for-sail
4.4 ドキュメントとバージョン:ドキュメントサイトが 3 回更新、VERSION は依然 0.1.15(10-22)
tilelang.github.io は 3 回 Update docs、TileFoundry.github.io は 1 件 render TileFoundry v0.0.3、TileOPs.github.io は 1 件。メインリポジトリの VERSION は 0.1.15 で停止しており、直近のリリースは依然 09-30 の v0.1.15 である。リポジトリは 8742 スター、915 fork、オープン issue 389 件。
ソース:https://github.com/tile-ai/tilelang.github.io
5. トレンド観察
1. メインリポジトリの欠陥タイプが層を変えた。 前週にメインリポジトリで片付けたのは「形状と幾何検証」だった——cast の丸め注釈が静かに破棄される、arch なしの CUDA target が sm_50 へ静かにフォールバックする、スパースメタデータの形状が検証されない。今回の調査期間では 9 件のうち 4 件がアナライザとソース位置というコンパイル期の状態に関するものだった:analyzer が書き換え後に再利用される、分割後に変数が再バインドされる、eager フロントエンドが Python 位置を失う。この種の欠陥に共通するのは「生成された kernel は依然として正しいかもしれないが、コンパイラがあなたに言うことと実際に行うことが一致しない」という点であり、TileLang が agent による自動 kernel 記述(4.2)を接続しつつある今、コンパイル期の状態一貫性が agent の反復ループを収束させられるかどうかを直接左右する。
2. 昇騰側の進展はオペレータ数から基層のデータ経路へ移った。 #3530 は 18 ファイルに及ぶ 1 件で、NZ 圧縮レイアウトを一部の UB→L1 コピーにおける行間隔表現として扱うもの;#3581 と #3523 はそれぞれコンパイル期の述語爆発と SIMT ベクトルローカル変数の割り当てセマンティクスを修正する。3 件はいずれもオペレータを新規追加するものではなく、既存経路をより多くの形状でコンパイル可能かつ正しく実行可能にするものだ。CANN コミュニティが別途 ops-tilelang オペレータリポジトリを立ち上げた(4.1)ことと合わせると、分業が読み取れる:オペレータ資産はコミュニティリポジトリが担い、メインリポジトリはバックエンドセマンティクスと codegen に集中する。
3. TileOPs は 2 日間で連続 20 件、作業の重心は SM90 特化と SM89 の補完にある。 調査期間内に 3 つの SM90 warp 特化 kernel(dense decode、因果スライディングウィンドウ band、FP8 varlen オーバーラップ)を新規追加するとともに、SM89 の SSD chunk scan、paged DSA、NSA top-k という 3 つの小規模経路も併せて推進した。注目すべきは 2 件の NaN 修正(#2708、#2699)がいずれも warp 特化経路に現れている点だ:共有メモリプランナの生存期間判定と fast math 下の FMA 融合は、この種の手動編成 kernel で最もエラーが起きやすい 2 か所であり、まさに TileLang がユーザに代わって管理することを約束している部分である。
4. 逆方向の指標の一つがベンチマーク層だ。 #2692 は参照実装を float32 計算に変更した。原因は MTT S5000 上で GroupNorm と InstanceNorm が失敗したことであり、問題は kernel 自体ではなく参照実装が入力 dtype に追随していたことにあった。同日、TileFoundry #256 は「デバイスが起動を拒否する」をサイレントから例外送出に変えた。2 件は同じ方向にある:「通ったように見える」を「確かに通った」に変えることであり、これらのツールに依存してチップ適応を行うチームにとって、この種の修正の実際の価値はいくつかの性能数値よりも高い。
5. 配布と産業の 2 本の線が加速している。 Sunrise 配布版は 0.1.15 と TANG バックエンドの wheel をパッケージ化してリリース PR とし、TileLang に既に複数の下流配布形態が存在することを示している;産業側では SCMP が SemiAnalysis の判断を転載し、「6 リポジトリのうち利用可能なのは TileLang のみ」という見方を主流メディアへ押し出した。技術のリズムとしては依然として穏やかな一日である(メインリポジトリはリリースなし、breaking change なし)が、エコシステム層の組織的な動きはコード層よりも明らかに多い。
付録:素材と検証の説明
情報源検証表
| 情報源 | 検証結果 |
|---|---|
| tile-ai 組織リポジトリ | 29 リポジトリ、調査期間内に 10 件プッシュ;35 件マージ、審査中更新 71 件 |
| tilelang メインリポジトリ | 9 件マージ、1 件は単行修正、1 件は 18 ファイルの昇騰変更;VERSION 0.1.15、新規リリースなし |
| TileOPs | 20 件マージ、すべて Attention / MoE / Elementwise / Reduction の kernel とテスト |
| TileFoundry | 3 件マージ、別途 4 件審査中;0.0.3 の 10-21 リリースの説明は前号で既に収録済み |
| tilelang-ascend | 2 件マージ(1 件はサンプル追加、1 件はサンプル巻き戻し)、11 件審査中 |
| tilelang-hygon | main はコミットなし;マージブランチ gcy_merge_main にスイッチ整理 1 件 |
| tilelang-sunrise | candidate ブランチを新規作成し 0.1.15+sunrise.1.2.0 リリース PR を起票、未マージ |
| tilelang-metax / musa / mlir-ascend | 調査期間内にプッシュなし、それぞれ 10-16 / 09-30 / 10-19 で停止 |
| TileRT / tilescale / DeepStack | 調査期間内にプッシュなし |
| 採用側 TileKernels / FlashQLA | いずれも 10-14 で停止、プッシュなし |
| Google News(中英語の複数組、プロキシ経由) | ヒット 3 件、うち 2 件は同一の SemiAnalysis レポートのメディア転載、1 件は DeepSeek の資金調達で主題と無関係のため除外済み |
| Hacker News | TileLang 関連の議論なし |
| arXiv | TileLang をタイトルまたは要旨のキーワードとして検索、調査期間内 0 件 |
| コミュニティ新規リポジトリ | 新規または活発なリポジトリ 5 個、いずれも作成日とプッシュ時刻を確認済み |
| 会員またはベンダー公式サイト | 本主題には該当なし |
完全な情報源リスト
- [1] メインリポジトリ ベクトル化前クローンアナライザ(#3612)— https://github.com/tile-ai/tilelang/pull/3612
- [2] メインリポジトリ 並列ループ分割後アナライザバインディング(#3226)— https://github.com/tile-ai/tilelang/pull/3226
- [3] メインリポジトリ deepseek_v32 top-k 候補バッファ(#3280)— https://github.com/tile-ai/tilelang/pull/3280
- [4] メインリポジトリ CUDA 16ビット数学デグレード(#3172)— https://github.com/tile-ai/tilelang/pull/3172
- [5] メインリポジトリ sm8x スパースメタデータ形状検証(#3125)— https://github.com/tile-ai/tilelang/pull/3125
- [6] メインリポジトリ eager フロントエンドソース位置(#3558)— https://github.com/tile-ai/tilelang/pull/3558
- [7] メインリポジトリ 審査待ちキュー — https://github.com/tile-ai/tilelang/pulls
- [8] 昇騰 NZ バッファ部分 UB→L1 コピー(#3530)— https://github.com/tile-ai/tilelang/pull/3530
- [9] 昇騰 PTO MAD 方向述語増加(#3581)— https://github.com/tile-ai/tilelang/pull/3581
- [10] 昇騰 PTO SIMT ベクトルローカル変数(#3523)— https://github.com/tile-ai/tilelang/pull/3523
- [11] 昇騰リポジトリ Dilation2D サンプル(#1916)— https://github.com/tile-ai/tilelang-ascend/pull/1916
- [12] 昇騰リポジトリ depthwise_conv_2d ロールバック(#1919)— https://github.com/tile-ai/tilelang-ascend/pull/1919
- [13] 昇騰リポジトリ sparse_mla_fwd(#1929)— https://github.com/tile-ai/tilelang-ascend/pull/1929
- [14] 昇騰リポジトリ sparse_mla_bwd(#1928)— https://github.com/tile-ai/tilelang-ascend/pull/1928
- [15] 海光 マージブランチコミット — https://github.com/tile-ai/tilelang-hygon/commits/gcy_merge_main
- [16] 沐曦 適応リポジトリ — https://github.com/tile-ai/tilelang-metax
- [17] SCMP 転載 SemiAnalysis レポート — https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report
- [18] TileOPs FP8 ページド共有メモリアライアス(#2708)— https://github.com/tile-ai/TileOPs/pull/2708
- [19] TileOPs dense prefill 永続化(#2706)— https://github.com/tile-ai/TileOPs/pull/2706
- [20] TileOPs SM90 dense decode(#2694)— https://github.com/tile-ai/TileOPs/pull/2694
- [21] TileOPs 因果スライディングウィンドウ band kernel(#2696)— https://github.com/tile-ai/TileOPs/pull/2696
- [22] TileOPs ページド decode チューニング block table(#2703)— https://github.com/tile-ai/TileOPs/pull/2703
- [23] TileOPs masked softmax NaN(#2699)— https://github.com/tile-ai/TileOPs/pull/2699
- [24] TileOPs varlen RoPE ループ境界(#2700)— https://github.com/tile-ai/TileOPs/pull/2700
- [25] TileOPs FP8 grouped GEMM 段階的 s
cale(#2704)— https://github.com/tile-ai/TileOPs/pull/2704
- [26] TileOPs gate/up ピンポン warp group(#2693)— https://github.com/tile-ai/TileOPs/pull/2693
- [27] TileOPs 述語ブロードキャストベクトル読み込み(#2702)— https://github.com/tile-ai/TileOPs/pull/2702
- [28] TileOPs ブールオペランド evict-first(#2697)— https://github.com/tile-ai/TileOPs/pull/2697
- [29] TileOPs SM89 SSD パイプライン階層化(#2705)— https://github.com/tile-ai/TileOPs/pull/2705
- [30] TileOPs SM89 paged DSA 分割(#2691)— https://github.com/tile-ai/TileOPs/pull/2691
- [31] TileOPs NSA top-k 列 tile(#2690)— https://github.com/tile-ai/TileOPs/pull/2690
- [32] TileOPs リダクション折り畳み(#2689)— https://github.com/tile-ai/TileOPs/pull/2689
- [33] TileOPs 空 DSA 選択検証(#2698)— https://github.com/tile-ai/TileOPs/pull/2698
- [34] TileOPs 参照実装を float32 に変更(#2692)— https://github.com/tile-ai/TileOPs/pull/2692
- [35] TileFoundry sm_89 アーキテクチャと RTX 6000 Ada(#259)— https://github.com/tile-ai/TileFoundry/pull/259
- [36] TileFoundry gmem 価格計算と次元バインディング(#258)— https://github.com/tile-ai/TileFoundry/pull/258
- [37] TileFoundry 起動失敗時のエラー送出(#256)— https://github.com/tile-ai/TileFoundry/pull/256
- [38] TileFoundry 0.0.3 リリースノート — https://github.com/tile-ai/TileFoundry/releases/tag/v0.0.3
- [39] Sunrise 0.1.15+sunrise.1.2.0 リリース PR — https://github.com/tile-ai/tilelang-sunrise/pull/9
- [40] 採用元 TileKernels — https://github.com/deepseek-ai/TileKernels
- [41] CANN コミュニティ ops-tilelang オペレータリポジトリ — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [42] tilelang/agent オペレータ最適化 skill — https://github.com/tilelang/agent
- [43] Sail ラッパーリポジトリ — https://github.com/flytiger-eco/tilelang-wrapper-for-sail
- [44] リモート自動チューニング agent リポジトリ — https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent
- [45] メインリポジトリドキュメントサイト — https://github.com/tile-ai/tilelang.github.io