GPU での CUDA または PTX コードの実行

`CUDAKernel` ワークフローの概要

このページでは、実行可能カーネルを CUDA^® C++ ソースファイル (CU) から作成し、MATLAB^® でそのカーネルを GPU 上で実行する方法を説明します。MATLAB ではカーネルは CUDAKernel オブジェクトで表されます。このオブジェクトは、ホストメモリに保存されている配列または GPU 配列を処理できます。

CUDAKernel の一般的なワークフローの手順は以下のとおりです。

mexcuda を使用して、CU ファイルから並列スレッド実行 (PTX) ファイルをコンパイルします。mexcuda を使用した PTX ファイルのコンパイルには、CUDA ツールキットは必要ありません。
R2023a より前: 関数 mexcuda ではなく、NVIDIA^® CUDA ツールキット内の nvcc コンパイラを使用して PTX ファイルをコンパイルします。
関数 parallel.gpu.CUDAKernel を使用して、CU ファイルおよび PTX ファイルから CUDAKernel オブジェクトを作成します。CUDAKernel には GPU 実行可能コードが含まれています。
CUDAKernel のプロパティを設定し、GPU での実行を制御できるようにします。
必要な入力を使用して CUDAKernel で feval を呼び出し、GPU でカーネルを実行します。

この手順に従って生成された MATLAB コードの例を以下に示します。

% 1. Compile a PTX file.
mexcuda -ptx myfun.cu

% 2. Create CUDAKernel object.
k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu");

% 3. Set object properties.
k.GridSize = [8 1];
k.ThreadBlockSize = [16 1];

% 4. Call feval with defined inputs.
g1 = gpuArray(in1); % Input gpuArray.
g2 = gpuArray(in2); % Input gpuArray.

result = feval(k,g1,g2);

この後の各節では、これらのコマンドとワークフローの手順の詳細について説明します。

`CUDAKernel` オブジェクトの作成

GPU で実行する CU ファイルがある場合は、まずそれをコンパイルして PTX ファイルを作成しなければなりません。PTX ファイルをコンパイルするには、-ptx フラグを指定して CU ファイルを mexcuda に渡します。

mexcuda -ptx myfun.cu

これにより、PTX ファイル myfun.ptx が生成されます。

CU ファイルと PTX ファイルを使用して、CUDAKernel オブジェクトを作成します。

k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu");

メモ

CUDAKernel オブジェクトに対し save または load を実行することはできません。

`CUDAKernel` オブジェクトのプロパティ

終了のセミコロンを指定しないで CUDAKernel オブジェクトを作成した場合、またはコマンドラインでオブジェクト変数を入力した場合、カーネルオブジェクトのプロパティが MATLAB により表示されます。

k = parallel.gpu.CUDAKernel("conv.ptx","conv.cu")

k = 
  parallel.gpu.CUDAKernel handle
  Package: parallel.gpu

  Properties:
     ThreadBlockSize: [1 1 1]
  MaxThreadsPerBlock: 512
            GridSize: [1 1 1]
    SharedMemorySize: 0
          EntryPoint: '_Z8theEntryPf'
  MaxNumLHSArguments: 1
     NumRHSArguments: 2
       ArgumentTypes: {'in single vector'  'inout single vector'}

CUDAKernel オブジェクトのプロパティは、そのオブジェクトの実行動作の一部を制御します。変更可能なプロパティを変更するには、ドット表記を使用します。オブジェクトのプロパティについては、CUDAKernel を参照してください。設定可能なプロパティを変更する一般的な理由として、以下のようにスレッド数を指定することがあげられます。

エントリポイントの指定

1 つの PTX ファイルには、さまざまなカーネルへの複数のエントリポイントを含めることができます。こうした各エントリポイントの名前は一意でなければなりません。各エントリポイントの名前は (C++ でのマングルと同様に) マングルされますが、CU ファイルから取得した元の関数名を常に含みます。たとえば、CU ファイルによりカーネル関数が次のように定義されているとします。

__global__ void simplestKernelEver( float * x, float val )

この場合、PTX コードには _Z18simplestKernelEverPff のようなエントリが含まれます。

複数のエントリポイントがある場合は、カーネルを生成するために parallel.gpu.CUDAKernel を呼び出す際に、特定のカーネルのエントリ名を指定します。

k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu","myKernel1");

メモ

関数 parallel.gpu.CUDAKernel は PTX ファイル内でエントリ名を検索し、出現するすべての部分文字列を照合します。したがって、エントリポイントの名前は他のエントリポイントの名前の部分文字列にしないでください。

元のエントリ名を制御できない場合、各エントリについて派生する一意のマングルに注意してください。たとえば、以下の関数テンプレートについて考えます。

template <typename T>
__global__ void add4( T * v1, const T * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

テンプレートを浮動小数点および倍精度小数点用に拡張すると、エントリポイントが 2 つ作成され、その両方に部分文字列 add4 が含まれます。

template __global__ void add4<float>(float *, const float *);
template __global__ void add4<double>(double *, const double *);

PTX には以下のように該当するエントリがあります。

_Z4add4IfEvPT_PKS0_
_Z4add4IdEvPT_PKS0_

エントリポイント add4If は浮動小数点バージョン、add4Id は倍精度小数点バージョンに使用します。

k = parallel.gpu.CUDAKernel("test.ptx","double *, const double *","add4Id");

スレッド数の指定

以下の 2 つのオブジェクトプロパティを設定して、CUDAKernel の計算スレッド数を指定します。

GridSize — 3 つの要素があるベクトルで、その積によりブロック数が決定されます。
ThreadBlockSize — 3 つの要素があるベクトルで、その積によりブロックあたりのスレッド数が決定されますこの積が MaxThreadsPerBlock プロパティの値を超えることはできません。

いずれのプロパティも既定値は [1 1 1] ですが、ここでは 500 個のスレッドを使用して、500 個の要素があるベクトルで要素単位の処理を並列に実行する場合を考えてみます。要素の積が 500 になるように ThreadBlockSize を設定します。

k = parallel.gpu.CUDAKernel("myfun.ptx","myfun.cu");
k.ThreadBlockSize = [500,1,1];

一般的に、グリッドとスレッドのブロックサイズは入力のサイズに基づいて設定します。スレッド階層と複数次元のグリッドとブロックの詳細については、『NVIDIA CUDA C Programming Guide』を参照してください。

C プロトタイプ入力による CUDAKernel オブジェクトの構築. PTX ファイルに対応する CU ファイルがない場合は、CU ファイルの代わりに C カーネルの C プロトタイプを指定できます。以下に例を示します。

k = parallel.gpu.CUDAKernel("myfun.ptx","float *, const float *, float");

C プロトタイプ入力は、ソースコードでサポートされているデータ型の名前が変更されていて、その変更が認識されていない場合にも使用されます。カーネルは以下のコードで構成されているものとします。

typedef float ArgType;
__global__ void add3( ArgType * v1, const ArgType * v2 )
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

ArgType 自体はサポートされているデータ型として認識されていないため、これが含まれている CU ファイルを MATLAB で CUDAKernel オブジェクトを作成する場合に入力として直接使用することはできません。ただし、add3 カーネルの入力としてサポートされているデータ型は、CUDAKernel コンストラクターへの C プロトタイプ入力として指定できます。以下に例を示します。

k = parallel.gpu.CUDAKernel("test.ptx","float *, const float *","add3");

サポートされるデータ型. サポートされている標準の C/C++ データ型は次の表のとおりです。

浮動小数点型整数型 boolean 型と文字型

浮動小数点型	整数型	boolean 型と文字型
`double`, `double2` `float`, `float2`	`short`, `unsigned short`, `short2`, `ushort2` `int`, `unsigned int`, `int2`, `uint2` `long`, `unsigned long`, `long2`, `ulong2` `long long`, `unsigned long long`, `longlong2`, `ulonglong2` `ptrdiff_t`, `size_t`	`bool` `char`, `unsigned char`, `char2`, `uchar2`

double, double2

float, float2

short, unsigned short, short2, ushort2

int, unsigned int, int2, uint2

long, unsigned long, long2, ulong2

long long, unsigned long long, longlong2, ulonglong2

ptrdiff_t, size_t

bool

char, unsigned char, char2, uchar2

また、プログラムで tmwtypes.h ヘッダーファイルを使用する場合は以下の整数型もサポートされます。

整数型

整数型
`int8_T`, `int16_T`, `int32_T`, `int64_T` `uint8_T`, `uint16_T`, `uint32_T`, `uint64_T`

int8_T, int16_T, int32_T, int64_T

uint8_T, uint16_T, uint32_T, uint64_T

出荷状態ではヘッダーファイルは matlabroot/extern/include/tmwtypes.h です。次の行を使用すると、プログラムにこのファイルを含めることができます。

#include "tmwtypes.h"

引数の制限. 入力はすべてスカラーまたはポインターとすることができ、const を使用して定数値のラベルを付けることができます。

カーネルの C 宣言は必ず次の形式を取ります。

__global__ void aKernel(inputs ...)

カーネルは何も返さず、入力引数 (スカラーまたはポインター) のみを扱わなければなりません。
カーネルはどのような形式のメモリも割り当てることができないため、すべての出力はカーネルの実行前に割り当てておかなければなりません。このため、カーネルを実行する前にすべての出力のサイズを把握しておかなければなりません。
原則的に、カーネルに渡される、const のラベルが付けられていないすべてのポインターは出力データを格納できます。カーネルの多数のスレッドでそのデータが変更される可能性があるためです。

C 内のカーネルの定義を MATLAB に変換する際は、以下のルールが該当します。

C 内のすべてのスカラー入力 (double、float、int など) は、MATLAB のスカラーまたはスカラーの (つまり 1 要素の) gpuArray 変数でなければなりません。
C 内のすべての定数ポインター入力 (const double * など) は、MATLAB のスカラーまたは行列にすることができます。これらは正しい型にキャストされ、デバイスにコピーされて、最初の要素のポインターがカーネルに渡されます。元のサイズについての情報はカーネルに渡されません。カーネルが mxArray にある mxGetData の結果を直接受け取ったような形となります。
C 内のすべての非定数ポインター入力は、そのまま非定数ポインターとしてカーネルに転送されます。ただし、非定数ポインターはカーネルによって変更される可能性があるため、カーネルからの出力と見なされます。
MATLAB ワークスペーススカラーや配列からの入力は、目的の型にキャストされた後でカーネルに渡されます。ただし、gpuArray 入力は自動的にキャストされないので、その型や実数/複素数は想定されるものと正確に一致しなければなりません。

これらのルールにはいくつかの意味があります。最も注目すべきは、カーネルからの出力がすべて同時にカーネルへの入力でもなければならないということで、これは、(GPU にメモリを割り当てられないことから) 入力によってユーザーが出力のサイズを定義できるためです。

`CUDAKernel` の実行

入力と出力の対応の決定

GPU で CUDAKernel を評価するには、関数 feval を使用します。

いくつかのカーネルが既に作成されており、それらを MATLAB で使用して GPU で実行するとします。2 つのベクトルの畳み込みを行うカーネルが 1 つあるので、2 つの乱数入力ベクトルと共に読み込み、実行します。

k = parallel.gpu.CUDAKernel("conv.ptx","conv.cu");

result = feval(k,rand(100,1),rand(100,1));

出力は gpuArray になります (入力がそうなっていない場合でも)。ただし、カーネルを実行する際は、入力として gpuArray オブジェクトを使用すると効率が高まることがあります。

k = parallel.gpu.CUDAKernel("conv.ptx","conv.cu");

i1 = rand(100,1,"single","gpuArray");
i2 = rand(100,1,"single","gpuArray");

result1 = feval(k,i1,i2);

出力は gpuArray であるため、GPU メモリとホストメモリの間での追加転送なしに、この入出力データを使用して他の演算を実行できるようになります。

入力と出力の対応の決定

[out1, out2] = feval(kernel,in1,in2,in3) を呼び出すときに、入力 in1、in2 および in3 は CU ファイル内の関数の各入力引数に対応します。出力 out1 と out2 には、カーネルの実行後に、関数の 1 番目と 2 番目の非定数ポインター入力引数の値が格納されます。

たとえば、CU ファイル内のカーネルに次のシグネチャがある場合を考えます。

void reallySimple( float * pInOut, float c )

この場合、MATLAB 内の対応するカーネルオブジェクト (k) は以下のプロパティをもちます。

MaxNumLHSArguments: 1
   NumRHSArguments: 2
     ArgumentTypes: {'inout single vector'  'in single scalar'}

したがって、このコードのカーネルオブジェクトを feval で使用するには、feval に (カーネルオブジェクトに加えて) 2 つの入力引数を指定しなければなりません。出力引数は 1 つ使用できます。

y = feval(k,x1,x2)

入力値 x1 と x2 は、関数のプロトタイプ内の pInOut と c に対応します。出力引数 y は、カーネル実行後における、関数のプロトタイプ内の pInOut の値に対応します。

以下はこれよりもやや複雑な例で、定数ポインターと非定数ポインターの組み合わせを示しています。

void moreComplicated( const float * pIn, float * pInOut1, float * pInOut2 )

この場合、MATLAB 内の対応するカーネルオブジェクトは以下のプロパティをもちます。

MaxNumLHSArguments: 2
   NumRHSArguments: 3
     ArgumentTypes: {'in single vector'  'inout single vector'  'inout single vector'}

3 つの入力引数および 2 つの出力引数を使用して、このコードのカーネル (k) に対して feval を使用できます。

[y1,y2] = feval(k,x1,x2,x3)

3 つの入力引数 x1、x2 および x3 は、関数に渡される 3 つの引数に対応します。出力引数 y1 および y2 は、カーネル実行後における pInOut1 と pInOut2 の値に対応します。

2 つの数値の追加

次の例では、GPU に 2 つの double 値を同時に追加しています。

これを行う CU のコードは、以下のようになります。
```
__global__ void add1( double * a, double b ) 
{
    *a += b;
}
```
命令 __global__ は、これがカーネルへのエントリポイントであることを示します。このコードではポインターを使用して、結果を入出力兼用の a に入れて送信しています。このコードを現在のディレクトリにある test.cu というファイルに保存します。
mexcuda を使用して CU のコードをコンパイルして、test.ptx という PTX ファイルを生成します。
```
mexcuda -ptx test.cu
```
MATLAB でカーネルを作成します。現時点では、この PTX ファイルには 1 つのエントリしかないため、エントリを指定する必要はありません。PTX ファイルに複数のカーネルエントリポイントが含まれていた場合は、エントリポイントとして add1 を指定することになります。
```
k = parallel.gpu.CUDAKernel("test.ptx","test.cu");
```
2 つの数値入力を指定して、カーネルを実行します。既定では、カーネルは 1 つのスレッドで実行されます。
```
result = feval(k,2,3)
```
```
result = 
    5
```

2 つのベクトルの追加

次の例は前の例を拡張したもので、2 つのベクトルを同時に追加しています。簡略化のため、ベクトルの成分とちょうど同数のスレッドがあり、スレッドブロックは 1 つだけであるとします。

CU のコードは前の例と少し異なります。入力は両方ともポインターですが、一方は変更がなされないため定数です。各スレッドでは、単純にスレッドインデックスに要素を追加します。スレッドインデックスは、このスレッドがどの成分を追加するかを解明するために使用しますこうしたスレッド固有およびブロック固有の値の取得は、CUDA のプログラミングでは極めて一般的なパターンです。
```
__global__ void add2( double * v1, const double * v2 ) 
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}
```
このコードを test.cu のファイルに保存します。
mexcuda を使用して CU ファイルから test.ptx という PTX ファイルをコンパイルします。
```
mexcuda -ptx test.cu
```
このコードを最初の例のコードと同じ CU ファイルに格納する場合、今回は違いがわかるようエントリポイント名を指定します。
```
k = parallel.gpu.CUDAKernel("test.ptx","test.cu","add2");
```
カーネルを実行する前に、追加するベクトルのスレッド数を正確に設定する必要があります。
```
N = 128;
k.ThreadBlockSize = N;
in1 = ones(N,1,"gpuArray");
in2 = ones(N,1,"gpuArray");
result = feval(k,in1,in2);
```

CU および PTX ファイルを使用する例

CUDA の操作方法を説明し、試用できる CU ファイルおよび PTX ファイルを提供している例は、GPU コンピューティングへの 3 つのアプローチの説明: マンデルブロ集合を参照してください。

参考

mexcuda | CUDAKernel | feval