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);

制限

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

参考

| | | |

関連するトピック