このページの内容は最新ではありません。最新版の英語を参照するには、ここをクリックします。
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 型と文字型 |
---|---|---|
|
|
|
また、プログラムで tmwtypes.h
ヘッダー ファイルを使用する場合は以下の整数型もサポートされます。
整数型 |
---|
|
出荷状態ではヘッダー ファイルは
です。次の行を使用すると、プログラムにこのファイルを含めることができます。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