Main Content

このページの翻訳は最新ではありません。ここをクリックして、英語の最新版を参照してください。

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

CUDA ツールキットの要件

CUDA® カーネル オブジェクトを CU コードから生成する場合や、GPU Coder™ を使用して CUDA 互換のソース コード、ライブラリ、および実行可能ファイルをコンパイルする場合は、CUDA ツールキットをインストールする必要があります。CUDA ツールキットには、CUDA のライブラリおよびコンパイル用ツールが含まれています。

メモ

GPU 上で MATLAB® 関数を実行する場合や、CUDA 対応の MEX 関数を生成する場合は、このツールキットは必要ありません。

タスク要件

  • gpuArray と GPU 対応の MATLAB 関数を使用する。

  • GPU Coder または mexcuda を使用して CUDA 対応の MEX 関数をコンパイルする。

NVIDIA Driver Downloads から、最新のグラフィックス ドライバーを入手します。

CUDA ツールキットは不要です。

  • CU コードから CUDA カーネル オブジェクトを作成する。*

  • GPU Coder を使用して CUDA 互換のソース コード、ライブラリおよび実行可能ファイルをコンパイルする。

ご利用の MATLAB リリースでサポートされているバージョンの CUDA ツールキットをインストールします。

* MATLAB で CUDA カーネル オブジェクトを作成するには、CU ファイルおよび対応する PTX ファイルの両方が必要です。CU ファイルから PTX ファイルをコンパイルするには、CUDA ツールキットが必要です。対応する PTX ファイルが既にある場合は、ツールキットは不要です。CUDA でサポートされているコンパイラの一部は、MATLAB でサポートされていません。

必要なツールキットのバージョンは、使用している MATLAB のバージョンによって異なります。以下の表で、使用している MATLAB のバージョンと互換性のあるツールキットのバージョンを確認してください。推奨されるベスト プラクティスは、サポートされているツールキットの最新バージョン (NVIDIA® からのすべての更新やパッチを含む) を使用することです。

MATLAB リリースCUDA ツールキットのバージョン
R2022b11.2
R2022a11.2
R2021b11.0
R2021a11.0
R2020b10.2
R2020a10.1
R2019b10.1
R2019a10.0
R2018b9.1
R2018a9.0
R2017b8.0
R2017a8.0
R2016b7.5
R2016a7.5
R2015b7.0
R2015a6.5
R2014b6.0
R2014a5.5
R2013b5.0
R2013a5.0
R2012b4.2
R2012a4.0
R2011b4.0

CUDA ツールキットの詳細、およびサポートされているバージョンのダウンロードについては、CUDA Toolkit Archive (NVIDIA) を参照してください。

CUDAKernel ワークフローの概要

このトピックでは、実行可能カーネルを CU または PTX (並列スレッド実行) ファイルから作成し、そのカーネルを MATLAB から GPU 上で実行する方法を説明します。MATLAB ではカーネルは CUDAKernel オブジェクトで表されます。このオブジェクトは MATLAB 配列または gpuArray 変数を処理できます。

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

  1. コンパイルされた PTX コードを使用して CUDAKernel オブジェクトを作成します。このオブジェクトには GPU 実行可能コードが含まれています。

  2. CUDAKernel オブジェクトのプロパティを設定し、GPU での実行を制御できるようにします。

  3. 必要な入力を使用して CUDAKernel でfeval を呼び出し、GPU でカーネルを実行します。

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

% 1. Create CUDAKernel object.
k = parallel.gpu.CUDAKernel('myfun.ptx','myfun.cu','entryPt1');

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

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

result = feval(k,g1,g2);

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

CUDAKernel オブジェクトの作成

CU ファイルからの PTX ファイルのコンパイル

GPU で実行する CU ファイルがある場合は、まずそれをコンパイルして PTX ファイルを作成しなければなりません。これを行う 1 つの方法は、NVIDIA CUDA ツールキットの nvcc コンパイラを使うことです。たとえば、CU ファイルが myfun.cu という名前の場合は、次のシェル コマンドを使ってコンパイルした PTX ファイルを作成できます。

nvcc -ptx myfun.cu

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

CU ファイル入力による CUDAKernel オブジェクトの構築

MATLAB では .cu ファイルと .ptx ファイルを使用して CUDAKernel オブジェクトを作成できます。このオブジェクトはカーネルの評価に使用できます。

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

メモ

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

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 型と文字型

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

出荷状態ではヘッダー ファイルは 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 ポインター入力 (const double * など) は、MATLAB のスカラーまたは行列にすることができます。これらは正しい型にキャストされ、デバイスにコピーされて、最初の要素のポインターがカーネルに渡されます。元のサイズについての情報はカーネルに渡されません。カーネルが mxArray にある mxGetData の結果を直接受け取ったような形となります。

  • C 内のすべての非定数ポインター入力は、そのまま非定数ポインターとしてカーネルに転送されます。ただし、非定数ポインターはカーネルによって変更される可能性があるため、カーネルからの出力と見なされます。

  • MATLAB ワークスペース スカラーや配列からの入力は、目的の型にキャストされた後でカーネルに渡されます。ただし、gpuArray 入力は自動的にキャストされないので、その型や実数/複素数は想定されるものと正確に一致しなければなりません。

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

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 オブジェクトのリファレンス ページを参照してください。設定可能なプロパティを変更する一般的な理由として、以下のようにスレッド数を指定することがあげられます。

エントリ ポイントの指定

PTX ファイルに複数のエントリ ポイントがある場合は、k カーネル オブジェクトで参照するよう、myfun.ptx 内の特定のカーネルを特定できます。

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

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

__global__ void simplestKernelEver( float * x, float val )

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

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

メモ

関数 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 個の要素があるベクトルで要素単位の処理を並列に実行する場合を考えてみます。この処理を実行する簡単な方法として、CUDAKernel を作成してプロパティを適切に設定することがあげられます。

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

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

CUDAKernel の実行

GPU で CUDAKernel を評価するには、関数 feval を使用します。以下の例では、MATLAB ワークスペース変数と gpuArray 変数を使用してカーネルを実行する方法を説明します。

ワークスペース変数の使用

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

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

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

入力が MATLAB ワークスペース データの定数または変数であっても、出力は gpuArray になります。

gpuArray 変数の使用

カーネルを実行する際は、入力として gpuArray オブジェクトを使用すると効率が高まることがあります。

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

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

result1 = feval(k,i1,i2);

出力は gpuArray であるため、MATLAB ワークスペースと GPU の間での追加転送なしに、この入出力データを使用して他の演算を実行できるようになります。GPU 計算がすべて完了したら、最終結果のデータを MATLAB ワークスペースに集めます。

result2 = feval(k,i1,i2);

r1 = gather(result1);
r2 = gather(result2);

入力と出力の対応の決定

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

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

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 は、C 関数のプロトタイプ内の pInOut と c に対応します。出力引数 y は、C カーネル実行後における、C 関数のプロトタイプ内の 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'}

このコードのカーネル (k) に対し、feval を次の構文で使用できます。

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

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

カーネル ワークフローの全体

2 つの数値の追加

次の例では、GPU に 2 つの double 値を同時に追加しています。NVIDIA CUDA ツールキットをインストールし、デバイス用に CUDA 対応のドライバーを用意しておく必要があります。

  1. これを行う CU のコードは、以下のようになります。 

    __global__ void add1( double * pi, double c ) 
{
    *pi += c;
}

    命令 __global__ は、これがカーネルへのエントリ ポイントであることを示します。このコードではポインターを使用して、結果を入出力兼用の pi に入れて送信しています。このコードを現在のディレクトリにある test.cu というファイルに格納します。

  2. シェル コマンド ラインで CU のコードをコンパイルして、test.ptx という PTX ファイルを生成します。

    nvcc -ptx test.cu

  3. MATLAB でカーネルを作成します。現時点では、この PTX ファイルには 1 つのエントリしかないため、エントリを指定する必要はありません。複数のカーネルを格納する場合は、エントリとして add1 を指定することになります。

    k = parallel.gpu.CUDAKernel('test.ptx','test.cu');

  4. 2 つの数値入力を指定して、カーネルを実行します。既定では、カーネルは 1 つのスレッドで実行されます。

    result = feval(k,2,3)
    result = 
    5

2 つのベクトルの追加

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

  1. CU のコードは前の例と少し異なります。入力は両方ともポインターですが、一方は変更がなされないため定数です。各スレッドでは、単純にスレッド インデックスに要素を追加します。スレッド インデックスは、このスレッドがどの成分を追加するかを解明するために使用します(こうしたスレッド固有およびブロック固有の値の取得は、CUDA のプログラミングでは極めて一般的なパターンです)。

    __global__ void add2( double * v1, const double * v2 ) 
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];
}

    このコードを test.cu のファイルに保存します。

  2. nvcc を使用して、前と同様にコンパイルします。

    nvcc -ptx test.cu

  3. このコードを最初の例のコードと同じ CU ファイルに格納する場合、今回は違いがわかるようエントリ ポイント名を指定する必要があります。

    k = parallel.gpu.CUDAKernel('test.ptx','test.cu','add2');

  4. カーネルを実行する前に、追加するベクトルのスレッド数を正確に設定する必要があります。

    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 つのアプローチの説明: マンデルブロ集合を参照してください。

参考

|

関連するトピック