GPUカーネルをPythonで書くOSS「TileLang」の仕組み

Triton でカーネルを書き始めると、最初の数週間はうまくいきます。PyTorch の演算を数十行に畳み込めて、素朴な実装よりもずっと速い。ところが、ある地点から先に進まなくなります。shared memory への載せ方もレジスタへの割り付けもコンパイラが決めているため、プロファイラが示すボトルネックに対して、コード側から打てる手が見つからないのです。

かといって CUTLASS や CuTe への書き直しは、多くのチームにとって現実的ではありません。テンプレートの記述量が跳ね上がり、カーネル 1 本のメンテナンスコストが別物になります。「Triton では足りないが、CUDA の手書きまでは踏み込めない」という中間地帯で止まっているチームは珍しくありません。

この中間地帯を正面から狙っているのが、本記事で取り上げる TileLang（リポジトリ名: tile-ai/tilelang）です。設計上のポジションを一言でまとめるなら、メモリ階層のバッファは開発者が明示的に宣言し、スレッドごとのレイアウトはコンパイラが推論するという分担になります。制御したい部分だけを手元に残し、面倒な部分は手放せるように境界を引き直した DSL だと捉えると、既存の選択肢との違いが見えやすくなります。

なお本記事は、README・公式ドキュメント・GitHub API から取得した公開情報をもとに整理した内容です。実行環境でのインストールやベンチマークの再実行は行っておらず、性能に関する数値はすべて開発元や論文による測定結果として紹介します。

本記事では、TileLang が埋めようとしている課題、Python ライクなカーネル記述の実際のコード、メモリ階層制御とレイアウト推論の仕組み、対応バックエンドの広がりとサポートレベルの差、Triton・CuTe DSL・ThunderKittens との違い、そして採用判断に必要なプロジェクトの健全性とライセンス上の注意点を順に解説します。

TileLangとは｜GPUカーネル開発のどの課題を埋めるOSSか

TileLang（公式ドキュメント上の正式名称は Tile Language、別称 tile-lang）は、GEMM・Dequant GEMM・FlashAttention・LinearAttention といった高性能な GPU/CPU/NPU カーネルの開発を簡潔にするためのドメイン固有言語です。Apache TVM の上にコンパイラ基盤を持ち、そこに Python 的な構文を載せることで、「生産性を犠牲にせず、最先端の性能に必要な低レベル最適化も手放さない」ことを目標に掲げています（出典: GitHub リポジトリの README）。

TileLangがGPUカーネル開発で狙う位置づけ

設計思想としてより踏み込んだ説明は、公式ドキュメントの概要ページにあります。そこでは、中身を覗けないコンパイラ最適化に頼るのではなく、メモリ配置とデータ移動を開発者が直接制御できるようにするという方針が示されています。物理メモリ空間やアクセラレータ固有の構造に対応する組み込み関数を、ユーザーが触れる形で露出させている点が特徴です。

この位置づけは、第三者による解説でも同様に整理されています。Triton は書きやすい代わりにレイアウトの判断をコンパイラに委ねる部分が大きく、CUTLASS/CuTe は完全な制御と引き換えにテンプレートの記述量が膨らむ。TileLang はその中間を取るもの、という見立てです（出典: atlascloud の TileLang 解説記事）。

TileLangリポジトリの基本情報

採用判断の出発点として、リポジトリの客観的な数字を押さえておきます。以下はいずれも 2026 年 10 月時点で GitHub API から取得した値です。

アーカイブ済みでもフォークでもなく、最終更新が記事執筆時点の直近であることは、少なくとも「放置されたリポジトリを掴まされる」リスクが低いことを示します。ただしスター数 8,606 という規模は、後述する Triton（20,372）と比べると 2 分の 1 以下です。エコシステムの厚みという観点では差があることを前提に読み進めてください。

TileLangのコード例で見るPythonライクなカーネル記述

抽象論よりも、実際の書き味を見たほうが移植コストの見積もりは早くなります。インストールは PyPI から行えます。

pip install tilelang

（出典: tile-ai/tilelang README）

GEMM + ReLU のコード例

README の Quick Start には、FP16 の GEMM を FP32 で累算し、ReLU エピローグを融合させたカーネルの定義・コンパイル・実行・検証までを 1 本にまとめた例が掲載されています。以下は改変を加えずにそのまま引用したものです。

import torch
import tilelang
import tilelang.language as T


@tilelang.jit
def matmul_relu(A, B, block_M: int = 128, block_N: int = 128, block_K: int = 32):
    M, N, K = T.const("M, N, K")
    A: T.Tensor((M, K), T.float16)
    B: T.Tensor((K, N), T.float16)
    C = T.empty((M, N), T.float16)

    with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
        A_shared = T.alloc_shared((block_M, block_K), T.float16)
        B_shared = T.alloc_shared((block_K, block_N), T.float16)
        C_local = T.alloc_fragment((block_M, block_N), T.float32)

        T.clear(C_local)
        for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):
            T.copy(A[by * block_M, k * block_K], A_shared)
            T.copy(B[k * block_K, bx * block_N], B_shared)
            T.gemm(A_shared, B_shared, C_local)

        for i, j in T.Parallel(block_M, block_N):
            C_local[i, j] = T.max(C_local[i, j], 0)

        T.copy(C_local, C[by * block_M, bx * block_N])

    return C


M = N = K = 1024
a = torch.randn((M, K), device="cuda", dtype=torch.float16)
b = torch.randn((K, N), device="cuda", dtype=torch.float16)
c = matmul_relu(a, b)
torch.testing.assert_close(c, torch.relu(a @ b), rtol=1e-2, atol=1e-2)

（出典: tile-ai/tilelang README「Quick Start」）

コード例に登場するプリミティブの役割

1 行ずつ追う代わりに、役割の単位で整理します。

• @tilelang.jit: カーネルを初回使用時に、入力の形状とコンパイル時引数で特殊化します。ターゲットは実行環境から自動的に選ばれます
• T.Kernel: 実行コンテキストを定義します。ブロックインデックス（bx / by）とスレッド数をここで宣言し、コンパイラがメモリアクセスと計算を推論・最適化する足場になります
• T.alloc_shared / T.alloc_fragment: 前者はオンチップの高速メモリ（NVIDIA GPU の shared memory に対応）、後者はアキュムレータを置くフラグメントメモリ（同じくレジスタファイルに対応）を確保します
• T.Pipelined(..., num_stages=3): global から shared へのデータ転送をステージングし、転送と計算をオーバーラップさせます
• T.gemm: タイル単位の行列積を、ターゲットバックエンドの実装へマップします
• T.Parallel: 要素ごとの処理（この例では ReLU エピローグ）をスレッド間で並列に実行します
• T.copy: global メモリとハードウェア固有メモリのあいだのデータ移動を担います

注目したいのは、shared memory とレジスタに相当するバッファが、ソースコード上に明示的な宣言として現れている点です。どのデータがどの階層に載るかをコードから読み取れるため、プロファイラの結果とソースを突き合わせやすくなります。一方で、スレッドごとのレイアウトやインデックス計算はコードに現れません。ここがコンパイラの担当範囲です。

より体系的な構文の説明は、プログラミングガイドの language basics にまとまっています。実際のカーネル群は examples ディレクトリ にあり、GEMM・grouped GEMM・FP8 GEMM・FlashAttention・Flash Decoding・block-sparse attention・DeepSeek MLA など、用途別の実装を読み比べられます。

メモリ階層の明示制御とレイアウト推論の仕組み

TileLang が「中間地帯」を成立させている機構を、公式ドキュメントの記述に沿って分解します。

3つの抽象レベル（Beginner / Developer / Expert）

公式ドキュメントは、TileLang のプログラミングインターフェースを 3 つの層に分けて説明しています。重要なのは、これらが 1 つのカーネル内で混在できることです。

（出典: The Tile Language: A Brief Introduction）

採用を検討する立場から見ると、現時点で実質的な入り口は Developer 層という理解が妥当です。「初心者でもハードウェアを意識せず書ける」という触れ込みを額面どおり受け取ると、期待と実態がずれます。逆に、GPU のメモリ階層をすでに理解しているチームであれば、Developer 層で書き始めて詰まった箇所だけを Expert 層に降ろす、という段階的な最適化が成立します。

レイアウト推論とメモリ階層プリミティブ

T.alloc_fragment で確保するアキュムレータのサイズは、スレッドブロック全体のレジスタ使用量を指します。そこから各スレッドへの割り当てを決めるのが layout inference パスで、T.Fragment レイアウトを導出してスレッドごとのレジスタ配置を確定させます。T.Parallel による並列代入の挙動も、このレイアウト推論の結果に依存します。

この分担が、先ほど触れた設計思想の実装にあたります。どの階層にどのバッファを置くかは人が決め、そのバッファをスレッドにどう割るかは機械が決める。Triton が後者だけでなく前者も含めて判断する範囲が広いのに対し、TileLang は前者を front-end の宣言として露出させた、という整理ができます（この差分は論文 TileLang: A Composable Tiled Programming Model for AI Syste… でも TileLang 側の主張として述べられています）。

コンパイルフローの6段階

公式ドキュメントは、記述したプログラムが実行形式になるまでを 6 段階で説明しています。

1. Tile Program: 高レベルの計算仕様
2. Tile Program with Tile Library: ライブラリ呼び出しに展開された状態
3. Tile Program with Thread Primitives: 手書きの低レベルスレッド構造を含む状態
4. IRModule: ハードウェア詳細を含む中間表現へのロワリング
5. Source Code Generation: ターゲット別のソース（CUDA / HIP / LLVM 等）の生成
6. Hardware-Specific Executable: ターゲットデバイス向け実行形式へのコンパイル

段階 2 と 3 が分かれていることが、前述した 3 層の混在を支えています。ライブラリで済むところは任せ、必要な箇所だけスレッドプリミティブで書き下ろす、という操作がコンパイルフロー上に位置づけられているわけです。

なお TileLang は、各パスを覗くためのツールも用意しています。pass visualizer（変換の構造ツリー閲覧）、IR lower trace（全パスにわたる IR 変化の確認）、pass diff（パス間の IR 比較）などで、生成コードが期待どおりかを追跡できます。「コンパイラの判断を覗けるようにする」という思想がツール側にも一貫している点は、評価の材料になります。

対応バックエンドとハードウェアの広がり

TileLang は TileLang-X と呼ばれるマルチバックエンドコンパイラへの移行を進めており、モジュール化されたバックエンド抽象を中心に設計されています。コンパイル対象は Target オブジェクトで表現し、既定の auto ターゲットは CUDA / HIP / Metal / Ascend デバイスを自動検出します。

サポートレベル別のバックエンド一覧

（出典: tile-ai/tilelang README「Platform and Backend Support」。ターゲット指定の構文やアーキテクチャオプションは target guide に詳細があります）

サポートレベルの読み方

表を眺めると対応ハードウェアの広さが目を引きますが、採用判断ではレベルごとの意味の違いのほうが重要です。

• Primary（CUDA）: リリース wheel と CI の両方でカバーされる中核バックエンド。ただし TMA・WGMMA・TMEM といった機能は、対応する GPU アーキテクチャが前提になります
• Supported: 本体リポジトリで実装されています。ただし保証の粒度には差があります。たとえば AMD は Linux wheel に同梱される一方、CI は self-hosted の gfx942（MI300X）ランナーで動いており、gfx950 はまだカバーされていません。Ascend 950 は USE_ASCEND=ON でのソースビルドが必要で、CANN と torch_npu を前提とします。Metal 4 の cooperative tensors は対応する M5 システムに限られます
• Experimental: LLVM CPU はソースビルド（USE_LLVM=ON、LLVM 15 以降）が必要で、WebGPU はコード生成とランタイム統合が発展途上と README に明記されています
• Ecosystem: 別リポジトリで開発され、TileLang のリリース wheel には含まれず、互換性スケジュールも独立しています。国産アクセラレータ向けのアダプタはここに属します

「対応表に自社のハードウェアが載っている」ことと「pip install tilelang で明日から使える」ことは別だ、という点が実務上の分かれ目です。prebuilt wheel が用意されているのは Linux x86-64/AArch64、Windows x86-64、macOS arm64 で、それ以外はソースビルドや別リポジトリの導入が前提になります。自社の対象がどのレベルに該当するかを Installation guide とあわせて確認しておくと、見積もりのずれを防げます。

類似OSS（Triton・CuTe DSL・ThunderKittens）との違い

TileLang 単体で評価するより、同じ領域の OSS と並べたほうが輪郭がはっきりします。以下の数値はいずれも 2026 年 10 月時点の GitHub 実測値です。

主要3プロジェクトとの比較表

Tritonのままでよいケース、TileLangを検討すべきケース

比較表から読み取れる実務上の切り分けは次のとおりです。

Triton を維持すべきケース: 単純な elementwise 処理や軽い fusion が中心で、現状の性能に不満がない場合です。この領域では Triton のほうが素直に書ける、という評価が第三者の解説でも示されています（出典: atlascloud の解説記事）。エコシステムの厚みも Triton が明確に上です。

TileLang を検討すべきケース: 複雑な fused kernel で、細粒度の制御が足りずに性能が頭打ちになっている場合です。TileLang の論文（ICLR 2026 採択）はまさにこの問題設定を出発点としており、H100 において Triton 比で最大 5 倍の高速化を報告しています。ただしこれは著者ら（開発元）による測定結果であり、対象カーネルやハードウェア構成に依存する値です（出典: ICLR 2026 ポスターページ）。自社ワークロードでの再現性は別途確認が必要という前提で受け取ってください。

CuTe を選ぶべきケース: NVIDIA GPU 専用で構わず、最後の数パーセントまで詰める必要があり、テンプレート記述のコストを許容できる場合です。

ThunderKittens が合うケース: C++ の既存 CUDA コードベースがあり、そこにタイル抽象だけを持ち込みたい場合です。独立したコンパイラ基盤を持たないヘッダライブラリであるため、TileLang のようなオートチューナやマルチバックエンドは期待できません。

競合ではなく階層が違うケース（CuTe DSL バックエンド）

CuTe DSL については、「どちらを選ぶか」という構図にならないケースがある点に注意が必要です。TileLang は cutedsl を Experimental バックエンドとして持っており、nvidia-cutlass-dsl を導入すれば CuTe DSL を経由したコンパイルが可能です。上位の記述を TileLang で行い、下位のコード生成に CuTe を使うという構成が成り立つため、純粋な競合というより階層が異なる関係にあります。

採用判断のためのプロジェクト健全性とライセンスの注意点

ここまでの機能面の整理に加えて、本番採用に耐えるかを判断する材料を揃えます。

メンテナンス状況の実測値

2026 年 10 月時点で GitHub API から取得した数値は次のとおりです。

更新頻度とコントリビューターの広がりは、個人メンテナンスの小規模プロジェクトとは明確に異なる水準です。Releases 一覧からは、月 1 本前後のペースでマイナーリリースが続いていることも確認できます。

一方で、バージョンが 0.1.x 系にとどまっている点は見落とせません。たとえば v0.1.13（2026-08-03）ではマルチバックエンド言語ダイアレクトの導入と同時に複数のレガシー API が削除されており、README にもアップグレード前に互換性ノートを読むよう明記されています。カーネルをプロダクションに載せる場合、バージョン固定の方針と、アップグレード時の互換性確認を運用に組み込む前提で見積もる必要があります。

ライセンス表記の読み解き（NOASSERTIONの理由）

GitHub API 上のライセンス判定は NOASSERTION（表示上は Other）になっています。これは「ライセンスが不明」という意味ではありません。LICENSE ファイルの本文は MIT License（Copyright (c) Tile-AI）ですが、冒頭に 2024-12-01 から 2025-03-14 までの期間について Microsoft Corporation との追加的な協業条項に関する追記があり、標準的な MIT 全文と差分が生じているため、GitHub の自動判定が Other に倒れている、という構造です。

この追記の背景は README の謝辞とも整合します。初期版は LeiWang1999・chengyupku・nox-410 の各氏が北京大学の Zhi Yang 教授の指導下で開発し、一部は Microsoft Research でのインターンシップ中に行われたと記載されています。

実務上の結論は明快です。MIT ベースではあるものの標準 MIT と完全一致ではないため、商用利用を前提とする場合は LICENSE 原文を法務確認の対象に含めるのが安全です。自動判定の表示だけで「MIT だから問題ない」と片付けないほうがよいケースにあたります。

学習リソースと開発者体験

初見で試す際のハードルを下げる資産が、周辺リポジトリとして公開されています。

• TileLang Puzzles: 段階的に難度が上がる演習 10 問。手を動かしながら言語仕様を把握する用途に向きます
• TileLang LSP: バッファ形状・dtype・スコープ・推論されたレイアウトを inlay hints で表示し、hover 詳細と診断を提供する Language Server。レイアウト推論の結果をエディタ上で確認できます
• デバッグ・解析ツール群: pass visualizer、IR lower trace、pass diff、AutoDD（自動デルタデバッグ）、IKET プロファイラ統合など
• Discord コミュニティ: README からリンクされています

ただし公式ドキュメント自体に「実験的・不完全な可能性がある」旨の注記があり、ドキュメントの成熟度はまだ発展途上という前提で読む必要があります。

向くケースと慎重に判断すべきケース

ここまでの材料を、判断の形に圧縮します。

TileLang が向くケース

• Triton でカーネルを書いた経験があり、メモリ階層の制御不足で性能が頭打ちになっている
• FlashAttention 系・MLA 系・量子化 GEMM など、examples に近い形のカーネルを扱っている
• CUDA 以外（ROCm・Metal・Ascend 950）にも展開する可能性があり、記述の共通化に価値を感じる
• カーネル開発に専任で時間を割けるメンバーがいて、0.1.x 系の API 変更に追随できる

慎重に判断すべきケース

• カーネル記述の経験が浅く、ハードウェアを意識せず書ける層（Beginner インターフェース）に期待している。この層は公式に未完成と明記されています
• 対象ハードウェアが Experimental または Ecosystem レベルに該当する（wheel 非同梱・互換性スケジュール独立のリスクを負う）
• バージョン固定とアップグレード検証の運用を回す余力がない
• ライセンスの法務確認プロセスに時間がかかる組織で、短期の導入判断を求められている

まとめ

TileLang は、メモリ階層のバッファを開発者が明示的に宣言し、スレッドごとのレイアウトをコンパイラが推論するという分担で、Triton と CUTLASS/CuTe のあいだにある空白を埋めようとしている OSS です。Developer 層と Expert 層を 1 カーネル内で混在できる構造と、コンパイラの判断を追跡できるツール群が、この分担を実務で成立させる支えになっています。

採用を検討する際の確認ポイントは 3 つに整理できます。第一に、自社のハードウェアが Primary / Supported / Experimental / Ecosystem のどのレベルに該当するか。第二に、0.1.x 系ゆえの API 変更に追随する運用を組めるか。第三に、NOASSERTION と判定されているライセンスの原文を法務確認できるか。

次の一歩としては、language basics で構文の全体像を掴んだうえで、examples ディレクトリから自社のワークロードに最も近いカーネルを 1 本選んで読み込むのが効率的です。書き味と制御粒度が要件に合うかは、そこで判断がつきます。

関連情報

GPU 推論基盤の最適化や機械学習まわりの開発体制づくりをご検討中の方は、お問い合わせフォームからご相談ください。要件が固まりきっていない構想段階からのご相談にも対応しています。