TileLang デイリーレポート (2026-10-27)
調査期間:2026-10-26 07:00 ~ 2026-10-27 07:00(北京時間、24 時間)。
情報源:GitHub tile-ai 組織のリポジトリ別調査(tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / ドキュメントサイト / 言語サービスリポジトリ / 夜間ベンチマークリポジトリ)+ マージ済み・審査中実値検索(マージ済み 30 件、審査中更新 49 件)+ 新規 issue 全件確認(11 条)+ コミットレベル diff とリリースページ確認 + PyPI 成果物状態確認 + 採用側リポジトリ確認(TileKernels / FlashQLA)+ コミュニティ新規リポジトリ検索 + Google News 中国語・英語複数クエリ(プロキシ経由)+ Hacker News + arXiv + 産業メディア報道。
本期インデックス
- 本日の重点:24時間で30件をマージ、審査中更新49件、新規issue 11件;メインラインは「ストリームとバリアのセマンティクス」に関する3件のサイレントエラーを中心に展開;TileOPsの初回リリースは依然としてPyPIの最後の1枠で停滞;海光の大型マージ2件
-
- コアプロジェクトの進捗
- 1.1 メインリポジトリ:out_idx出力を呼び出し側ストリーム上での割り当てに変更、H100で再現した5回中3回の破損がゼロに(#3649)(10-26)
- 1.2 メインリポジトリ:NVRTCデバイス側コンパイルをTVM-FFI実行パスに接続、コンパイルキー項目を新設(#3285)(10-26)
- 1.3 メインリポジトリ:バッチTMAロードのcache hintにテンプレートパラメータが1つ不足、DeepSelectのTMAパスがコンパイル不能(#3683)(10-26)
- 1.4 メインリポジトリ:リダクションのFP8/FP4非有限定数が不正リテラルを出力、サブバイトアキュムレータを早期拒否に変更(#3189)(10-26)
- 1.5 メインリポジトリ:SM120ブロック単位スケーリングGEMMでKアトムを明示展開、FullCol構成が51%向上(#3286)(10-26)
- 1.6 メインリポジトリ:JIT出力形状内のスカラパラメータを割り当て前に解析(#3602)(10-26)
- 1.7 メインリポジトリ:sm8xスパースメタデータのオペランドdtypeがフォールバックなしで流入しないように(#3335)(10-26)
- 1.8 メインリポジトリ:DeepSelect Top-Kサンプルをマージ、129構成中torch.topk比で最大8.00倍(#3304)(10-26)
- 1.9 審査中:TIRxリファクタリング規模がさらに1194ファイルに拡大、パッチ未適用の上流バージョンにアンカー(#3666)(10-26)
- 1.10 新規欠陥:ThreadSyncの3種類のフェンス位置ずれ、審査中修正3件を併設(#3680 / #3682 / #3681)(10-26)
- 1.11 新規欠陥:sm_100以上のリダクションbatchとレイアウトの競合、審査中修正1件あり(#3667 / #3684)(10-26)
- 1.12 審査中:転置ターゲット形状検証、リージョン化killの再順序付け(#3677 / #3670)(10-26)
- 1.13 クローズ済み:ファジングテストシリーズのコンパイル時クラッシュ6件(10-26)
- コアプロジェクトの進捗
-
- マルチバックエンド適応(昇騰 / 海光 / 沐曦)
- 2.1 昇騰メインリポジトリ:バージョン付きオンチップバッファの精確アドレッシングインターフェース(#3607)(10-26)
- 2.2 昇騰審査中:手動スケジューリングのFlashMLA、BF16プリフィルとFP8デコード(#3678)(10-26)
- 2.3 昇騰審査中:オンチップメモリオーバーフローの事前拒否、E4M3のUBからL1へのコピー、ビューレイアウトとポインタ単位(#3655 / #3639 / #3623)(10-24~26)
- 2.4 昇騰審査中:バッチスケジューリングのオプション化、クロスコア同期のバッチモード要求、L0Cデュアルコピー直接転送(#3675 / #3614 / #3468)(10-26)
- 2.5 昇騰リポジトリ:インストーラのSHMEMセクションを二重修正、空のwheelが偽成功しなくなる(#1943 / #1946)(10-26)
- 2.6 昇騰リポジトリ:W4A4のint4行列積結果が誤り、根本原因はニブルパッキングを1バイトとして扱っていたこと(#1731)(10-26)
- 2.7 昇騰リポジトリ:2次元超のGMからUBへのコピーでDMA行数に従い先頭次元を折り畳み(#1805)(10-26)
- 2.8 昇騰新規欠陥:Cythonラッパーが呼び出しごとにアナライザを生成、約0.28ミリ秒(#1945)(10-26)
- 2.9 海光:レジスタパイプラインとwarp divergenceをマージ、28ファイル(#18)(10-26)
- 2.10 海光:ネイティブB16ストレージと動的WASP GEMMをマージ、scaleバッファのデュアルバンクは審査中(#28 / #29)(10-26)
- 2.11 沐曦:組織ビューにプッシュ記録あり、6ブランチいずれも期間内の新規コミットなし(10-26)
- マルチバックエンド適応(昇騰 / 海光 / 沐曦)
-
- エコシステムと採用側(TileOPs / TileFoundry)
- 3.1 TileOPs:初回リリースは依然として最後の1枠、アーティファクトは依然として2025-10-09の開発版(#2738)(10-26)
- 3.2 TileOPs:D128稠密フォワードでkeyを分割、prefill-ropeが31.04マイクロ秒から17.09に短縮(#2777)(10-26)
- 3.3 TileOPs:D128永続化フォワードをスレッド単位で書き換え、LSEを出力(#2774)(10-26)
- 3.4 TileOPs:融合加算のLayerNormで行をCTA間に分割、精度欠陥は別途新規チケット(#2768 / #2773)(10-26)
- 3.5 TileOPs:要素単位とリダクションの4件、clampのbf16行が5.9%高速化(#2766 / #2767 / #2769 / #2770)(10-26)
- 3.6 TileOPs:呼び出しテンソルのアライメント判定をオペレータ層に移動、3件で収束(#2772 / #2778 / #2780)(10-26)
- 3.7 TileOPs:オフセットストレージの3種修正、SM90グループGEMMと出力ABI(#2761 / #2765 / #2762)(10-26)
- 3.8 TileOPs審査中:FP8クエリのページドプリフィルをe4m3によるvalue乗算に変更(#2775)(10-26)
- 3.9 TileOPs審査中:スパース長シーケンスのGD
- エコシステムと採用側(TileOPs / TileFoundry)
N プレフィル分割スキャン(#2779)(10-26)
- 3.10 TileFoundry:動的共有メモリをカーネル実際のビューに合わせてサイズ決定(#271)(10-26)
- 3.11 TileFoundry:ウィンドウ全体をオペランド自体で発行し、sm_89 の GEMM がコンパイル可能に復帰(#273)(10-26)
- 3.12 TileFoundry の新欠陥:分析側が共有メモリのピークを二重計上し、しかも TIR に伝播しない(#272)(10-26)
-
- コミュニティ、チュートリアル、イベント
- 4.1 言語サービス:実験的 CUDA デバッグを公開、cuda-gdb と Nsight に対応(10-26)
- 4.2 ドキュメントサイト:1 日で 8 件の更新(10-26)
- 4.3 コントリビューション・プロセス整備の 4 連発:PR 説明の規範化、純ドキュメント変更時のバックエンドテスト省略、Agent タイトル確認、レビューボットのノイズ低減(#3671 / #3672 / #3673 / #3674)(10-26)
- 4.4 コミュニティ新リポジトリ:RISC-V コミュニティハブリポジトリ、CANN オペレータリポジトリ、Hexagon バックエンドリポジトリ、デバッグツールリポジトリ(10-26)
- 4.5 産業側:昇騰と DeepSeek のエコシステム格差が縮小しているとの報道(10-26)
- コミュニティ、チュートリアル、イベント
-
- トレンド観察
- 付、素材と検証説明
今日の重点
24 時間以内に tile-ai 組織で 30 件がマージされ(tilelang 14、TileOPs 14、tilelang-hygon 2)、審査中の更新は 49 件、新規 issue は 11 件、クローズは 12 件。昨日の 72 時間ウィンドウでは 78 件だった量が単日で 30 件に落ちたが、ペースは正常であり、内容の重心は明らかに「性能清掃」から「ストリームとバリアのセマンティクス」へ移っている。メインリポジトリにマージされた出力割り当て修正(#3649)は、H100 上で 5 回中 3 回出力が破損する問題をゼロにした。原因は、3 つのアダプタが呼び出し側の stream を受け取っているのに、現在のストリーム上で out_idx 出力を割り当てていたことにある。キャッシュアロケータは割り当てストリーム順に再利用するため、別のストリームでキューに入っている処理がカーネル結果を上書きしていた。同日浮上した 3 件の ThreadSync 欠陥(#3680)は、フェンスが発散分岐内に置かれていること、分岐内のフェンスが外側の保留依存を誤ってクリアすること、スレッド数でカウントするループ内のフェンス位置を指摘しており、#3682 と #3681 の 2 件の審査中修正が対応している。さらに tvm フォークリポジトリの 1 件では、ブール演算のビット演算が論理演算として翻訳されていたため、証明器が 10000 ステップの予算内で諦めなくなった(#88)。リリース側は進展なし:TileOPs の初回リリーストラッキング issue(#2738)は TestPyPI の試行と最初のタグを残すのみで、PyPI 上の最新は依然 2025-10-09 の 0.0.1.dev2 であり、リポジトリにはタグも release もない。バックエンド側では海光が 1 日で 2 件の大型マージを実施:レジスタパイプラインに warp 発散を追加(#18、28 ファイル +5312/-64)とネイティブ B16 ストレージに動的 WASP GEMM を追加(#28)。昇騰メインラインにはバージョン付きオンチップバッファを備えた精密アドレッシングインターフェースがマージされた(#3607、14 ファイル +982/-26)。審査中で最も重いのは依然 #3666 で、規模は昨日の 1000 強から 1194 ファイル +106648/-104497 に拡大している。
1. コアプロジェクトの進展
1.1 メインリポジトリ:out_idx 出力を呼び出し側ストリーム上で割り当てるよう変更、H100 で再現した 5 回中 3 回の破損がゼロに(#3649)
10-26 22:36 にマージ、5 ファイル +183/-31、#3648 を修正。cython、nvrtc、cutedsl の 3 つのアダプタはいずれも stream= を受け取りそのストリーム上で起動するが、out_idx 出力は torch.empty で現在の torch ストリーム上に割り当てられていた。キャッシュアロケータは「割り当てストリーム」順にのみ再利用するため、キューに処理が残っている VRAM ブロックが空きブロックとして返され、その処理がその後カーネル結果を上書きしていた。修正は output_allocation_stream(stream, device) を新設:呼び出し側ストリームが現在のストリームと異なる場合は呼び出し側ストリームを返し、そうでなければ空を返す。ハンドルがゼロの場合はデフォルトストリームにマッピングする。旧版 PyTorch の ExternalStream(0) はデフォルトストリームではなくプール化ストリームを返すためである。H100 80GB 上の再現は 4/5、5/5、5/5 の破損から、3 つのアダプタそれぞれ 0/5 に変化し、新規テストファイルは 10 項目通過。作者は wait_stream を追加する書き方を明確に否定している。それでは毎回の呼び出しで起動ストリームが現在のストリームの全処理の後ろに並んでしまうからだ。
出典:https://github.com/tile-ai/tilelang/pull/3649
1.2 メインリポジトリ:NVRTC デバイス側コンパイルを TVM-FFI 実行パスに接続、コンパイルキー項目を追加(#3285)
10-27 00:02 にマージ、12 ファイル +133/-133。tl.cuda_compiler="nvrtc" を新規追加し、既存の TVM-FFI ホストランタイムと組み合わせる。NVRTC がデバイス側 CUDA ソースを cubin にコンパイルし、共有 CUDA ランタイムモジュールがこれらのバイトをインポートする。ホスト側のバインディング、スカラー変換、条件、ループ、起動準備は引き続き TVM-FFI が生成し、デフォルト値は依然 nvcc、旧 execution_backend="nvrtc" アダプタは元の実装を維持する。バイナリキャッシュキーにコンパイラ名と NVCC/NVRTC バージョンを追加し、ターゲット詳細と有効なコンパイルオプションとともにハッシュに参加させる。アーキテクチャトークン 90a や 100f はそのまま保持し、非互換ターゲットはコンパイル前に診断を出す。正常に作成された各 NVRTC program は finally で破棄される。この PR の CI は H20 で 5506 件通過、491 件スキップ、サンプル 76 件通過、ROCm 7.2 チャネル 3220 件通過、昇騰チャネル 1797 項通過を記録した。
情報源:https://github.com/tile-ai/tilelang/pull/3285
1.3 メインリポジトリ:バッチ TMA ロードの cache hint にテンプレート引数が 1 つ足りず、DeepSelect の TMA パスがコンパイルできない(#3683)
10-27 00:11 にマージ、2 ファイル +20/-11。eviction_policy="evict_first" または "evict_last" のバッチ TMA ロードが main 上でコンパイル失敗する。コード生成が CacheHintSm90 テンプレート実引数を 1 つ発行するが、生ポインタ版の tl::tma_load オーバーロードはバリア型しか受け付けない。この欠陥は DeepSelect の TMA パスを直接押さえ込み、GB200 と CUDA 13.1 上で当該サンプルの ring および replay ケースで再現した。修正方法は #3296 からキャッシュヒント部分を抽出し、欠けているテンプレート引数を補って選択したポリシーを PTX 命令に渡す。共有メモリアラインメントの部分は引き続き #3296 に残す。修正後、同一マシン上で DeepSelect 全ファイル 144 項通過、TMA コピーテスト 22 項通過、CUDA 12.9 の SM90 コンパイルチェックでも 3 種のポリシーの PTX が適切なオペランドを備えていることを確認した(#3296 の残りの変更は未マージ)。
情報源:https://github.com/tile-ai/tilelang/pull/3683
1.4 メインリポジトリ:リダクションの FP8/FP4 非有限定数が不正リテラルを出力、サブバイトアキュムレータを早期拒否に変更(#3189)
10-26 15:44 にマージ、7 ファイル +204/-4、#2998 を修正。T.reduce_max と T.reduce_min が出力 FP8 E4M3 時に不正な inff リテラルを発行していた。非有限の E4M3、E8M0、FP4 定数は既存の変換コンストラクタ経由に変更し、正当な CUDART_INF_F と CUDART_NAN_F 式を渡す。飽和と NaN の挙動は維持される。同じ変更で通常リダクションと遅延リダクションがサブバイト数値アキュムレータ(FP4、int4 など)を早期に拒否し、明確な診断を出す。FP4 入力に FP16 または FP32 アキュムレータの組み合わせは引き続きサポートされ、ブールリダクションは免除、E8M0 リダクションのサポートは今回の範囲外。B300 と CUDA 13.1 上で 28 項のフォーカステスト全通過、完全なリダクションテストファイル 150 項通過、他に 3 項のバッチ検証失敗は main 上の既存問題。
情報源:https://github.com/tile-ai/tilelang/pull/3189
1.5 メインリポジトリ:SM120 ブロックスケーリング GEMM で K 原子を明示展開、FullCol 構成が 51% 向上(#3286)
10-26 16:36 にマージ、4 ファイル +308/-2。#3257 が T.gemm_blockscaled と T.mma_gemm_blockscaled を導入した後の仕上げ。A オペランドをレジスタフラグメントから取得する場合、K 原子ループは従来 T.serial で、各フラグメントアクセスが実行時添字となり、nvcc は自己展開しないパイプライン化・warp 特化カーネルでフラグメントをローカルメモリに残してしまう。つまり MMA パスを溢れさせる。T.unroll(..., explicit=True) に変更後、各アクセスはコンパイル時定数となる。第二の点は SFA と SFB がともに共有メモリにある場合、従来は MMA ペアごとに共有メモリを 3 回読んでいたが、現在はレーンごと K 原子ごとにコンパクト選択パケットを 1 回ロードするだけで、行主序とブロック K 主序の両レイアウトをカバーする。RTX PRO 6000 Blackwell 上で 143 項通過、4 項スキップ、出力は float32 リファレンスとビット単位で一致。レジスタ A ベンチマークでは FullRow 2.7% 向上、Square 11.1%、FullCol 51%(348.5 から 526.7 TFLOPS)、ptxas 報告の溢れ書きは 664 バイトと 484 バイトから 320 と 160 に減少。作者は同時に注記:block_K = 128 と行主序共有 scale のパスは main 上で元々検証を通らず、今回の変更とは無関係で、別途調査待ち。
情報源:https://github.com/tile-ai/tilelang/pull/3286
1.6 メインリポジトリ:JIT 出力形状内のスカラパラメータを割り当て前に解析(#3602)
10-26 10:06 にマージ、7 ファイル +242/-3、#3406 を修正。Cython および PTO ラッパーはこれまで、出力を割り当てる際に明示的なスカラ PrimFunc パラメータで表現された出力次元を解析できなかった。現在はアダプタが「出力次元からスカラパラメータ位置への」マッピングを渡し、ラッパーは実引数値を直接用い、重複する起動パラメータを追加しない。リグレッションは両バックエンド、ゼロ次元と非ゼロ次元、スカラと多次元出力、出力順序、デバイス配置、起動パラメータ個数をカバーする。同族の欠陥に関する后续の 2 件はなお審査中:#3636 は (n, 2*n) のような複合次元式を扱い、#3638 は同名だが互いに独立した動的シンボルが名前で統合される問題を扱う。
出典:https://github.com/tile-ai/tilelang/pull/3602
1.7 メインリポジトリ:sm8x スパースメタデータのオペランド dtype がフォールバックなしで通過しなくなる(#3335)
10-26 16:54 にマージ、1 ファイル +12/-5。スパースメタデータのレイアウトコンストラクタはアーキテクチャごとにディスパッチし、MMA オペランド dtype とメタデータバッファ幅を対応付ける。sm90 分岐はサポートされない dtype を拒否するが、sm8x 分岐は 2 つのホワイトリストチェックとして書かれておりフォールバックがなく、そのため他の綴りはすべて下の配置ロジックに落ち、そのロジックはバッファ幅からのみ group と interweave を推論する。結果として float8_e4m3fn(torch の正規の綴り)が 16 ビットメタデータとして受け入れられ、fp16 の group 32 と interweave 4 の規約が適用された。未認識の綴り(タイプミスを含む)も任意の幅で受け入れられた。現在はサポートされない dtype は直接エラーとなる。
出典:https://github.com/tile-ai/tilelang/pull/3335
1.8 メインリポジトリ:DeepSelect Top-K サンプルをマージ、129 組の設定で torch.topk に対し最大 8.00 倍(#3304)
10-26 16:53 にマージ、4 ファイル +1265/-0。サンプルは DeepSeek の DeepSelect を TileLang で書き直したもので、インライン PTX を含まない。著者の実測では 129 組の設定、10 組が使用不可で、RTX 5090(170 SM)に CUDA 12.9 と PyTorch 2.7 を組み合わせ、重複入力とキャッシュクリアの 2 手法で実行した。重複入力の条件では、BF16、バッチ 6、N が 16384、K が 512 の組でレイテンシ 5.059 マイクロ秒に対し torch.topk は 40.448、加速 8.00 倍;N が 1048576 の組は 15.503 対 100.254 で 6.47 倍。FP32 と K が 4096 の組み合わせはこのパッケージでは使用不可、K が 1 のネイティブパスにはアラインメントアドレスエラーがあり、著者は表でゼロではなく「使用不可」と横線で注記している。このサンプルが依存するキャッシュヒントの修正は、まさに同日マージされた #3683 である。
出典:https://github.com/tile-ai/tilelang/pull/3304
1.9 審査中:TIRx リファクタリングの規模が 1194 ファイルに再拡大、未パッチの上流バージョンにアンカー(#3666)
10-25 に作成、10-26 に 3 度目の更新、現在 1194 ファイル +106648/-104497。方針は TileLang を内蔵 TVM から独立リリースの TIRx へ移行し、固定した上流ベースラインにゼロパッチを当てること:3rdparty/tvm とそのサブモジュール項目を削除し、3rdparty/ には CUTLASS、HIP ヘッダー、無視ファイルのみを残す;未修正の apache/tvm-tirx ソースを 9af5544d に固定し、apache-tvm-tirx==0.1.0a2 と apache-tvm-ffi==0.1.15.post0 に依存、従来の +tilelang.1 サフィックス付き統合パッチは不要となる。これは現在のキューで最大規模の一笔であり、今後の TVM 側修正(例えば同日の tvm フォークリポジトリの Z3 ブールビット演算)がどのようにメインラインへ入るかも決める。
出典:https://github.com/tile-ai/tilelang/pull/3666
1.10 新規欠陥:ThreadSync の 3 種類のミスアラインされたフェンス、審査中の 3 件の修正を伴う(#3680 / #3682 / #3681)
#3680 は 10-26 19:56 に起票され、tl.ThreadSync が 3 箇所でフェンスを誤って配置していると指摘した。スレッドごとの値を取るためのガードをブロック内で一致していると見なしたため、フェンスが分岐内に残り、分岐をスキップするスレッドを待ってしまう。if の某一分岐内のフェンスが外側シーケンスの保留依存を消してしまい、ブロックがその分岐をスキップする場合に if 前の書き込みと if 後の読み取りが衝突する。ループの行程数がスレッドごと(例えばスレッド番号による剰余)である場合、フェンスがループ体内に残り、部分同期の下降が直接「スレッド矩形を証明できない」で失敗する。#3682(3 ファイル +39/-7 レベルの変更、12 項の新テスト)は一巡の前置解析でスレッドごとの値を取り得るバッファを収集し、同期要約を分岐とループの到達可能性に応じて縮小する。作者はこの変更が競合を報告するだけで競合を排除せず、分岐分割も行わないと宣言している。#3681 は別の同源のクラッシュを修正する。sm_100 以降で float32 ループが 256 ビットにベクトル化され、8 レーンにブロードキャストする際に生成された式が一時オブジェクトのアドレスを取ってしまい、NVRTC が「左辺値でなければならない」としてコンパイルを拒否するため、__float_as_uint とシフトで直接ビットを組み立てるよう変更した。併せて tvm フォーク PR #88 がブールビット演算を論理演算に翻訳し、証明器が 10000 ステップの予算内で同一の結論を出せるようにした。
出典:https://github.com/tile-ai/tilelang/pull/3682
1.11 新規欠陥:sm_100 以上のリダクション batch とレイアウトの衝突、審査中の修正が 1 件(#3667 / #3684)
#3667 は 10-26 08:31 に起票された:T.reduce_* は sm_100 と sm_120 上で batch が N に等しい場合に下降できない。フラグメントコピーが sm_100 以上で 256 ビットロードを通るため、スレッドごとに保持する出力数が減り、sm_90 では合法だった batch がハード検証に引っかかる。#3684 は 10-27 05:44 に起票され修正方法を添付した:ReduceOp は N の batch を超えない最大の真因数を取って警告を出し、finalize_reducer はそのテストが挙動を固定しているため厳格のままとする。新規の sm_100a と sm_120a の下降テストは main で失敗し、本ブランチで通過する。実測機は RTX 5070 Ti。同じリダクション問題の別件 #3341(バッチ全リダクションのスレッド範囲解析)は依然審査中。
出典:https://github.com/tile-ai/tilelang/pull/3684
1.12 審査中:転置ターゲット形状の検証、領域化 kill の再配置(#3677 / #3670)
#3677 は tl.transpose のターゲット形状検証を TransposeNode の構築に組み込み、NormalizeToAccessRegion の後に実行するため、Python ラッパーを迂回して直接構築してもカバーされる。相等性は Python の次元リストを比較するのではなく arith::Analyzer::CanProveEqual で判定し、(3, n+1) と独立に構築した (n+1, 3) は依然受け入れられ、定数の不一致は軸番号と両側の長さでエラーを報告する。従来は不一致のターゲットが越界インデックスに下降し、誤った結果や越界書き込みを生じていた。CPU ターゲットの 13 項のテストが通過。#3670 は共有メモリの生存解析に残っていた scope_level の仕組みを置き換える。この値は SeqStmt を含む For 体内でのみ増加し、While では決して増加しないため、両方向で誤る——While の場合は kill 点がバックエッジを越えず、パッカーが互いに排他的な領域を重複させ、sm90 での実測結果が毎回倍増する。兄弟ループの場合は kill が領域境界全体を越えてしまう。新規則は実際の外側 For と While のチェーンおよび各ループの連続区間で計算し、越界と早期停止の両方向が挟み込まれる。
出典:https://github.com/tile-ai/tilelang/pull/3677
1.13 クローズ済み:ファジングシリーズのコンパイル時クラッシュ 6 件(10-26)
10-26 にファジング由来のコンパイル時クラッシュまたは誤ったコードクラスの欠陥 6 件を一括でクローズした:int16 と uint16 ターゲット上の T.atomic_add が nvcc に 16 ビットオーバーロードがないため直接中止する(#2572)。非ブールのバッファ要素を直接 if や assert の条件とする場合に内部の is_bool() アサートで中止する(#3009)。浮動小数点オペランドを整数のみ受け付ける div と mod 族に渡した場合も同様に内部アサート(#3008)。int2、uint2 のような 1、4、8、16、32、64 の列にない整数型幅で CUDA コード生成が中止する(#2996)。reduce_max と reduce_min を FP8 または FP4 のスライスに書き込むとクラッシュする(#2998)。そしてバッファ全体を目的地とするアトミックリダクションが 0 番スレッド以外の全スレッドを黙って捨てる(#2959)。最後の 1 件にはドキュメント側の修正が別途マージされ(#3340)、例の中の目的地をバッファ全体から名前付き要素に変更し、例自体が誤った書き方を教え続けることを避けた。
出典:https://github.com/tile-ai/tilelang/issues/2959
2. マルチバックエンド適配(昇騰 / 海光 / 沐曦)
2.1 昇騰メインリポジトリ:バージョン付きオンチップバッファの精確アドレッシングインターフェース(#3607)
10-26 10:04 にマージ、14 ファイル +982/-26。新規 T.annotate_buffer_addresses を追加し、スケジューリング後の昇騰パイプラインにおいて、UB、L1、L0A、L0B、L0C と SS の各プール内の割り当てを正確なバイトアドレス指定できるようにした。アドレス指定を入れるバッファは、まず T.annotate_buffer_versions で固定コピー数を宣言する必要がある。検証はバージョン展開とレイアウト補完の後に行われ、コピー数の欠落や非対応、重複、アラインメント違反、動的サイズ、容量オーバーフローをカバーする。アロケータはさらに、SIMT の UB コンテキスト用の予約と AIV プライベートアドレス空間への譲渡を行い、L0 の scale factor ハンドルはそのデータ割り当てに追従する。部分マッピングは指定された区間だけを固定し、残りのバッファはアラインメント後の空き空間に配置され、既に占有されたアドレスは再利用できない。アドレス指定アノテーションのないプールは従来の動作を保つ。サブコア投影時は、アクティブなストレージコピー上のアノテーションのみを保持する。
出典:https://github.com/tile-ai/tilelang/pull/3607
2.2 昇騰審査中:手動スケジューリングの FlashMLA、BF16 プリフィルと FP8 デコード(#3678)
10-26 19:31 に起票、Ascend 950 上の BF16 プリフィルとコンパクト FP8 デコードをカバーし、オプションの FP8 と FP4 の追加 KV cache、attention sink、不正インデックス、クエリごとの長さを含む。アテンションループ、オンライン softmax、逆量子化、スケジューリングはすべて TileLang で表現し、小さな intrinsic ヘッダのみが単一のハードウェア命令をラップする。手動実装は AIC と AIV のパイプラインと同期を明示的に制御する。永続化バージョンは T.Stage とコンパイラ管理のバッファバージョン・同期を使い、QK と PV は同一の L0B ストレージリングの異なるレイアウトビューを使用する。
出典:https://github.com/tile-ai/tilelang/pull/3678
2.3 昇騰審査中:オンチップメモリのオーバーフローを事前拒否、E4M3 の UB から L1 へのコピー、ビューレイアウトとポインタ単位(#3655 / #3639 / #3623)
#3655 は MergeUBAllocations が下降期に目標物理容量に基づいてオーバーフローを拒否するようにする:2 つの独立したループがそれぞれ合法的に 32 KB の L0A スライスを 2 つ選択し、マージ後に 64 KB がちょうど満杯になる場合、以前は物理バッファの末尾を越えるポインタオフセットをサイズメタデータなしで発行していたが、現在は短い診断メッセージを返す。#3639 は E4M3 の UB から L1 へのコピーを有効化する:ND 入力は既存の scatter とポストコピーパイプライン経由で自動再配置されるか、直接事前配置済みの NZ 入力へコピーされ、既存の幾何チェックを再利用し、フットプリントとレイテンシ分析に FP8 要素サイズの処理を追加する。#3623 は T.view と T.reshape が生成するバッファがそれぞれレイアウトを保持するようにし、共有データ変数経由でレイアウトを解決する際に誤った Cube マッピングや DMA アドレスを選択するのを避け、生ポインタオフセットを自身の要素単位で再マッピングする。
出典:https://github.com/tile-ai/tilelang/pull/3655
2.4 昇騰審査中:バッチスケジューリングのオプション化、クロスコア同期要件のバッチモード、L0C デュアルコピー直接変換(#3675 / #3614 / #3468)
#3675 は昇騰にオプションのバッチスケジューリング(schedmode 1)を追加する:通常モードではスケジューラは必要なコアリソースが部分的に利用可能であればオペレータを起動し、複数ストリームの全コア同期オペレータが交互に実行される可能性がある。バッチモードは全リソースの準備が整ってから全体を起動し、全コア同期シナリオではこれを使う必要がある。#3614 は同じメカニズムの mode 0 における症状と修正方法を記述する。#3468 は A5 上で FixPipe がデュアルターゲットモードと quant_pre を同時に有効化できない制限を処理し、T.dual_copy が L0C の fp32 から f16/bf16 への変換を直接行うようにし、Vector 変換 1 回と fp32 の UB バッファ 1 塊を省略し、L0C から UB への fp32 帯域も省略する。
出典:https://github.com/tile-ai/tilelang/pull/3675
2.5 昇騰リポジトリ:インストーラの SHMEM セクション二重修正、空殻 wheel が偽の成功をしなくなる(#1943 / #1946)
2 つの審査中 PR が同じ欠陥(#1728)を指す:install_ascend.sh --enable-shmem は build.sh の失敗後に戻り値をチェックせず bdist_wheel を続行し、package_data のワイルドカードがファイルにマッチしなくてもエラーを出さないため、__init__.py だけの空殻パッケージを作成して Successfully installed shmem-1.0.0 と表示し、20 数行後に source ../../../install/set_env.sh の箇所で「shmem 環境の設定に失敗」と報告し、トラブルシューティングの方向を環境設定に向けさせてしまう。#1946 は 2 箇所だけ変更する:戻り値チェックを追加して即座に終了し、set_env.sh をスクリプトディレクトリベースの絶対パスに変更し、shmem セクションを逐語的に抽出する制御されたハーネスで検証する(修正前は 10 項中 6 失敗 4 通過、修正後は 10 項すべて通過)。#1943 は別の道を取る:既存パッケージをアンインストールする前に、ネイティブビルド、環境スクリプト、wheel ビルド、wheel の存在を順にチェックし、10 項の CPU 制御フローテストを追加し、これが修正するのはインストーラの継続行と診断であり、上流の SHMEM コンパイル失敗自体を含まないことを説明する。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1946
2.6 昇騰リポジトリ:W4A4 の int4 行列積で結果が誤り、根本原因は半ニブルパッキングを 1 バイトとして扱っていたこと(#1731)
#1731 は 10-26 に更新され、特定と再現がかなり完全な形で示されている。C++ テンプレートスタック全体が sizeof(AscendC::int4b_t) == 1 でフラクタル幾何を導出しているが、この型は半ニブルパッキングであり、32 バイトの C0 ブロックの正しい要素数は 32 ではなく 64 である。そのため 3 層の転送すべてに 2 倍のずれが生じる:L1 から L0A/L0B への LoadData パラメータが誤った定数で計算され、データがオンチップ上で誤った位置に積み重なる;B オペランドの転置ロードは load_cbuf_to_cb_s4 によってハードコードで無効化され、フラグが黙って破棄されるため、B は mad_s4 が期待する nZ レイアウトを永遠に得られない;UB と GM 間のバイト長は要素あたり 1 バイトで計算され、各行が 2 倍コピーされ隣接行が混入する。計算命令自体は正しいため、症状は「桁は正常、数値は全滅」(元報告では 99.6% 不一致)となる。修正はすべて src/tl_templates/ascend/common.h と 2 か所の変換に収まり、W4A4 の例と 10 項目の int4 回帰が新規追加され、1024 キューブのエンドツーエンドは 99.6% 不一致からゼロ不一致となり、既存の gemm_v0 スイート 48 項目、l1_to_l0 と gm_to_ub 263 項目、要素ごと 492 項目がすべて通過した。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1731
2.7 昇騰リポジトリ:2 次元を超える GM から UB へのコピーで DMA 行数に先頭次元を折り畳む(#1805)
#1805 は #1231 を修正する。copy_gm_to_ub と copy_ub_to_gm の DMA モデルは 2 次元矩形であり、AscendCopy::Lower は region のアクティブ次元から最後の 2 つだけを行と列として取るため、アクティブ次元が 2 つを超えると先頭次元が黙って破棄される:3D ブロックコピー A[t:t+16, :, :](region は 16、4、8、計 512 要素)は 4 と 8 に降格され、DMA は 32 要素しか転送せず、UB シャードの残り 15 トークンは未初期化となり、後続の計算がゴミ値を読む。修正方法は先頭次元を DMA の行数に折り畳むことで、両方向とも同じく変更される。同日、昇騰リポジトリにはドキュメントのファクトチェック PR の更新が別途一批あり(sin と cos、axpy、clamp、bitwise シリーズ、clear、fill)、うち clamp の分はついでに PTO 上の T.tile.clamp の下限が無効となる問題も修正した。
出典:https://github.com/tile-ai/tilelang-ascend/pull/1805
2.8 昇騰の新規不具合:Cython ラッパーが呼び出しごとにアナライザを生成、約 0.28 ミリ秒(#1945)
10-26 15:52 に起票され、報告者は層ごとの特定と因果検証を示している。CythonKernelWrapper.forward() は呼び出しごとに無条件で TVM の Analyzer() を構築し、ローカル中央値は約 0.28 ミリ秒だが、これが使われるのは動的 shape 分岐内の 1 回だけである;静的 shape カーネルの dynamic_symbolic_map は空辞書であり、simplify() は決して実行されない。因果検証の方法は、モジュール内の Analyzer グローバルを「simplify を呼ぶと即座にエラーを投げる」スタブに差し替えることで、すべての呼び出しは通常どおり成功し、全呼び出しの所要時間は 442〜467 マイクロ秒から 104〜107 マイクロ秒に低下し、かつカーネル規模に依存しない(1 ブロックと 32 ブロックが同程度)。実測環境は Ascend 910B3 と CANN 9.2.0。これは小規模オペレータを密に呼び出す推論経路にとって固定オーバーヘッドであり、修正方針の行方が注目に値する。
出典:https://github.com/tile-ai/tilelang-ascend/issues/1945
2.9 海光:レジスタパイプラインと warp ダイバージェンスをマージ、28 ファイル(#18)
10-26 13:33 にマージ、28 ファイル +5312/-64、10-08 に起票され 2 週間余りで収束。変更は tilelang-hygon にレジスタパイプラインを追加し、既存の共有メモリパイプラインと協調でき、対応するコミットとバリア命令がパイプライン段に応じて挿入される;warp 数が 4 を超えると warp ダイバージェンスが有効になる。作者の結論は example_gemm が rocblas と同等の性能に達するというもの。同じリポジトリで当日注意すべきは #18 のマージ順序である:これは #28 より先にマージされ、主干の変更をブランチに取り込んだ後の大きなマージに属するため、行数が #28 よりはるかに多い。
出典:https://github.com/tile-ai/tilelang-hygon/pull/18
2.10 海光:ネイティブ B16 ストレージと動的 WASP GEMM をマージ、scale バッファのデュアル bank は審査中(#28 / #29)
#28 は 10-26 16:23 にマージ、6 ファイル +340/-230、三つの事項:MLS の Let パラメータ発射順序を保持し、GEMM の warp 分割を行列ストレージへ伝播し、パックされた B16 から B32 への出力エイリアスを拒否する。同時に WASP 次元を動的にし、軸ごとに対齊を推論し、pong の K 末尾フィルタを修正し、WASP のエンドツーエンド回帰を基本 MLS テストから分離した。#29 は 10-26 17:13 に起票され審査中:scale バッファの行計画に LDS bank と分かれたデュアル bank パーティションを設け、割り当て、scale コピー、MMAC 下降が同一の行マッピングを共有し、物理区間のライフサイクルと m0 対齊ページングを保持し、アイドル wave の作業を二つの bank に分散し、手動 warp 特化下のタイル化スケーリング GEMM に用いる。
出典:https://github.com/tile-ai/tilelang-hygon/pull/28
2.11 沐曦:組織ビューにプッシュ記録があるが、六つのブランチはいずれも調査期間内の新規コミットなし(10-26)
組織ページは tilelang-metax が調査期間内にプッシュのタイムスタンプを持つことを示している(北京時間 10-26 15:08)が、ブランチごとに照合した結果、対応するコミットは見つからなかった:デフォルトブランチ main の最新は依然 2026-09-29 のバージョン番号引き上げ、dev ブランチの最新は 2026-10-16 のマージ、0.1.14 ブランチは 2026-09-17、0.1.12 と race はさらに早い。イベントストリームにも調査期間内のプッシュイベントはない。これに基づき「実質的な変更なし」として処理し、本期の条目には計上せず、その性質も推測しない。当該リポジトリの直近の正式リリースは依然 10-16 の 0.1.15 である。
出典:https://github.com/tile-ai/tilelang-metax
3. エコシステムと採用側(TileOPs / TileFoundry)
3.1 TileOPs:初回リリースは依然最終マス、成果物は依然 2025-10-09 の開発版(#2738)
初回リリース追跡チケットは 10-26 に更新され、八項目のクリーンアップのうち六項目が完了した:命名とマニフェストパラメータ、パッケージングとリポジトリファイル、Python 3.13、公開 API と拡張ポイント、オペレータの動作規約、入力検証生成、ドキュメントサイト。残り二マス:夜間ベンチマーク(壊れた二つのベースラインを修正 #2763 はマージ済み;新しい行をすべての比較器の上へ引き上げ、累積スキャンを単回起動に変更 #2766 はマージ済み、GQADenseWithLSEFwdOp の D128 プログラム #2774 はマージ済みだが、MHA 行は依然 FA3 の 0.996 のみで、チューニングを継続)とリリースそのもの(第三者声明、PyPI インストールドキュメント、開発者依存グループ、CHANGELOG とリリース手順 #2753、およびメンテナの実行が必要な TestPyPI 試走と最初の tag)。PyPI とリポジトリを照合:tileops の最新は依然 2025-10-09 の 0.0.1.dev2 で、全三バージョン、リポジトリにはいかなる tag と release もない。初回リリースは「ドキュメント準備完了、成果物未アップロード」の最終マスにあり、既に複数日続いている。
出典:https://github.com/tile-ai/TileOPs/issues/2738
3.2 TileOPs:D128 稠密フォワードで key を分割、prefill-rope が 31.04 マイクロ秒から 17.09 へ(#2777)
10-26 22:45 にマージ、2 ファイル +527/-124。D128 稠密呼び出しで 128 行のタイルが SM を遊ばせる場合、各タイルの key を min(SM 数割るタイル数, 最短タイルの key ステップ数) 份に分割し、かつ分割が最長タイルから二ステップ分の key を均せるときのみ分割する;各份は本行のスケーリング最大値と和を持つ未正規化 fp32 出力を書き出し、マージカーネルは最大の最大値で重み付けするため、相殺項と大きな logit も精度を失わない。タイル最後のステップ key において、ある warpgroup が最後に必要とする key が低い 64 個の key 内に収まる場合、softmax はその 64 個のみにマスクと指数を適用し、warpgroup 0 の因果対角ステップは必ず命中する。H200 1500 MHz 上で prefill-rope の f16 と bf16 は 31.04 と 31.01 から 17.09 と 17.09 マイクロ秒へ低下し、FA3 に対して 2.324 と 2.328 倍;qwen3-8b のチャンクプリフィルは 31.30 から 17.79 へ低下。テストノードは純増 14。
出典:https://github.com/tile-ai/TileOPs/pull/2777
3.3 TileOPs:D128 永続化フォワードをスレッドごとに書き換え LSE を出力(#2774)
10-26 18:22 にマージ、2 ファイル +656/-48。根本原因は GQADenseWithLSEFwdOp が以前、非永続化の warp 特化カーネル上で動作していたことで、永続化カーネルが return_lse を拒否するため、夜間ベンチマークの実測では FA3 の 0.73 から 0.78 しかなかった。今回、TileLang のスレッド単位プリミティブ(TMA マッピング、mbarrier、WGMMA 記述子、warp shuffle、stmatrix)で D128 永続化稠密フォワードを書き直し、log-sum-exp を直接書き出す。同一の実装が GQADenseFwdOp と GQADenseWithLSEFwdOp の両方に同時にサービス提供する。すべての key と value が L2 の半分を超える場合、1 ヘッドのクエリブロックを隣接配置し、並行タイルが L2 内の K と V を共有できるようにする。各タイルの最初の QK は前のタイルのエピローグ内で起動される。公開インターフェースとマニフェストは変更なし。マージ後、MHA 行は依然として FA3 の 0.996 であり、追跡単标注チューニングは継続中で、これは #2777 とは 2 本の独立した最適化ラインに属する。
出典:https://github.com/tile-ai/TileOPs/pull/2774
3.4 TileOPs:融合加算の LayerNorm 行を CTA 横断で分割、精度欠陥は別チケットに(#2768 / #2773)
#2768 は 10-26 15:45 にマージ、5 ファイル +525/-24。出発点は夜間ベンチマークのリグレッション(llama-70b デコード行が 0.0037 から 0.0044 ミリ秒に上昇)で、調べたところコード変更ではなくチューニングの再選択であった。チューニングキー変更後に FusedAddLayerNormKernel が再チューニングされ、チューナが L2 バックツーバック計時の候補から 1024 スレッドを選択し、コールドスタートが 256 スレッドより 19% 遅かった。新たに書かれた FusedAddLayerNormHeldKernel は 1 行(SM90 では 1 つの cluster CTA のシェア)をレジスタに保持し、最も忙しい SM のコストがより低い場合に 1 行を最大 8 つの CTA に分割し、ブロックサイズは全 warp 数のうちシェアのアクセス量を整除できるものを取る。ディスパッチは「キュー内の行数」と「レジスタが 255 を超える」という互いに拒否する 2 条件で選択し、チューニングはチューニング対象ビルドのシェアを埋められる幅のみを提供するため、デコード行が再選択されることは不可能になった。llama-70b デコードは 2.368 マイクロ秒で flash-attn の 4.800 に対して。同一の変更で露呈した精度問題は別途 #2773 を起票:行がゼロから遠い場合、fp32 合計の平均と分散が精度を失う(定数 10000.1 の行は正規化後におよそ 0.7 となりバイアスにならない)。1e36 付近では二乗和がオーバーフローして全体が NaN となり、著者の提案は Welford またはスケーリング 2 パス法への変更である。
出典:https://github.com/tile-ai/TileOPs/pull/2768
3.5 TileOPs:要素単位とリダクションの 4 本、clamp の bf16 行が 5.9% 高速化(#2766 / #2767 / #2769 / #2770)
4 本が同日マージされ、いずれも「まず根本原因を特定してから変更する」に属する。clamp(#2767)は自身の精度内で比較するように変更された。根本原因は旧実装が各要素を fp32 に拡張し、各境界で NaN 判定と選択を 1 回ずつ行い、要素あたり約 5 命令であったこと。現在は要素ペアあたり 1 命令の HMNMX2.NAN となり、16384 平方の bf16 が 0.25704 ミリ秒から 0.25024 に低下し、torch-compile に対して 1.0587 倍となり、全 16 ビットモード、NaN、正負の無限大、符号付きゼロの境界で torch の fp32 比較とビット単位で一致することを検証済み。累積スキャン(#2766)は dtype=float32 の cumsum と cumprod を「まず 1 コピーに拡張してからスキャン」から 1 回の起動で直接ターゲット精度に書き込む方式に変更し、hidden-state-f32 ワークロードで転送量が 112 MiB から 48 MiB に減少、cumprod は 47.36 から 14.62 マイクロ秒、cumsum は 47.86 から 14.62 に。GELU 乗算(#2769)は共有 gelu() を経由して fp32 で評価し 1 回だけ丸めるように変更された。旧 erf パスは NaN 選択分岐、3 回の 16 ビット丸め、より長い依存チェーンを伴っていた。小ゲート付きグリッドのブロック数は 56 から 14 に減少。二項演算(#2770)は 2 バイトオペランドのスレッドあたり要素数を 8 から 16 に引き上げ、2 本の 16 バイトロードを in-flight に保ち、34 行のデフォルト設定がそれに伴って変化し、MinimumFwdOp の llama-7b 行は 8.42 から 8.00 マイクロ秒に低下した。
出典:https://github.com/tile-ai/TileOPs/pull/2767
3.6 TileOPs:呼び出しテンソルのアライメント判定をオペレータ層に移動、3 本で締めくくり(#2772 / #2778 / #2780)
#2772 は 10-26 15:28 にマージ、55 ファイル +373/-224:実行器は実装の実行前に、ストレージ起点が 16 バイト境界にない呼び出しテンソルをアラインメントされたストレージへコピーし、成功後に書き込まれた那几个をコピーし戻し、呼び出し元のテンソルを返す。コピーされたテンソルが他の実引数と重複する場合(かつ少なくとも一方が書き込み対象である場合)、コピーがカーネル自身のチェックを覆い隠すため、実行前に直接拒否する。アドレスのないストレージ(偽テンソルや読み取りを拒否するラッパー)はコピーしない。カーネル側のアラインメントと一時退避ロジック(MoE グループ化 GEMM、因果畳み込み状態、融合 QK-norm RoPE の qkv を含む)はすべて削除された。代償として、ホットパス上のアラインメント呼び出しのホスト時間が 1 から 2 パーセントポイント増加し、MoEPostPermuteFwdOp で out を伴う場合は 4.4 パーセントポイント増加する。#2778 はこの判断を生成されたオブザーバに折り込み、計画キャッシュにコピーを決定させ、ホットパスはその決定を読むだけにし、実行器の従来の独立したアドレス走査を削除した。#2780 は 10-26 に起票:現状では四つのカーネルのみが out と入力の重複を検査し、残りの out を伴うオペレータは静かに競合し、しかも各検査は毎回 data_ptr を重複して読む(MoEPostPermuteFwdOp は一回の呼び出しで 12 回多く読む)。方向性は、リスト内でオペレータごとに重複を許可するかどうかを宣言し、オペレータ層で統一的に検査することである。
出典:https://github.com/tile-ai/TileOPs/pull/2772
3.7 TileOPs:オフセットストレージの三種類の修正、SM90 グループ化 GEMM と出力 ABI(#2761 / #2765 / #2762)
#2761 はグループ化 GEMM の MMA 実装が起点にオフセットを持つ連続テンソルを受け入れ、ベクトルアクセスをアラインメントし、結果をオフセット出力に書き戻す際に識別子と未タッチのエキスパート行を保持し、コピー前に重複を検査するようにした。元の実装ではこの種のビューに遭遇すると CUDA アドレスの非アラインメントで失敗していた。新たに 16 個の回帰を追加し、FP16 と BF16 それぞれの a、b、bias、out のオフセットをカバーし、コンパクトとマスクの二種類のレイアウトを用意した。#2765 は同じ能力を SM90 の MoE グループ化 GEMM に補完した:ストレージ起点が 16 バイト境界にない場合に a、b をコピーし out を一時退避し、vector_aligned は今や「書き込まれるテンソルは一時退避し起動後にコピーし戻す」という規則を明記し、BF16 と FP16 の MMA カーネルはこれを使うように変更され、FP8 カーネルは SM90 FP8 カーネルから継承する。#2762 は出力 ABI を修正:TVM FFI アダプタが被呼び出し側による出力の割り当てを宣言する場合、TileLang 課実行体が自ら割り当てて返すようにし、他のアダプタは元の割り当て経路を維持する。報告された回帰は、旧実装では新しいビルド上で 27 個の Mamba ケースが全敗し、候補実装は公開呼び出し経路で通過するというものである。
出典:https://github.com/tile-ai/TileOPs/pull/2765
3.8 TileOPs 審査中:FP8 クエリのページドプリフィルを e4m3 で value を乗算するよう変更(#2775)
10-26 17:20 に起票、10-27 07:00 時点でも更新中。GQAPagedFP8WSKernel は FP8 クエリの value 乗算を e4m3 内で行うよう変更した:変換を担う warp は各 value タイルを K 主順の e4m3 オペランドに転置し(ldmatrix 転置、ワード単位の一度のバイト置換、stmatrix)、各 16 内で key の順序を整え、各スレッドのスコアがまさにそれが供給する確率になるようにし、さらにレジスタ内パッキングで確率を e4m3 に丸める。softmax は FA3 の FP8 形式に従う:2 の 8 乗のオフセットにより小さな確率が e4m3 の非正規数に入らないようにし、行最大値はより大きなスコアごとに移動し、key タイルは最後から逆順にスキャンする。H200 上で llama-8b-fp8-prefill は 369.3 から 339.8 マイクロ秒に低下(FA3 比で 1.000 から 1.087)、qwen3-fp8-ragged-odd は 130.9 から 123.9 に低下。著者は精度許容差を 2 の負 4 乗から 2 の負 3 乗に緩和したことを正直に明記しており、その理由は e4m3 確率自体に 2 の負 4 乗の半分のオーダーの丸めがあり、FA3 も同種の入力で同様に超えるためである。
出典:https://github.com/tile-ai/TileOPs/pull/2775
3.9 TileOPs 審査中:疎な長序列の GDN プリフィルの分割スキャン(#2779)
10-26 に起票。SM89 上で疎な長序列 GDN プリフィルのグリッドに対し、既存の状態スキャンと独立出力プログラムを再利用する:適用条件は密で key 主順の D128 入力、少なくとも 16384 トークン、かつ状態スキャン CTA が SM 数の四分の一を超えないことであり、他の呼び出しは元の実装を維持する。分解方式はプログラムキャッシュ識別子に書き込み、末尾ブロックと初期状態の常駐回帰を追加した。実測(CUDA 13、PyTorch 2.13、TileLang 0.1.15)では 32768 長 bf16 が 1.0384 ミリ秒対 FLA の 1.0656、65536 が 2.1178 対 2.1786、16384 では依然として FLA の方が速い(0.4695 対 0.4409)。著者は、利益はこの四つの value head の組でのみ成立し、八つと十六の value head の既存モデル形状ではこの経路を選ばず、候補も一般に FLA より速いわけではないと明確に述べている。
出典:https://github.com/tile-ai/TileOPs/pull/2779
3.10 TileFoundry:動的共有メモリをカーネルの実際のビューに応じてサイズ決定(#271)
10-26 に起票され審査中、12 ファイル +340/-32。合成されたホスト入口は以前、各 TIR カーネルに dynamic_smem=0 を渡していたため、整数の TensorView オフセットで共有メモリにアクセスするカーネルは動的共有メモリなしで実行され、48 KiB を超えても cudaFuncAttributeMaxDynamicSharedMemorySize を昇格させることがなかった。作者は目標容量ではなくカーネルが実際に到達する最遠のバイト数でサイズを決めることを選んだ。理由は二つ。CUDA は静的共有メモリと動的共有メモリを同じブロックあたりの上限で計算するが、このランタイムはいくつかの静的共有メモリ(ProgramMetaData または共有の割当テンソル)を自前で持っており、満容量で申請すると起動できなくなる。ブロックあたりの申請量は一つの SM にいくつ配置できるかも決め、sm_89 で 99 KiB を申請するとたとえ 512 バイトしか使わなくても一つしか配置できない。コード生成も同時に、最初の文の前に一度だけ smem_base を宣言するよう変更した。loop_outputs の中では以前、最初のループのブロック内に置かれていたため二番目のループから見えず、nvcc が未定義識別子を報告していた。H200 と sm_89 で二つの GPU テストが通過し、パッチなしでは前者が不正メモリアクセス、後者がコンパイル失敗となる。
出典:https://github.com/tile-ai/TileFoundry/pull/271
3.11 TileFoundry:ウィンドウ全体をオペランド自体で下ろし、sm_89 の GEMM がコンパイル可能に復帰(#273)
10-26 に起票され審査中、5 ファイル +54/-24。schedule finalize が sm_89 GEMM に対して書き出すコードがコンパイルできない。下降がアキュムレータ全体をゼロから始まりその全長を覆うウィンドウ T.ptr_of(acc[0:0+16, 0:0+8]) に包み、CUDA コード生成は分割レジスタタイルのポインタスライスに対して「通常のソースレイアウトが必要」と報告する。さらに下では MMA がレジスタ断片を lhs_frag_dst[0:0+16, 0:0+32] のような生のスライスとして読み、同様に処理分岐がない。これらのウィンドウは実際にはすべてオペランド全体である。修正は下降に置かれる。ConvertHIRToTIR が一度ウィンドウがそのオペランドを覆うか(開始が静的ゼロ、サイズがオペランド形状に等しい)を判定し、ウィンドウを構築する三箇所がその判定を共有し、最終的に T.ptr_of(acc) を書き出し、MMA が断片を直接読む。TIR ツインを持つ 49 個のスケジュールフィクスチャのうち 29 個の finalize 出力がこれに従って変化し、sm_89 GEMM と mesh ビューがコンパイル可能に復帰し、二つの wgmma 出力はポインタスライス拒否を越えて次のビューで停止し、その箇所のレイアウトは main 上でも同様に拒否される。
出典:https://github.com/tile-ai/TileFoundry/pull/273
3.12 TileFoundry の新欠陥:分析側が共有メモリのピークを過大計算し、TIR に伝えない(#272)
10-26 に起票、作者は #271 のレビュー中に発見した。二つのギャップはどちらも同じ結果を指す。分析が通るプログラムが起動時に拒否される、または最終 TIR がどれだけ共有メモリを必要とするかを自己記述できない。第一に、メモリ分析は smem を facts.memory.shared.per_cta(sm_89 で 101376 バイト)と比較するが、CUDA はこの上限を静的共有メモリと動的共有メモリの両方に作用させ、ランタイム自身は 12 バイトの __shared__ ProgramMetaData を保持している。この占有は runtime.cuh にだけ書かれており分析には知らされていないため、ちょうど 12 バイト以内の配置は分析を通るが起動時に拒否される。第二に、分析は CTA あたりの共有メモリ高水位を RegionMemoryMetadata.peaks に記録するが、ConvertHIRToTIR はそれを出力 TIR に書き込まず、PrimFunction.metadata は印刷も解析もされないため、finalize の後に再計算しなければならない(#271 はまさに全整数アドレスの共有メモリビューを走査して再計算していた)。また作者が書いたホスト入口が過少申請でも実行前に発見されない。
出典:https://github.com/tile-ai/TileFoundry/issues/272
4. コミュニティ、チュートリアルとイベント
4.1 言語サービス:実験的 CUDA デバッグが公開、cuda-gdb と Nsight に対応(10-26)
tilelang-lsp は 10-26 17:11 に最初のデバッグコミットをプッシュした(78a5761、タイトルは実験的 TileLang CUDA デバッグの追加)。内容は四つ。新規の tlsp.debug ランチャーを追加し、Python 行番号マッピング、-G と -lineinfo を付けたコンパイル、キャッシュの隔離、生成ソースの取得を行う。Nsight の cuda-gdb を VS Code に接続し、Python ブレークポイント、起動前チェック、生成 CUDA ソースの閲覧をサポートする。単体テスト、起動構成、オプションの GPU 受け入れテスト(保留ブレークポイント、デバイス側シングルステップ、スレッド検視、数値をカバー)を補う。中国語・英語のドキュメントと CI を同時に更新する。これはこのリポジトリの 8 月以来初の実質的コミットであり、直前の最終更新は 8-04 の CI テストだった。
出典:https://github.com/tile-ai/tilelang-lsp/commit/78a5761404
4.2 ドキュメントサイト:単日八件の更新(10-26)
tilelang.github.io は調査期間内に8回プッシュされた(北京時間 07:32、10:16、10:21、16:51、17:25、22:42、翌日 00:08、00:16)。コミットタイトルは統一して Update docs であり、メインリポジトリのマージに伴い自動同期されるドキュメントビルドに属する。同日にはサードパーティのローカライズリポジトリ(tilelang-docs-l10n)の1回のプッシュと、中国語 RISC-V コミュニティハブリポジトリへの初回コンテンツコミットもあり、ドキュメント同期の経路が正常に稼働していることを示している。
出典:https://github.com/tile-ai/tilelang.github.io
4.3 貢献フローガバナンスの4連:PR 説明の規範、純ドキュメントのバックエンドテストスキップ、Agent タイトル確認、レビューボットのノイズ低減(#3671 / #3672 / #3673 / #3674)
4件が同日にマージされ、いずれも「レビューコスト」を巡るもの。#3671 は CONTRIBUTING.md に、PR 説明には問題・方案・トレードオフを明記すべきこと、証拠はそれが支持する論点のそばに貼ること、欠陥修正には修正前後の挙動または回帰テストの意図と結果を書くこと、性能変更には比較可能な計測と条件を示すことを明記した。#3672 は純ドキュメント PR に対して CUDA、ROCm、Metal、Ascend、CuTeDSL のビルドとテストをスキップさせる。方法は既存 lint ジョブのチェックアウトでマージベースの差分を見るもので、変更パスがすべて Markdown、reStructuredText、または docs/_static/ 配下の場合にスキップする。ホワイトリストは意図的に狭めており、実行可能なサンプル、ビルドファイル、ワークフローは対象外とする。#3673 はコミット技能に確認ステップを追加する。最終 diff、検証、PR 説明が整った後、まず使用予定のタイトル原文を貢献者に見せ、確認または差し替えを求める。#3674 はレビューボットの自動要約と docstring カバレッジチェックを無効化し、共有レビュー指示を追加して、指摘を現在の diff と実際の呼び出し経路に紐付けること、回帰と既存の制約を区別すること、既存実装を比較してから修正案を出すことを求めた。
出典:https://github.com/tile-ai/tilelang/pull/3672
4.4 コミュニティの新規リポジトリ:RISC-V コミュニティハブリポジトリ、CANN オペレータリポジトリ、Hexagon バックエンドリポジトリ、デバッグツールリポジトリ(10-26)
調査期間内に登録に値する外部の進展がいくつか現れた。TongmingLAIC/tilelang-riscv は 10-23 に作成され、10-26 にプッシュされた。位置付けは RISC-V 上の TileLang コミュニティハブであり、コンパイラバックエンド、AI オペレータ、開発ツール、ドキュメントを集約する。現時点ではまだ空の骨組みだが、方向性は国内の RISC-V AI スタックの要求と一致している。Atlas-Ascend-Fork/ops-tilelang は CANN コミュニティが昇騰 NPU 向けに提供する TileLang オペレータリポジトリで、カーネル実装、自動化ケース、標準負荷、性能ベースラインの統一管理を強調し、10-26 にプッシュがあった。xwhzz/tilelang-hexagon は Qualcomm Hexagon バックエンドを手掛け、リポジトリは 32 MB で、10-26 にプッシュがあった。superAngGao/tilelang-debugger は H200 上でソースに基づく位置特定を行うデバッグワークスペースである。ほかに RuneFang/tvm_tilelang_cookbook(37 スター)と localizethedocs/tilelang-docs-l10n の定例更新がある。
出典:https://github.com/TongmingLAIC/tilelang-riscv
4.5 産業側:昇騰と DeepSeek のエコシステム格差縮小に関する報道(10-26)
10-26 にメディアが SemiAnalysis の 10-22 のレポートを伝え、DeepSeek と昇騰の深い連携がチップハードウェアとソフトウェアエコシステムにおける NVIDIA との格差を縮めつつあり、DeepSeek は 2027 年または 2028 年には主に昇騰技術スタックに依存するようになると予測した。報道では特に、この判断が DeepSeek による昇騰エコシステム向けソフトウェアツールの公開から数週間後に出たこと、その中には NVIDIA プラットフォームと一対一で対応する6つのソフトウェアモジュールが含まれることが言及されている。この消息自体は TileLang の機能変更には関わらないため、本期はエコシステムの背景としてのみ登録する。この6つのモジュールのうち高級言語コンパイルツールがすなわち tilelang-ascend であり、昇騰リポジトリ当日のインストーラと int4 行列積の修正はまさに同じ経路上にある。
5. トレンド観察
-
正確性欠陥の重心が「性能後退」から「コンパイル期に報告されず、特定のアーキテクチャやストレージドメインでのみ誤る」へ移った。当日の主線となる3件はいずれもこの類に属する:出力が誤ったストリームに割り当てられる(#3649、非デフォルトストリームかつ現在のストリームにキューされた作業がある場合にのみ破損)、TMA キャッシュヒントにテンプレート実引数が1つ足りない(#3683、eviction_policy 付きのバルクロードでのみ失敗)、sm8x のスパースメタデータにフォールバックがない(#3335、torch の正規の綴りでのみ認識され、さらに fp16 のレイアウト規約を誤って適用する)。この類の欠陥に共通するのはハードウェアまたはパラメータの前提があり、通常の CI のカバレッジ外にあるため、ますますファジングテスト系列によって駆動されるようになっている——本章で当日クローズされた6件のうち5件がこの系列に由来する。
-
同期とメモリ計画が主戦場に。当日浮上した #3680 の ThreadSync 欠陥3件、#3682 と #3681 の2件の修正、tvm フォークリポジトリ #88 の証明器修正、メインリポジトリ #3670 のリージョン化 kill 再配列、TileFoundry #272 の共有メモリピークが TIR に落ちない問題は、いずれも同じことを指している。コンパイラの同期とメモリ判断が「ヒューリスティック」から「証明可能」へと移行しつつあるということだ。これが #3666 の重要性を説明してもいる。組み込み TVM を独立パッケージ化された TIRx に置き換えたことで、以降この種の証明器修正がフォークリポジトリに留まらず、通常のリリースサイクルでメインラインに入れるようになった。
-
バックエンド適応は「精密制御」段階に。昇騰の
T.annotate_buffer_addressesにより UB、L1、L0A、L0B、L0C、SS の各プールをバイト単位でアドレス指定でき、重なり・アライメント・容量の検証も行える。海光のレジスタパイプライン、warp 発散、scale バッファのデュアルバンク分割も、「動く」を「ハードウェア構造に沿って配置できる」へと押し上げるものだ。これらの機能は直近2週間に集中しており、国産バックエンドが可用性の閾値を超え、性能上限の問題に取り組み始めたことを示している。 -
リリースエンジニアリングは依然として弱点。TileOPs の初回リリースは、ドキュメント、マニフェスト、パッケージング、API 安定性、入力検証がすべて整った後も、TestPyPI の試走と最初の tag で数日間止まっており、PyPI の成果物は 2025-10-09 のまま。196 スター、当日14コミットを合併したオペレータライブラリにとって、インストール可能性はさらなる1パーセントの最適化よりも外部採用速度を左右する。追跡チケットで唯一未完了の技術項目は MHA 行 0.996 であり、管理層はこれに tag を阻害させ続けるべきではない。
-
外部エコシステムは同期的に拡大中。RISC-V コミュニティハブリポジトリ、CANN オペレータリポジトリ、Hexagon バックエンドリポジトリが同日に出現または更新されたことは、TileLang が「tile-ai 組織のプロジェクト」から、複数プラットフォームがそれぞれリポジトリを建てるエコシステムへと変わりつつあることを示している。国内の読者にとって、RISC-V 方向は特に追跡に値する。これは従来の AI 眼鏡やエンドサイド推論のハードウェアルートと直接関係するが、現時点では空の骨格のみであり、実質的な進展は今後、実行可能なオペレータとバックエンドが着地するかどうかを見る必要がある。
付、素材と検証説明
情報源検証表
| 情報源 | 検証結果 |
|---|---|
| GitHub tile-ai 組織(29 リポジトリ) | 期間内に7リポジトリへプッシュ:tilelang、TileOPs、TileOPs-nightly、tilelang-hygon、tilelang-metax、tilelang-lsp、tilelang.github.io |
| 合併真値(is:merged) | 30件、1件ずつ合併時間が期間内であることを確認 |
| 審査中真値(is:open 更新) | 49件、更新時間順に全件確認 |
| 新規 issue | 11件、本文を1件ずつ読了 |
| クローズ issue | 12件、うち6件はファジングテストシリーズ |
| PyPI(tileops / tilelang) | tileops 最新 0.0.1.dev2(2025-10-09);tilelang 最新 0.1.15(2026-09-30) |
| リリースページ tag 検証 | TileOPs は tag なし release なし;tilelang 最新 v0.1.15(2026-09-30) |
| tilelang-ascend リポジトリ | 期間内プッシュなし(前回 10-23)、ただし PR と issue の活動は活発 |
| tilelang-metax リポジトリ | 組織ビューにプッシュタイムスタンプがあるが、6ブランチいずれも期間内コミットなし、実質変更なしとして処理 |
| 採用側(TileKernels / FlashQLA) | いずれも期間内プッシュなし(10-14) |
| Google News 中英語複数クエリ | 24時間以内に TileLang コンポーネントの正確なヒットなし;産業側報道1件は 4.5 を参照 |
| Hacker News | 期間内にフロントエンド項目1件、TileLang と無関係、除外 |
| arXiv | キーワード検索で期間内の新論文なし、直近1件は 2026-07-24 |
| 会員および採用側リポジトリコミュニティ検索 | 6つの外部リポジトリに期間内プッシュあり、4.4 を参照 |
完全情報源リスト
- [1] tilelang 呼び出し元ストリーム単位の出力割り当て(#3649)— https://github.com/tile-ai/tilelang/pull/3649
- [2] tilelang NVRTC の TVM-FFI 接続(#3285)— https://github.com/tile-ai/tilelang/pull/3285
- [3] tilelang バッチ TMA キャッシュヒント修正(#3683)— https://github.com/tile-ai/tilelang/pull/3683
- [4] tilelang リダクションの非有限定数とサブバイトアキュムレータ(#3189)— https://github.com/tile-ai/tilelang/pull/3189
- [5] tilelang SM120 タイルドスケーリング GEMM(#3286)— https://github.com/tile-ai/tilelang/pull/3286
- [6] tilelang JIT 出力形状スカラーパラメータ(#3602)— https://github.com/tile-ai/tilelang/pull/3602
- [7] tilelang sm8x スパースメタデータ dtype(#3335)— https://github.com/tile-ai/tilelang/pull/3335
- [8] tilelang DeepSelect Top-K サンプル(#3304)— https://github.com/tile-ai/tilelang/pull/3304
- [9] tilelang TIRx リファクタリング(#3666)— https://github.com/tile-ai/tilelang/pull/3666
- [10] tilelang ThreadSync フェンス修正(#3682)— https://github.com/tile-ai/tilelang/pull/3682
- [11] tilelang ThreadSync 不具合報告(#3680)— https://github.com/tile-ai/tilelang/issues/3680
- [12] tilelang float32 ブロードキャストのアドレス取得(#3681)— https://github.com/tile-ai/tilelang/pull/3681
- [13] tilelang リダクション batch 適応(#3684)— https://github.com/tile-ai/tilelang/pull/3684
- [14] tilelang 転置ターゲット形状検証(#3677)— https://github.com/tile-ai/tilelang/pull/3677
- [15] tilelang 領域化 kill 再配置(#3670)— https://github.com/tile-ai/tilelang/pull/3670
- [16] tilelang アトミックリダクション整数バッファ不具合(#2959)— https://github.com/tile-ai/tilelang/issues/2959
- [17] tilelang アトミックリダクションサンプル修正(#3340)— https://github.com/tile-ai/tilelang/pull/3340
- [18] tilelang 昇騰 精度バッファアドレッシング(#3607)— https://github.com/tile-ai/tilelang/pull/3607
- [19] tilelang 昇騰 手動スケジューリング FlashMLA(#3678)— https://github.com/tile-ai/tilelang/pull/3678
- [20] tilelang 昇騰 オンチップメモリオーバーフロー拒否(#3655)— https://github.com/tile-ai/tilelang/pull/3655
- [21] tilelang 昇騰 E4M3 の UB から L1 へのコピー(#3639)— https://github.com/tile-ai/tilelang/pull/3639
- [22] tilelang 昇騰 ビューレイアウトとポインタ単位(#3623)— https://github.com/tile-ai/tilelang/pull/3623
- [23] tilelang 昇騰 オプションのバッチスケジューリング(#3675)— https://github.com/tile-ai/tilelang/pull/3675
- [24] tilelang 昇騰 コア間同期バッチモード(#3614)— https://github.com/tile-ai/tilelang/pull/3614
-
[25] tilelang 昇騰 L0C デュアルコピー直接変換(#3468)— <https://github.com/tile-ai
- [26] tilelang-ascend インストーラ SHMEM 修正(#1946)— https://github.com/tile-ai/tilelang-ascend/pull/1946
- [27] tilelang-ascend インストーラ 失敗時即停止(#1943)— https://github.com/tile-ai/tilelang-ascend/pull/1943
- [28] tilelang-ascend W4A4 int4 行列積(#1731)— https://github.com/tile-ai/tilelang-ascend/pull/1731
- [29] tilelang-ascend 多次元コピーの先頭次元折り畳み(#1805)— https://github.com/tile-ai/tilelang-ascend/pull/1805
- [30] tilelang-ascend ラッパーの呼び出しごとのオーバーヘッド(#1945)— https://github.com/tile-ai/tilelang-ascend/issues/1945
- [31] tilelang-hygon レジスタパイプライン(#18)— https://github.com/tile-ai/tilelang-hygon/pull/18
- [32] tilelang-hygon ネイティブ B16 と動的 WASP(#28)— https://github.com/tile-ai/tilelang-hygon/pull/28
- [33] tilelang-hygon scale バッファのデュアルバンク(#29)— https://github.com/tile-ai/tilelang-hygon/pull/29
- [34] tilelang-metax リポジトリ — https://github.com/tile-ai/tilelang-metax
- [35] TileOPs 初回リリース追跡チケット(#2738)— https://github.com/tile-ai/TileOPs/issues/2738
- [36] TileOPs D128 稠密キー分割(#2777)— https://github.com/tile-ai/TileOPs/pull/2777
- [37] TileOPs D128 永続化 LSE 書き込み(#2774)— https://github.com/tile-ai/TileOPs/pull/2774
- [38] TileOPs 融合加算 LayerNorm 行分割(#2768)— https://github.com/tile-ai/TileOPs/pull/2768
- [39] TileOPs LayerNorm 精度欠陥(#2773)— https://github.com/tile-ai/TileOPs/issues/2773
- [40] TileOPs clamp 自体の精度(#2767)— https://github.com/tile-ai/TileOPs/pull/2767
- [41] TileOPs 累積スキャン単一起動(#2766)— https://github.com/tile-ai/TileOPs/pull/2766
- [42] TileOPs GELU 乗算(#2769)— https://github.com/tile-ai/TileOPs/pull/2769
- [43] TileOPs 2バイト二項演算(#2770)— https://github.com/tile-ai/TileOPs/pull/2770
- [44] TileOPs オペレータ層アライメント(#2772)— https://github.com/tile-ai/TileOPs/pull/2772
- [45] TileOPs アライメント判定をオブザーバへ(#2778)— https://github.com/tile-ai/TileOPs/pull/2778
- [46] TileOPs 出力オーバーラップ宣言(#2780)— https://github.com/tile-ai/TileOPs/issues/2780
- [47] TileOPs グループ GEMM オフセットストア(#2761)— https://github.com/tile-ai/TileOPs/pull/2761
- [48] TileOPs SM90 MoE オフセットストア(#2765)— https://github.com/tile-ai/TileOPs/pull/2765
- [49] TileOPs 出力 ABI(#2762)— https://github.com/tile-ai/TileOPs/pull/2762
-
[50] TileOPs FP8 ページドプリフィル e4m3(
- [51] TileOPs 疎長GDNプリフィル(#2779)— https://github.com/tile-ai/TileOPs/pull/2779
- [52] TileFoundry 動的共有メモリのサイズ確定(#271)— https://github.com/tile-ai/TileFoundry/pull/271
- [53] TileFoundry ウィンドウ全体をオペランド単位で発行(#273)— https://github.com/tile-ai/TileFoundry/pull/273
- [54] TileFoundry 共有メモリ解析のギャップ(#272)— https://github.com/tile-ai/TileFoundry/issues/272
- [55] tilelang-lsp CUDAデバッグのコミット — https://github.com/tile-ai/tilelang-lsp/commit/78a5761404
- [56] tilelang ドキュメントサイト — https://github.com/tile-ai/tilelang.github.io
- [57] tilelang PR説明規範(#3671)— https://github.com/tile-ai/tilelang/pull/3671
- [58] tilelang ドキュメントのみのPRでバックエンドテストをスキップ(#3672)— https://github.com/tile-ai/tilelang/pull/3672
- [59] tilelang Agentタイトル確認(#3673)— https://github.com/tile-ai/tilelang/pull/3673
- [60] tilelang レビューボットのノイズ低減(#3674)— https://github.com/tile-ai/tilelang/pull/3674
- [61] tile-ai/tvm Z3ブールビット演算(#88)— https://github.com/tile-ai/tvm/pull/88
- [62] tilelang-riscv コミュニティハブリポジトリ — https://github.com/TongmingLAIC/tilelang-riscv
- [63] CANNコミュニティ ops-tilelangリポジトリ — https://github.com/Atlas-Ascend-Fork/ops-tilelang
- [64] tilelang-hexagon バックエンドリポジトリ — https://github.com/xwhzz/tilelang-hexagon
- [65] PyPI tileops — https://pypi.org/project/tileops/
- [66] 昇騰とDeepSeekのエコシステム報道 — https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report