Main Content

このページの翻訳は最新ではありません。ここをクリックして、英語の最新版を参照してください。

gpucoder.reduce

リダクション演算の最適化された GPU 実装

説明

S = gpucoder.reduce(A,FUN) は、指定された関数ハンドル FUN を使用して入力配列 A に含まれる値を単一値に集計します。出力 S はスカラーです。

S = gpucoder.reduce(A,{@FUN1,@FUN2,...}) は入力配列と関数ハンドルの cell 配列を受け入れます。cell 配列で提供された関数ハンドルごとに入力配列に含まれる値を単一値に集計します。出力のサイズは 1 行 N 列であり、N は関数ハンドルの数です。

コード ジェネレーターは shuffle 内部パラメーターを使用して GPU で効率的な削減を実行します。複数の関数ハンドルが GPU の単一カーネル内で集計されます。

すべて折りたたむ

この例では、CUDA® コードを生成して配列の要素の和と最大値を求めます。

1 つのファイル内に、行列入力 A を受け入れるエントリポイント関数 multireduce を作成します。関数 gpucoder.reduce を使用して、A の要素で 2 種類のリダクション演算を実行します。

function s = multireduce(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

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

codegen -config coder.gpuConfig('mex') -args {rand(1,1024,'double')} -report multireduce

以下は、生成コードのスニペットです。

...
cudaMalloc(&gpu_s, 16ULL);
cudaMalloc(&gpu_A, 8192ULL);
cudaMemcpy(gpu_A, (void *)&A[0], 8192ULL, cudaMemcpyHostToDevice);
multireduce_kernel1<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(*gpu_A, *gpu_s);
coder_reduce0<<<dim3(2U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_A, *gpu_s);
cudaMemcpy(&s[0], gpu_s, 16ULL, cudaMemcpyDeviceToHost);
...
static __inline__ __device__ real_T shflDown2(real_T in1, uint32_T offset,
  uint32_T mask)
{
  int2 tmp;
  tmp = *(int2 *)&in1;
  tmp.x = __shfl_down_sync(mask, tmp.x, offset);
  tmp.y = __shfl_down_sync(mask, tmp.y, offset);
  return *(real_T *)&tmp;
}
...

入力引数

すべて折りたたむ

リダクション演算を実行する入力配列。コード生成では、入力配列は数値または論理データ型でなければなりません。

ユーザー定義関数へのハンドル。FUN は関数ハンドルの cell 配列にすることもできます。関数ハンドルは二項関数であり、以下の要件を満たさなければなりません。

  • 2 つの入力を受け入れて、1 つの出力を返します。関数への入力と出力のタイプは入力配列 A と一致しなければなりません。

  • 関数は可換的かつ結合的でなければならず、それ以外の場合、動作は定義されません。

出力引数

すべて折りたたむ

リダクション演算の結果。削減中に、S は入力配列 A のいずれかの要素の値に初期化されます。次に、FUNAS のすべての要素に適用することで、リダクション演算が実行されます。

制限

  • gpucoder.reduce は複素数データ型の入力配列をサポートしません。

  • ユーザー定義関数は 2 つの入力を受け入れて、1 つの出力を返さなければなりません。関数への入力と出力のタイプは入力配列 A と一致しなければなりません。

  • ユーザー定義関数は可換的かつ結合的でなければならず、それ以外の場合、動作は定義されません。

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

R2019b で導入