coder.gpu.kernel

for ループを GPU カーネルにマッピングするプラグマ

ページ内をすべて折りたたむ

構文

coder.gpu.kernel()
coder.gpu.kernel(B,T)
coder.gpu.kernel(B,T,M,name)

説明

coder.gpu.kernel() は、for ループの直前に配置しなければならないループ レベルのプラグマです。このプラグマはカーネルを生成し、ループ パラメーターから起動パラメーターを計算します。

coder.gpu.kernel プラグマは、すべての並列ループ解析のチェックをオーバーライドします。このオーバーライドにより、GPU Coder™ は、すべての反復が依存していないことを並列ループ解析が証明できない状況で、ループを並列化できるようになります。並列ループ解析のチェックにパスした関数内のループを並列化するには、coder.gpu.kernelfun の使用を検討してください。

メモ

リダクションを含む for ループの前に coder.gpu.kernel プラグマを使用することは推奨されません。

例

coder.gpu.kernel(B,T) は、B および T で指定した次元のカーネルを生成します。B[Bx,By,Bz] は、次元 x および y (z は使用しない) に沿ったグリッドのブロック数を定義する配列です。T[Tx,Ty,Tz] は、次元 x、y、および z に沿ったブロックのスレッド数を定義する配列です。

B および T の値が -1 の場合は、GPU Coder で自動的にグリッドとブロックの次元を推測しなければならないことを示します。coder.gpu.kernel プラグマは、無効なグリッドとブロックの次元に対してエラーを生成します。

例

coder.gpu.kernel(B,T,M,name) は、オプションの引数 M と name を指定します。M は、ストリーミング マルチプロセッサあたりの最小ブロック数を指定する正の整数です。M を増やすとカーネル内でのレジスタ使用量が減少し、カーネルの占有率が改善する可能性があります。M の値が -1 の場合は、GPU Coder で既定値の 1 を使用しなければならないことを示します。name は、生成されたカーネルの名前のカスタマイズを可能にする文字配列です。

この関数はコード生成関数です。MATLAB® では効果がありません。

例

すべて折りたたむ

この例では、coder.gpu.kernel プラグマを使用して CUDA® カーネルを生成する方法について説明します。

1 行 1024 列の 2 つの入力ベクトル a と b に対して要素ごとの乗算を実行する multiplyVectors という名前の関数を作成します。この関数には、ベクトルの要素を乗算する for ループが含まれています。

function out = multiplyVectors(a,b) %#codegen
out = zeros(size(a));

for i=1:size(a,2)
    out(i) = a(i)*b(i);
end
end

for ループからカーネルを生成するには、for ループの前に coder.gpu.kernel プラグマを追加します。ループ パラメーターからカーネル起動パラメーターを計算するには、入力引数なしで coder.gpu.kernel プラグマを指定します。

function out = multiplyVectors(a,b) %#codegen
out = zeros(size(a));

coder.gpu.kernel();
for i=1:size(a,2)
    out(i) = a(i)*b(i);
end
end

codegen コマンドを使用して、multiplyVectors からコードを生成します。生成されたコードには、multiplyVectors_kernel1 という名前のカーネルが含まれています。

cfg = coder.gpuConfig("mex");
a = ones([1 1024]);
b = ones([1 1024]);
codegen -config cfg -args {a,b} -report multiplyVectors

この例では、coder.gpu.kernel プラグマを使用して CUDA カーネルを生成し、起動パラメーターを指定する方法について説明します。

1 行 4096 列の 2 つの入力 x と y を受け入れる addVectors という名前の関数を作成します。この関数には、x と y を加算する 1 つの for ループがあります。 

function out = addVectors(x,y) %#codegen
out = zeros(size(x));
    
for i=1:size(x,2)
    out(i) = x(i)+y(i);
end
end

カーネルを作成するには、ベクトル加算ループの直前に coder.gpu.kernel プラグマを配置します。ブロック数を自動的に決定するには、ブロック数に -1 を指定し、1 ブロックあたり 128 スレッドを指定します。 

function out = addVectors(x,y) %#codegen
out = zeros(size(x));
    
coder.gpu.kernel(-1,128);
for i=1:size(x,2)
    out(i) = x(i)+y(i);
end
end

codegen コマンドを使用して CUDA コードを生成します。

cfg = coder.gpuConfig("mex");
x = ones([1 4096]);
y = ones([1 4096]);
codegen -config cfg -args {x,y} -report addVectors

生成されたコードには、addVectors_kernel1 という名前のカーネルが含まれています。カーネルは、32 個のブロック (1 ブロックあたり 128 スレッド) で起動します。

addVectors_kernel1<<<dim3(32U, 1U, 1U), dim3(128U, 1U, 1U)>>>(*gpu_y, *gpu_x,
                                                                *gpu_out);

カーネル起動パラメーターを指定する際に変数または式を使用できます。たとえば、coder.gpu.kernel を使用して、addVectors 関数に T という名前の入力引数を追加し、T をスレッド数として指定できます。 

function out = addVectors(x,y,T) %#codegen
out = zeros(size(x));

coder.gpu.kernel(1,T);
for i=1:size(x,2)
    out(i) = x(i)+y(i);
end
end

codegen 関数を使用して CUDA コードを生成します。生成されたコードは、入力変数 T を使用して各ブロックのスレッド数を決定します。

cfg = coder.gpuConfig("dll");
x = ones([1 4096]);
y = ones([1 4096]);
T = 512;
codegen -config cfg -args {x,y,T} -report addVectors

バージョン履歴

R2017b で導入

参考

アプリ

関数

オブジェクト

トピック