Main Content

coder.gpu.nokernel

ループのカーネル作成を無効にするプラグマ

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

構文

coder.gpu.nokernel()

説明

例

coder.gpu.nokernel() はループ レベルのプラグマであり、for ループの直前に配置すると、コード ジェネレーターでそのループ内のステートメントの CUDA® カーネルが生成されなくなります。このプラグマは入力パラメーターを必要としません。

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

例

すべて折りたたむ

この例では、関数に nokernel プラグマを使用して、コード ジェネレーターでそのループ内のステートメントの CUDA カーネルが生成されないようにする方法を示します。

1 つのファイル内に、サイズ 32x512 の 2 つのベクトル入力 A,B を受け入れるエントリポイント関数 nestedLoop を記述します。この関数には反復の長さが異なる 2 つの入れ子にされた for ループがあります。1 つは列に沿った演算用で、もう 1 つは行に沿った演算用のループです。1 番目の入れ子にされたループでは 2 つのベクトル入力の合計を計算し、2 番目の入れ子にされたループではその合計を 3 倍にスケーリングしています。

function [C] = nestedLoop(A, B)
    G = zeros(32, 512);
    C = zeros(32, 512);        
   
    coder.gpu.kernelfun();
    % This nested loop will be fused
    for i = 1:32
       for j = 1:512
           G(i,j) = A(1,j) + B(1,j);
       end
    end

    coder.gpu.nokernel();  
    for i = 1:32
       for j = 1:512
           C(i,j) = G(i,j) * 3;
       end
    end    
end

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

cfg = coder.gpuConfig('mex');
cfg.GenerateReport = true;
codegen -config cfg -args {ones(1,512,'double'),ones(1,512,'double')} nestedLoop

GPU Coder は 2 つのカーネルを作成します。nestedLoop_kernel1 は、1 番目の入れ子にされたループの計算 G(i,j) = A(1,j) + B(1,j); を実行し、nestedLoop_kernel2 カーネルは、2 番目の入れ子にされたループの計算 C(i,j) = G(i,j) * 3; を実行します。2 番目のカーネルは、2 番目の入れ子にされたループの内部ループ用に作成します。noKernel プラグマは、そのステートメントの直後のループにのみ適用されます。生成されたカーネルのスニペットを示します。

static __global__ __launch_bounds__(512, 1) void nestedLoop_kernel1(const real_T
  B[512], const real_T A[512], real_T G[16384])
{
  uint32_T threadId;
  ...
  if (i < 32) {
    G[i + (j << 5)] = A[j] + B[j];
  }
}
static __global__ __launch_bounds__(512, 1) void nestedLoop_kernel2(real_T G
  [16384], int32_T i, real_T C[16384])
{
  uint32_T threadId;
  ...;
  if (j < 512) {
    C[i + (j << 5)] = G[i + (j << 5)] * 3.0;
  }

main 関数のスニペットでは、コード ジェネレーターによって 1 番目の入れ子にされたループが、カーネルの起動パラメーターで示されるとおりに融合されていることがわかります。前述のとおり、2 番目の入れ子にされたループの外側のループはカーネルにマッピングされません。そのため、コード ジェネレーターは、for-loop ステートメントを 2 番目の CUDA カーネル nestedLoop_kernel2 の呼び出しの直前に配置します。

void nestedLoop(const real_T A[512], const real_T B[512], real_T C[16384])
{
  int32_T i;
  ...
  //  These two loops will be fused
  cudaMemcpy(gpu_B, (void *)&B[0], 4096UL, cudaMemcpyHostToDevice);
  cudaMemcpy(gpu_A, (void *)&A[0], 4096UL, cudaMemcpyHostToDevice);
  nestedLoop_kernel1<<<dim3(32U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_B, *gpu_A, *
    gpu_G);
  for (i = 0; i < 32; i++) {
    nestedLoop_kernel2<<<dim3(1U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_G, i,
      *gpu_C);
    C_dirtyOnGpu = true;
  }
...
  cudaFree(*gpu_C);
}

バージョン履歴

R2019a で導入

参考

アプリ

関数

オブジェクト

トピック