CUDA Tile IR（CUDA タイル IR）

CUDA Tile IR は、MLIR をベースにした中間表現（IR）とコンパイラインフラストラクチャで、CUDA カーネルの最適化を行うためのものです。
NVIDIA Tensor Core ユニットを対象としたタイルベース計算パターンと最適化に焦点を当て、NVIDIA GPU 上でタイル化された計算を表現・最適化するための包括的なエコシステムを提供します。

• オープンソースリリースは CUDA Toolkit 13.1 と同梱。
• 詳細情報：https://developer.nvidia.com/cuda/tile

コアコンポーネント

CUDA Tile 仕様

仕様書は、NVIDIA GPU 上でタイルベース計算を行うための正式な意味論、操作、および型システムを定義しています。詳細（方言オペレーション、型、変換パス）は CUDA Tile Specification を参照してください。

CUDA Tile のビルド

前提条件

• CMake ≥ 3.20.0
• C++17 コンパイラ
• Python ≥ 3.6（バインディング用）
• MLIR/LLVM ソースまたは互換コミットのプリビルトライブラリ（

  cmake/IncludeLLVM.cmake

  を確認）
• Ninja（任意）

クイックスタート

# 設定
cmake -G Ninja -S . -B build \
  -DCMAKE_BUILD_TYPE=Release \
  -DLLVM_ENABLE_ASSERTIONS=OFF \
  -DCUDA_TILE_ENABLE_BINDINGS_PYTHON=ON

# ビルド
cmake --build build

# テスト実行
cmake --build build --target check-cuda-tile

MLIR/LLVM ソースは GitHub から自動で取得されます。

ビルド設定オプション

MLIR/LLVM のビルド方法

cmake -G Ninja -S . -B build -DCMAKE_BUILD_TYPE=Release

cmake -G Ninja -S . -B build -DCMAKE_BUILD_TYPE=Release -DCUDA_TILE_USE_LLVM_SOURCE_DIR=/path/to/llvm/sources

cmake -G Ninja -S . -B build -DCMAKE_BUILD_TYPE=Release -DCUDA_TILE_USE_LLVM_INSTALL_DIR=/path/to/llvm/install

Python バインディング

cmake -G Ninja -S . -B build \
  -DCMAKE_BUILD_TYPE=Release \
  -DCUDA_TILE_ENABLE_BINDINGS_PYTHON=ON

MLIR/LLVM をソースからビルドする場合、バインディングは自動で有効化されます。
プリビルト LLVM ライブラリを使用する場合は、

-DMLIR_ENABLE_BINDINGS_PYTHON=ON

Ccache

cmake -G Ninja -S . -B build \
  -DCMAKE_BUILD_TYPE=Release \
  -DCUDA_TILE_ENABLE_CCACHE=ON

テスト

CUDA Tile は LLVM の

lit

-DCUDA_TILE_ENABLE_TESTING=ON

cmake --build build --target check-cuda-tile

プロジェクトへの統合

オプション 1：プリビルトライブラリを利用

include_directories(${CUDA_TILE_INSTALL_DIR}/include)

# 必要なライブラリにリンク
target_link_libraries(your_target PRIVATE
  CudaTileDialect          # 方言オペレーション・型
  CudaTileBytecodeReader   # バイトコード読み込み
  CudaTileBytecodeWriter   # バイトコード書き出し
)

オプション 2：ソース統合

include(FetchContent)

FetchContent_Declare(
  cuda_tile
  GIT_REPOSITORY https://github.com/NVIDIA/cuda-tile.git
  GIT_TAG        main
  SOURCE_DIR     ${CMAKE_BINARY_DIR}/_deps/cuda_tile-src
  BINARY_DIR     ${CMAKE_BINARY_DIR}/_deps/cuda_tile-build
)

# ビルドオプション（必要に応じて設定）
set(CUDA_TILE_USE_LLVM_INSTALL_DIR ${YOUR_LLVM_INSTALL_DIR} CACHE PATH "")
set(CUDA_TILE_ENABLE_BINDINGS_PYTHON ON CACHE BOOL "")
set(CUDA_TILE_ENABLE_TESTING OFF CACHE BOOL "")

FetchContent_MakeAvailable(cuda_tile)

include_directories(${CUDA_TILE_SOURCE_DIR}/include)
include_directories(${CUDA_TILE_BINARY_DIR}/include)

target_link_libraries(your_target PRIVATE
  CudaTileDialect
  CudaTileBytecodeReader
  CudaTileBytecodeWriter
)

例：CUDA Tile IR プログラムの作成と実行

1. MLIR ソース (

example.mlir

)

cuda_tile.module @example_module {
  entry @example_kernel(%data_pr : tile<ptr<f32>>) {
    print "Running example module\n"
    %offsets = iota : tile<128xi32>
    %data_ptr_reshaped = reshape %data_pr : tile<ptr<f32>> -> tile<1xptr<f32>>
    %data_ptr_broadcasted = broadcast %data_ptr_reshaped : tile<1xptr<f32>> -> tile<128xptr<f32>>
    %data_ptr_tensor = offset %data_ptr_broadcasted, %offsets : tile<128xptr<f32>>, tile<128xi32> -> tile<128xptr<f32>>
    %data, %token = load_ptr_tko weak %data_ptr_tensor : tile<128xptr<f32>> -> tile<128xf32>, token
    print "Data: %f\n", %data : tile<128xf32>
    return
  }
}

2. バイトコードへの変換

cuda-tile-translate example.mlir \
  --bytecode-version=13.1 \
  --mlir-to-cudatilebc \
  --no-implicit-module -o example.tilebc

3. AoT（Ahead‑of‑Time）コンパイル（任意）

tileiras --gpu-name sm_100 example.tilebc -o example.cubin

(

sm_100

JIT コンパイルを起動時に行う場合は AoT をスキップし、直接

example.tilebc

を使用します。

4. ホストプログラム (

example_host.cpp

)

#include <cuda.h>
#include <cuda_runtime_api.h>
#include <stdio.h>
#include <stdlib.h>

#define CUDA_CHECK(call)                                          \
  do {                                                            \
    CUresult err = call;                                         \
    if (err != CUDA_SUCCESS) {                                   \
      const char *errStr;                                        \
      cuGetErrorString(err, &errStr);                            \
      fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, errStr); \
      exit(1);                                                   \
    }                                                             \
  } while (0)

float data[] = { /* 635 values */ };

int main() {
  CUdevice cuDevice; CUcontext cuContext; CUmodule cuModule;
  CUfunction example_kernel; CUstream stream;

  CUDA_CHECK(cuInit(0));
  CUDA_CHECK(cuDeviceGet(&cuDevice, 0));
  CUDA_CHECK(cuCtxCreate(&cuContext, NULL, 0, cuDevice));
  CUDA_CHECK(cuStreamCreate(&stream, CU_STREAM_DEFAULT));

  CUDA_CHECK(cuModuleLoad(&cuModule, "example.cubin"));
  CUDA_CHECK(cuModuleGetFunction(&example_kernel, cuModule, "example_kernel"));

  CUdeviceptr data_ptr;
  CUDA_CHECK(cuMemAlloc(&data_ptr, sizeof(data)));
  CUDA_CHECK(cuMemcpyHtoD(data_ptr, data, sizeof(data)));

  void *kernel_args[] = {&data_ptr};
  CUDA_CHECK(cuLaunchKernel(example_kernel,
                            1, 1, 1,
                            1, 1, 1,
                            0, stream,
                            kernel_args, NULL));
  CUDA_CHECK(cuCtxSynchronize());

  CUDA_CHECK(cuModuleUnload(cuModule));
  CUDA_CHECK(cuCtxDestroy(cuContext));

  return 0;
}

5. ホストビルド

g++ example_host.cpp -o example \
  -I/usr/local/cuda/include \
  -L/usr/local/cuda/lib64 -lcuda

必要に応じてコンパイラ、インクルード、およびライブラリパスを調整してください。

6. 実行

./example

期待される出力

Running example module
Data: [0.000000, 5.000000, 10.000000, … , 630.000000, 635.000000]

コントリビューションとサポート

現在の状態: プロジェクトは活発に開発中で、明確なロードマップを持っています。
外部からのコントリビューションは 現時点では受け付けていません。

バグ報告・フィードバック・体験共有は GitHub Issues で行ってください。皆様の入力が今後の優先順位決定に役立ちます。

ライセンス

CUDA Tile IR は Apache License v2.0（LLVM Exceptions） の下でライセンスされています。