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 プラグマを使用することは推奨しません。

coder.gpu.kernel(B,T) は、for ループの直前に配置しなければならないループ レベルのプラグマです。B および T で指定した次元のカーネルを生成します。B[Bx,By,1] は、次元 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) は、同じ引数 B および T を想定しています。オプション引数 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,1024) プラグマを配置します。このプラグマにより、1024 スレッドを含む 1 つのブロックが用意されたカーネルが作成されます。2 番目のループの外に coder.gpu.kernel() プラグマを配置します。

function [vout, sout1] = scalars(x,y,scale)
    sout1 = 0;
    vout = coder.nullcopy(zeros(size(x)));
    
    coder.gpu.kernel(1,1024);
    for i=1:1024
        vout(i) = x(i) + y(i);
    end
    
    coder.gpu.kernelfun();
    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 により 2 つのカーネルが作成されます。ベクトル加算用の scalars_kernel1 と、累積和用の scalars_kernel2 カーネルです。sout1=0 に初期化するためにカーネルは必要ありません。

  cudaMemcpy(*gpu_y, cpu_y, 32768ULL, cudaMemcpyHostToDevice);
  cudaMemcpy(*gpu_x, cpu_x, 32768ULL, cudaMemcpyHostToDevice);
  scalars_kernel1<<<dim3(1U, 1U, 1U), dim3(1024U, 1U, 1U)>>>(*gpu_y, *gpu_x,
                                                             *gpu_vout);
  cudaMemcpy(gpu_sout1, cpu_sout1, 8ULL, cudaMemcpyHostToDevice);
  scalars_kernel2<<<dim3(8U, 1U, 1U), dim3(512U, 1U, 1U)>>>(scale, *gpu_x,
                                                            gpu_sout1);
  cudaMemcpy(cpu_vout, *gpu_vout, 32768ULL, cudaMemcpyDeviceToHost);
  cudaMemcpy(cpu_sout1, gpu_sout1, 8ULL, cudaMemcpyDeviceToHost);
  cudaMemcpy(gpu_y, y, 32768U, cudaMemcpyHostToDevice);

scalars_kernel1 には、1 ブロックあたり 1024 スレッドを含む 1 つのブロックが用意され、それぞれが各要素を加算します。scalars_kernel2 カーネルには、1 ブロックあたり 512 スレッドを含む 8 つのブロックが用意され、合計では 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 で導入

参考

アプリ

関数

オブジェクト

トピック