Main Content

要素単位ループのカーネル

CUDA® カーネルの作成の最も単純なケースは、スカラー化された要素単位の数学演算を含む MATLAB® 関数から作成することです。要素単位の演算が for ループ本体に囲まれている場合、同時実行 CUDA スレッドを呼び出して各ループ反復を並列に計算できます。CUDA スレッドは特定の順には実行されず、互いに独立しているため、for ループ内の反復は他の反復の結果に依存しないようにする必要があります。

要素単位の数学の例

この例では、要素単位の数学演算を含む関数から CUDA カーネルを作成する方法を説明します。行列 x の各要素を 2 乗し、1/(i+j) の係数でスケーリングするとします。ここで、i,j は行と列のインデックスです。この例を MATLAB 関数として実装できます。

function [y] = myFun(x)

y = zeros(size(x));
for i = 1:size(x,1)
    for j = 1:size(x,2)
        y(i,j)=(x(i,j)^2)/(i+j);
    end
end
end

コード生成用の myFun の準備

関数 myFun の最初のステートメント zeros(size(A)) は、結果のベクトル y を 0 に初期化します。CUDA コード生成では、y のメモリを事前に割り当てることで、メモリを 0 に初期化するオーバーヘッドは発生しません。この行を coder.nullcopy(zeros(size(y))) に置き換えます。

ループから CUDA カーネルを作成するために、GPU Coder™ には別のプラグマ coder.gpu.kernel が用意されています。このカーネル プラグマを指定すると、すべての並列ループ解析がオーバーライドされます。パラメーターを指定しない場合、GPU Coder ではループ範囲と入力サイズに基づいてカーネル範囲が決定されます。これには、"スレッド""ブロック" のサイズなどのカーネル起動パラメーターを指定する方法が用意されています。ただし、これはループを並列化して問題がないことがわかっている場合にのみ使用します。myFun の例はシンプルで、カーネル起動パラメーターを指定する必要がないため、coder.gpu.kernelfun プラグマを利用して CUDA カーネルを生成できます。

これらの変更を加えることで、元の関数 myFun がコード生成に適したものになります。

function [y] = myFun(x) %#codegen

y = coder.nullcopy(zeros(size(x)));
coder.gpu.kernelfun();
for i = 1:size(x,1)
    for j = 1:size(x,2)
        y(i,j)=(x(i,j)^2)/(i+j);
    end
end
end

生成された CUDA コード

GPU Coder アプリまたはコマンド ラインを使用して CUDA コードを生成すると、GPU Coder によって、2 乗演算とスケーリング演算を実行する単一のカーネルが作成されます。以下は、myFun_kernel1 カーネル コードのスニペットです。

static __global__ __launch_bounds__(512, 1) void myFun_kernel1(const real_T *x,
  real_T *y)
{
...
threadId = ((((gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y) +
                blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) +
               threadIdx.z * blockDim.x * blockDim.y) + threadIdx.y * blockDim.x)
    + threadIdx.x;
  i = (int32_T)(threadId / 512U);
  j = (int32_T)(threadId - (uint32_T)i * 512U);
  if ((!(j <= 512)) && (!(i <= 512))) {
    y[i + (j << 9)] = x[i + (j << 9)] * x[i + (j << 9)] / ((real_T)(i + j) + 2.0);
  }
}

以下は、main 関数 myFun のスニペットです。myFun_kernel1 の呼び出しの前に、行列 x をホスト (x) からデバイス (gpu_x) に転送する単一の cudaMemcpy 呼び出しがあります。このカーネルには、1 ブロックあたり 512 スレッドを含む 512 個のブロックが用意されています。これは入力ベクトルのサイズと一致します。2 番目の cudaMemcpy 呼び出しは計算結果を元のホストにコピーします。

cudaMemcpy((void *)gpu_x, (void *)x, 2097152ULL, cudaMemcpyHostToDevice);
myFun_kernel1<<<dim3(512U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_x, gpu_y);
cudaMemcpy((void *)y, (void *)gpu_y, 2097152ULL, cudaMemcpyDeviceToHost);

制限

  • ループ範囲が符号なしのデータ型である場合、ループ範囲が有効かどうかを判断するための条件チェックがコード ジェネレーターによって追加される場合があります。これらの条件チェックにより、ソフトウェアによって実行される最適化が制限され、パフォーマンスの変化を引き起こす可能性のあるリダクション カーネルが導入される可能性があります。

参考

| | | |

関連するトピック