NVIDIA CUDA Tile を用いた C++ による高性能 GPU カーネルの開発

2026 年 5 月 26 日

AI が生成した要約

高評価

低評価

• NVIDIA CUDA Tile C++は、既存のC++コードベース内でタイルベースのGPUカーネルプログラミングを可能にし、明示的なスレッド管理なしに、NVIDIAアーキテクチャ全体にわたる低レベルのGPU並列処理、メモリアクセス、およびハードウェア機能を抽象化します。
• CUDA Tile C++プログラミングモデルは、多次元テンソルスパンとパーティションビューを使用し、固定サイズの配列タイル上で操作を行うことで、従来のSIMTカーネルよりも宣言的かつ効率的にベクトル加算や行列乗算などの並列計算を記述できるようにします。
• __restrict__ポインタ修飾子、16バイトアラインメントの仮定、マスク付きロード/ストア操作といった最適化により、パフォーマンスとメモリエフィシエンシーが向上し、タイルカーネルはブロックあたりスレッド1つで起動され、スレッドの実行詳細はコンパイラが処理します。
• CUDA Tile C++カーネルはNVIDIA Nsight Computeによるプロファイリングをサポートしており、従来のCUDA C++カーネルと同様に、タイル固有の詳細統計やソースレベルのパフォーマンスメトリクスを提供します。
• 行列乗算カーネルはタイルパーティションとNVIDIAの行列乗算累積（mma）演算を活用し、部分結果の効率的な蓄積を実現することで、複雑な線形代数ワークロードに対するCUDA Tilesの能力を実証しています。
• CUDA Tile C++を完全に活用するには、計算能力8.x以降のGPU、NVIDIA Driver R580以降、およびCUDA Toolkit 13.3以降が必要です。

AI 生成コンテンツは情報を不完全に要約する可能性があります。重要な情報は必ず確認してください。詳細はこちら

開発者は、既存の大規模な C++ GPU コードベース内で NVIDIA CUDA Tile プログラミングを利用できるようになり、タイルベースの抽象化を用いて高度に最適化された GPU カーネルを開発することが可能になりました。

NVIDIA CUDA 13.1 のリリースとともに導入された NVIDIA CUDA Tile は、GPU 向けのタイルベースプログラミングを提供します。これは上位レベルの言語レイヤーと、あらゆる高水準プログラミング言語がターゲットとする中間レイヤーを備えて設計されており、アプリケーションが直接それらにターゲットを設定する必要なく、NVIDIA ハードウェアの高度な機能（テンソルコア、共有メモリ、テンソルメモリアクセラレータ）を自動的に活用します。

タイルベース GPU アプリケーションで最初にサポートされた言語は Python でした。新たにリリースされたCUDA 13.3 では、C++ でタイルカーネルを記述するサポートが追加され、開発者が高度に最適化された GPU カーネルを構築できるようになりました。

CUDA Tile C++ とは？

CUDA Tile C++ は、CUDA Tile IR 仕様書 を基盤として構築された、C++ における CUDA Tile プログラミングモデルの表現です。これにより、開発者は C++ でタイルカーネルを記述し、単一命令多重スレッド (SIMT: Single Instruction, Multiple Threads) モデルに代わる、あるいはそれに加えて、タイルベースのモデルを用いて GPU カーネルを表現することが可能になります。

リフレッシュのため、タイルモデルにおける要点は以下の通りです。

• 多次元配列が主要なデータストレージとなります。
• タイルとは、カーネルが操作する配列の一部です。
• カーネルとは、ブロックによって並列実行される関数です。
• ブロックは GPU のサブセットであり、タイルに対する演算は各ブロック内のすべてのスレッドにわたって並列化されます。

CUDA Tile C++ は、ブロック内の並列処理、非同期処理、メモリアクセスの移動、および GPU プログラミングにおけるその他の低レベルの詳細を自動化します。CUDA Tile C++ は異なる NVIDIA GPU アーキテクチャ間で移植可能であり、開発者はコードを書き換えることなく最新のハードウェア機能を利用できます。

CUDA Tile C++ ベクトル加算の例

SIMT 向けの CUDA C++ に慣れている開発者であれば、おそらく標準的なベクトル加算カーネルに遭遇したことがあるでしょう。データがすでに GPU 上にあると仮定すると、CUDA SIMT におけるベクトル加算カーネルは 2 つのベクトルを受け取り、要素ごとに足し合わせて 3 つ目のベクトルを生成します。これは記述が最も簡単な CUDA カーネルの一つです。その様子は以下の通りです。

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)

{

int workIndex = threadIdx.x + blockIdx.x*blockDim.x;

if(workIndex {} と ct::shape{8_ic} はこの文脈において同等です。作成されるパーティションビューは、本質的に元の配列をタイルサイズである 8 のチャンクにスライスしたものです。

auto aView = ct::partition_view{aSpan, ct::shape{8_ic}};

auto bView = ct::partition_view{bSpan, ct::shape{8_ic}};

auto oView = ct::partition_view{oSpan, ct::shape{8_ic}};

• X 次元のブロックインデックスを ct::bid().x で取得して入力タイルを読み込みます。多次元ブロックを扱う場合は、Y および Z 次元も使用します。次に a と b のタイルを読み込みます。利便性のために auto を使用しますが、明示的に記述すれば、aTile と bTile は ct::tile> 型です。これらは要素型が float でサイズが 8 の 1D タイルです。パーティションビューを使用すると、ブロックインデックスを簡単に渡すことができます。load 関数は自動的に配列の正しいチャンクを取得し、それをタイルに読み込みます。

int bx = ct::bid().x;

auto aTile = aView.load(bx);

auto bTile = bView.load(bx);

• 結果を加算して格納します。これは、入力タイルに対して要素ごとの加算を実行し、その結果を出力タイルに格納する一行のコードです。この出力タイルを、X 次元における同じブロックインデックス bx を用いてインデックスされた oView パーティションビューへ格納します。

auto oTile = aTile + bTile;

oView.store(oTile, bx);

ベクトル加算の完全な例

以下の例では、C++ においてこのベクトル加算カーネルを呼び出す方法を示しています。これは完全に実行可能なコードの一部です。

コンパイラが最適化を行うのを助けるために、いくつか注意すべき点があります。

まず、最高のパフォーマンスを得るためには、カーネル実行中は入力および出力配列に対して、それぞれのポインタを通じてのみアクセスする必要があります。この条件が満たされる場合、他のポインタやシンボルを用いたアクセス（エイリアシング）は存在しません。配列ポインタに __restrict__ デコレータを付与することで、これをコンパイラに伝達できます。

基底ポインタが 16 バイト境界にアライメントされた配列を使用すると、コンパイラはより効率的なメモリアクセスパターンを生成できるようになります。カーネル引数のそれぞれに対して ct::assume_aligned を呼び出すことで、ポインタがアライメントされていることをコンパイラに伝達します。これらの呼び出しの戻り値を使用して、コンパイラがこのアライメントを活用できるようにします。cudaMalloc や同様の CUDA API によって返されるポインタは、常にこの条件を満たします（256 バイトのアライメントが保証されているため）。

最後に、タイルサイズは 8 よりもはるかに大きな値を使用してください。以下の実行可能なコードにはこれらの調整を適用し、タイルサイズで割り切れない可能性があるデータを処理する load_masked と store_masked の使用を追加してください。

以下に、カーネルとメイン関数を含む完全なコードを示します。適用された最適化と冗長性の削減にご注意ください。

#include

#include

#include "cuda_tile.h"

__tile_global__ void vectorAdd(float* __restrict__ a, float* __restrict__ b, float* __restrict__ out, size_t n) {

namespace ct = cuda::tiles;

using namespace ct::literals;

a = ct::assume_aligned(a, 16_ic);

b = ct::assume_aligned(b, 16_ic);

out = ct::assume_aligned(out, 16_ic);

int bx = ct::bid().x;

auto aTile = ct::partition_view{ct::tensor_span{a, ct::extents{n}}, ct::shape{1024_ic}}.load_masked(bx);

auto bTile = ct::partition_view{ct::tensor_span{b, ct::extents{n}}, ct::shape{1024_ic}}.load_masked(bx);

auto oTile = aTile + bTile;

auto oView = ct::partition_view{ct::tensor_span{out, ct::extents{n}}, ct::shape{1024_ic}};

oView.store_masked(oTile, bx);

}

#define checkCudaError(X) do {

auto ret = X;

if (ret != cudaSuccess) {

printf("\n error on line %d, CUDART error string : %s", __LINE__, cudaGetErrorString(ret));

exit(1);

}

} while (0)

int main() {

constexpr size_t N = 2ULL >>(d_a, d_b, d_out, N);

checkCudaError(cudaDeviceSynchronize());

checkCudaError(cudaMemcpy(h_out, d_out, sizeof(float) * N, cudaMemcpyDeviceToHost));

float max_err = 0.0f;

for (size_t idx = 0; idx >>(d_a, d_b, d_out, N);

タイルカーネルを起動する際、>>内の最初の引数はタイルブロック数（SIMT ではスレッドブロック数に相当）です。2 つ番目の引数は必ず 1 にする必要があります。カーネルの実行に使用されるスレッド数はコンパイラによって決定されますので、タイルカーネルを起動する際は常にこの引数に 1 を指定してください。

計算能力 8.0 で NVIDIA Ampere アーキテクチャ以降の GPU 上で CUDA 13.3 以降を実行している場合、これらのコマンドは以下の出力を生成します。

-arch sm_120 コマンドをアーキテクチャに合わせて調整し、cuda_tile.h を使用する場合は -std=c++20 を含め、タイルカーネルをコンパイルするには --enable-tile オプションを追加してください。

$ nvcc -std=c++20 --enable-tile -arch sm_120 -o vectorAdd vectorAdd.cu

$ ./vectorAdd

N: 67108864

Max error: 0.000000e+00

これで最初の CUDA Tile C++ プログラムが完了しました。

開発者ツール

タイル C++ カーネルは、SIMT カーネルと同様に NVIDIA Nsight Compute を用いてプロファイリングできます。以下のコマンドは、Nsight Compute を使用してプロファイルを作成する方法を示しています。

$ ncu -o VecAddProfile --set detailed ./vectorAdd

作成後、Nsight Compute のグラフィカルバージョンで開くと：

• ドロップダウンメニューから vectorAdd カーネルを選択します。
• Details タブを選択します。
• Tile Statistics リポートセクションを展開します。

図 1 は Nsight Compute から生成されたプロファイルを示しています。

image*図 1. Nsight Compute から生成されたプロファイルで、vectorAdd カーネルのタイル統計情報を示す*。

タイル統計情報レポートセクションには、指定されたタイルブロックの数、コンパイラによって選択されたブロックサイズ、およびその他のタイル固有情報が含まれていることに注意してください。

ソースページも、CUDA C++ カーネルと同様に、ソース行レベルでのタイルカーネルとパフォーマンスメトリクスをサポートしています。

行列乗算

以前の例では、パーティションビューの読み込みと保存の詳細を伴う vectorAdd を示しました。この行列乗算の例は、非常にシンプルなコードで行列乗算をどのように表現するかを示しています。

このカーネルは、MxK 行 KxN 列の行列乗算を実行して MxN 行列を計算します。本カーネルでは、M=8、N=16 であり、K は 8 の倍数であれば可変とします。ここでは K=24 と設定します。これらの非常に小さなサイズは、概念を示すためだけに使用されています。

完全なカーネルと主要ポイントの概要を以下に示します。

#include "cuda_tile.h"

__tile_global__ void kernel(float* __restrict__ a, float* __restrict__ b, size_t length, float* __restrict__ c) {

namespace ct = cuda::tiles;

using namespace ct::literals;

a = ct::assume_aligned(a, 16_ic);

b = ct::assume_aligned(b, 16_ic);

c = ct::assume_aligned(c, 16_ic);

auto aShape = ct::extents{8_ic, length};

auto bShape = ct::extents{length, 16_ic};

auto cShape = ct::extents{8_ic, 16_ic};

auto aSpan = ct::tensor_span{a, aShape};

auto bSpan = ct::tensor_span{b, bShape};

auto cSpan = ct::tensor_span{c, cShape};

auto aView = ct::partition_view{aSpan, ct::shape{4_ic, 8_ic}};

auto bView = ct::partition_view{bSpan, ct::shape{8_ic, 4_ic}};

auto cView = ct::partition_view{cSpan, ct::shape{4_ic, 4_ic}};

using f32x4x4 = ct::tile>;

auto accTile = ct::full(0);

auto [xBlock, yBlock, dummy] = ct::bid();

for (auto idx : ct::irange(0, 1 + int(length - 1) / 8)) {

auto aTile = aView.load_masked(xBlock, idx);

auto bTile = bView.load_masked(idx, yBlock);

accTile = ct::mma(aTile, bTile, accTile);

}

cView.store_masked(accTile, xBlock, yBlock);

}

• a、b、c 行列に対して ct::extents オブジェクトで範囲 (extents) を作成します。コンパイル時または実行時の値のいずれかを使用できます。M=8、N=16 ですが、K は可変です。これらは次のステップでテンソンスパンを作成するために使用されます。

auto aShape = ct::extents{8_ic, length};

auto bShape = ct::extents{length, 16_ic};

auto cShape = ct::extents{8_ic, 16_ic};

• テンソンスパンを作成します。これにより、a、b、c の情報が保持され、パーティションビューの作成に使用されます。

auto aSpan = ct::tensor_span{a, aShape};

auto bSpan = ct::tensor_span{b, bShape};

auto cSpan = ct::tensor_span{c, cShape};

• a を 4×8、b のビューを 8×4 のビューとするように、a、b、c の分割ビューを作成します。これらの調整は、a と b の値に適切に割り切れる場合に限り行えます。これらの次元はまた、c のビューが 4×4 であることを決定します。

auto aView = ct::partition_view{aSpan, ct::shape{4_ic, 8_ic}};

auto bView = ct::partition_view{bSpan, ct::shape{8_ic, 4_ic}};

auto cView = ct::partition_view{cSpan, ct::shape{4_ic, 4_ic}};

2 次元の分割は 2 つの次元でインデックス付けされます。a 行列は 8×24 であり、図 2 に示すように分割ビューは 4×8 です。

aView と bView の分割ビューサイズはまた、行列乗算中に結果を蓄積するために使用されるタイル accTile の形状も決定します。この例では、accTile は cView の形状と一致する 4×4 タイルです。

using f32x4x4 = ct::tile>;

auto accTile = ct::full(0);

• ct::bid を使用してループを実行し、3 つの次元におけるブロックインデックスを取得します。このループは 0 から長さ / 8 まで反復し、全体としての K 次元を 8 で割った値に対応します。8 での除算は、aView の K 次元と bView の K 次元が 8 であることに一致しています。ループ内では、load_masked を使用して a と b からタイルを読み込み、ct::mma を呼び出して行列乗算を実行し、その結果を accTile に蓄積します。

auto [xBlock, yBlock, dummy] = ct::bid();

for (auto idx : ct::irange(0, int(length / 8))) {

auto aTile = aView.load_masked(xBlock, idx);

auto bTile = bView.load_masked(idx, yBlock);

accTile = ct::mma(aTile, bTile, accTile);

}

• accTile の値を c のパーティションビューである cView に格納します。これで完了です。カーネルコードの大部分はデータのビュー設定と、データ読み込み・書き込みに費やされています。計算部分はシンプルです。

cView.store_masked(accTile, xBlock, yBlock);

• カーネルを起動します。cView の次元に合わせて dim3(2,4) を使用します。cView は 4×4 であり、各ブロックが C 行列の 4×4 チャンクを計算していることを意味します。C は 8×16 なので、cView の次元を C 行列の次元で割ります。8/4=2 かつ 16/4=4 となるため、dim3(2,4) でカーネルを起動します。

kernel>>(d_a, d_b, K, d_c);

CUDA Tile C++ を今日から始めよう

CUDA Tile C++ プログラムを実行するには以下の要件が必要です:

• 計算能力が 8.x 以降の GPU。

• NVIDIA ドライバ R580 以降。タイルカーネルに対して JIT コンパイルが必要な場合、NVIDIA ドライバのバージョンはコード生成に使用された CUDA Toolkit に関連付けられたバージョンと同等かそれより新しいものでなければなりません。例えば、CUDA Toolkit 13.3 には R610 ドライバ以降が必要です。

• CUDA Toolkit 13.3

タイルベースプログラミングのパワーが C++ 開発者にも利用可能になりました。