Main Content

スキャッター/ギャザー タイプ演算のカーネル

GPU Coder™ はリダクションの概念もサポートしています。この概念は、ループ反復が独立していなければならないという規則に対する重要な例外です。リダクション変数は、すべての反復に依存するが、反復の順序には依存しない値を累積します。リダクション変数は、加算、ドット積、並べ替えなどにおける代入ステートメントの両辺に使用されます。次の例は、リダクション変数 x の一般的な使用法を示しています。

x = ...; % Some initialization of x
for i = 1:n
  x = x + d(i);
end

各反復の変数 x は、ループに入る前に値を取得するか、ループの前回の反復から値を取得します。こうした連続タイプの実装は、逐次実行の中で依存が連鎖するため、並列実行に適しません。別の方法は、二分木ベースの方法を採用することです。

この木ベースの方法では、特定の通過回数を超えた木のすべての水平レベルを並列に実行できます。逐次実行と比較すると、二分木の方がより多くのメモリを必要とします。それは、通過ごとに一時的な値の配列を出力として必要とするためです。得られるパフォーマンス上のメリットはメモリ使用量の増加というコストを大幅に上回ります。GPU Coder は、この木ベースの方法を使用してリダクション カーネルを作成します。そこでは、スレッド ブロックごとにこの配列の一部がリダクションされます。並列リダクションでは、スレッド ブロック間で部分的な結果のデータを交換する必要があります。古い CUDA® デバイスでは、このデータ交換は共有メモリとスレッド同期を使って実行されていました。Kepler GPU アーキテクチャ以降、CUDA によりシャッフル (shfl) 命令と高速なデバイス メモリ アトミック操作が提供され、リダクションはさらに高速化されています。GPU Coder によって作成されるリダクション カーネルは、shfl_down 命令を使用してスレッドのワープ (32 個のスレッド) 全体に対してリダクションを実行します。次に、各ワープの最初のスレッドでアトミック操作命令を使用してリダクション済みの値を更新します。

命令の詳細については、NVIDIA® ドキュメンテーションを参照してください。

ベクトル和の例

この例では、GPU Coder を使用して CUDA リダクション タイプ カーネルを作成する方法を示します。ベクトル v を作成し、その要素の合計を計算するとします。この例を MATLAB® 関数として実装できます。

function s = VecSum(v)
    s = 0;
    for i = 1:length(v)
       s = s + v(i);
    end
end

カーネル作成用の vecSum の準備

GPU Coder は、リダクション カーネルを推定するために特殊なプラグマを必要としません。この例では、coder.gpu.kernelfun プラグマを使用して CUDA リダクション カーネルを生成します。変更された関数 VecSum を使用します。

メモ

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

function s = VecSum(v) %#codegen
    s = 0;
    
    coder.gpu.kernelfun();
    for i = 1:length(v)
       s = s + v(i);
    end
end

生成された CUDA コード

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

static __global__ __launch_bounds__(512, 1) void vecSum_kernel1(const real_T *v,
  real_T *s)
{
  uint32_T threadId;
  uint32_T threadStride;
  uint32_T thdBlkId;
  uint32_T idx;
  real_T tmpRed;
  ;
  ;
  thdBlkId = (threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x)
    + threadIdx.x;
  threadId = ((gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y) +
              blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) + thdBlkId;
  threadStride = gridDim.x * blockDim.x * (gridDim.y * blockDim.y) * (gridDim.z *
    blockDim.z);
  if (!((int32_T)threadId >= 512)) {
    tmpRed = 0.0;
    for (idx = threadId; threadStride < 0U ? idx >= 511U : idx <= 511U; idx +=
         threadStride) {
      tmpRed += v[idx];
    }

    tmpRed = workGroupReduction1(tmpRed, 0.0);
    if (thdBlkId == 0U) {
      atomicOp1(s, tmpRed);
    }
  }
}

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

  cudaMemcpy((void *)gpu_v, (void *)v, 4096ULL, cudaMemcpyHostToDevice);
  cudaMemcpy((void *)gpu_s, (void *)&s, 8ULL, cudaMemcpyHostToDevice);
  VecSum_kernel1<<<dim3(1U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_v, gpu_s);
  cudaMemcpy(&s, gpu_s, 8U, cudaMemcpyDeviceToHost);

メモ

パフォーマンスを向上させるために、GPU Coder はリダクションよりもカーネルの並列化を優先させます。アルゴリズムで並列ループ内にリダクションが含まれている場合、GPU Coder はそのリダクションを通常のループとして推測し、そのためのカーネルを生成します。

GPU での 1 次元リダクション演算

関数 gpucoder.reduce を使用して、1 次元リダクション演算を GPU で効率的に実行する CUDA コードを生成できます。生成されたコードは CUDA のシャッフル内部パラメーターを使用してリダクション演算を実装します。

たとえば、配列 Asum 要素と max 要素を求めるには、次のようにします。

function s = myReduce(A)
   s = gpucoder.reduce(A, {@mysum, @mymax}); 
end

function c = mysum(a, b)
   c = a+b;
end

function c = mymax(a, b)
   c = max(a,b);
end
コード生成では、関数 gpucoder.reduce には次の要件があります。

  • 入力は数値型または logical データ型でなければならない。

  • @handle を介して渡される関数は 2 つの入力を受け入れて 1 つの出力を返す二項関数でなければならない。入力と出力はデータ型が同じでなければならない。

  • 関数は可換的かつ結合的でなければならない。

メモ

入力が整数データ型である場合、関数 gpucoder.reduce 用に生成されたコードで、中間計算が飽和する可能性があります。このような場合、生成コードから得られる結果は MATLAB のシミュレーション結果と一致しない場合があります。

参考

| | | | |

関連するトピック