Main Content

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 プラグマでは、ソフトウェアで実行されるすべての並列ループ解析のチェックがオーバーライドされます。coder.gpu.kernel プラグマを使用する前にまず、coder.gpu.kernelfun を使用してください。

メモ

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

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® では効果がありません。

例

すべて折りたたむ

この例では、kernel プラグマを関数で使用し、CUDA® コードを生成する方法について説明します。

1 つのファイル内に、サイズ 1x4096 の 2 つのベクトル入力 (x および y) と、1 つのスカラー入力 scale を受け入れるエントリポイント関数 scalars を記述します。この関数には反復の長さが異なる 2 つの for ループがあります。1 つはベクトル加算用で、1 つは累積和を求めるためのループです。ベクトル加算ループの直前に coder.gpu.kernel(-1,128) プラグマを配置します。このプラグマは、ブロック数を既定に設定し、ブロックごとに 128 スレッドを割り当てるカーネルを作成します。ループ パラメーターから計算した次元のカーネルを生成するには、累積和ループの直前に coder.gpu.kernel() プラグマを配置します。

function [vout, sout1] = scalars(x,y,scale)
    sout1 = 0;
    vout = zeros(size(x));
    
    coder.gpu.kernel(-1,128);
    for i=1:1024
        vout(i) = x(i) + y(i);
    end
    
    coder.gpu.kernel();
    for i=1:4096
        sout1 = (x(i)*scale) + sout1;    
    end
end

関数 codegen を使用して CUDA MEX 関数を生成します。

codegen -config coder.gpuConfig('mex')...
 -args {ones(1,4096,'double'),ones(1,4096,'double'),coder.typeof(0)}...
 -report scalars

GPU Coder は、ベクトル加算用の scalars_kernel1 と累積和用の scalars_kernel2 カーネルの 2 つのカーネルを作成します。sout1=0 に初期化するためにカーネルは必要ありません。

  scalars_kernel1<<<dim3(8U, 1U, 1U), dim3(128U, 1U, 1U)>>>(*gpu_y, *gpu_x,
                                                            *gpu_vout);
  scalars_kernel2<<<dim3(4U, 1U, 1U), dim3(1024U, 1U, 1U)>>>(scale, *gpu_x,
                                                             gpu_sout1);

scalars_kernel1 には 1 ブロックあたり 128 スレッドを含む 8 つのブロック (合計 1024 スレッド) が用意され、それぞれが各要素を加算します。scalars_kernel2 には 1 ブロックあたり 1024 スレッドを含む 4 つのブロックが用意され、合計では 4096 スレッドになります。

カーネルの次元を指定する際に変数または式を使用できます。たとえば、グリッドとブロックの次元がコンパイル時に指定されるように、エントリポイント関数 scalars を書き換えることができます。 

function [vout, sout1] = scalars(x,y,scale, a)
    sout1 = 0;
    vout = zeros(size(x));
    
    coder.gpu.kernel(1,a);
    for i=1:1024
        vout(i) = x(i) + y(i);
    end
    
    coder.gpu.kernelfun();
    for i=1:length(x)
        sout1 = (x(i)*scale) + sout1;    
    end
end

関数 codegen を使用して CUDA MEX 関数を生成します。

codegen -config coder.gpuConfig('mex')...
 -args {ones(1,4096,'double'),ones(1,4096,'double'),20,1024}...
 -report scalars

バージョン履歴

R2017b で導入

参考

アプリ

関数

オブジェクト

トピック