Main Content

gpucoder.reduce

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

R2019b 以降

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

構文

S = gpucoder.reduce(A,FUN)
S = gpucoder.reduce(A,{@FUN1,@FUN2,...})
S = gpucoder.reduce(___,Name,Value)

説明

例

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

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

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

S = gpucoder.reduce(___,Name,Value) は、1 つ以上の Name,Value ペアの引数で指定されたオプションを使用して、入力配列に存在する値を集計します。

例

すべて折りたたむ

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

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

function [s1,s2] = multireduce(A,dim)
  [s1,s2] = gpucoder.reduce(A, {@mysum, @mymax},"dim",dim); 
end

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

function c = mymax(a, b)
  c = max(a,b);
end

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

inputArgs = {rand(32,32,"double"),coder.Constant(2)};
cfg = coder.gpuConfig('mex');
codegen -config cfg -args inputArgs -report multireduce

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

...
cudaMalloc(&gpu_iv, 8UL);
cudaMalloc(&gpu_s2, 256UL);
cudaMalloc(&gpu_s1, 256UL);
cudaMalloc(&gpu_A, 8192UL);
multireduce_kernel1<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(*gpu_iv);
cudaMemcpy(cpu_iv, *gpu_iv, 8UL, cudaMemcpyDeviceToHost);
validLaunchParams = mwGetLaunchParameters1D(static_cast<real_T>(cpu_iv[0U]),
                                            &grid, &block, 1024U, 65535U);
if (validLaunchParams) {
  cudaMemcpy(*gpu_A, cpu_A, 8192UL, cudaMemcpyHostToDevice);
  multireduce_kernel2<<<grid, block>>>(*gpu_A, *gpu_iv, *gpu_s1);
} else {
  cudaMemcpy(*gpu_A, cpu_A, 8192UL, cudaMemcpyHostToDevice);
}
coder_reduce0<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(*gpu_A, *gpu_s1);
multireduce_kernel3<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(*gpu_iv);
cudaMemcpy(cpu_iv, *gpu_iv, 8UL, cudaMemcpyDeviceToHost);
validLaunchParams = mwGetLaunchParameters1D(static_cast<real_T>(cpu_iv[0U]),
                                            &grid, &block, 1024U, 65535U);
if (validLaunchParams) {
  multireduce_kernel4<<<grid, block>>>(*gpu_A, *gpu_iv, *gpu_s2);
}
coder_reduce1<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(*gpu_A, *gpu_s2);
cudaMemcpy(cpu_s1, *gpu_s1, 256UL, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_s2, *gpu_s2, 256UL, cudaMemcpyDeviceToHost);
cudaFree(*gpu_A);
cudaFree(*gpu_s1);
cudaFree(*gpu_s2);
cudaFree(*gpu_iv);
...

入力引数

すべて折りたたむ

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

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

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

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

名前と値の引数

オプションの引数のペアを Name1=Value1,...,NameN=ValueN として指定します。ここで、Name は引数名で、Value は対応する値です。名前と値の引数は他の引数の後に指定しなければなりませんが、ペアの順序は重要ではありません。

R2021a より前では、コンマを使用して名前と値をそれぞれ区切り、Name を引用符で囲みます。

例: gpucoder.reduce(A, {@mySum, @myMax},'dim',2);

指定された次元に沿って削減を実行します。

例: gpucoder.reduce(A, {@mySum, @myMax},'dim',2);

リダクション演算を実行する前に、入力配列の要素に前処理関数を適用します。

例: gpucoder.reduce(A,@mySum,'preprocess',@myScale);

出力引数

すべて折りたたむ

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

制限

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

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

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

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

バージョン履歴

R2019b で導入

参考

アプリ

関数

オブジェクト

トピック