gpucoder.batchedMatrixMultiply

バッチ行列乗算演算の最適化された GPU 実装

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

構文

[D1,D2] = gpucoder.batchedMatrixMultiply(A1,B1,A2,B2)
[D1,...,DN] = gpucoder.batchedMatrixMultiply(A1,B1,...,AN,BN)
___ = gpucoder.batchedMatrixMultiply(___,Name,Value)

説明

[D1,D2] = gpucoder.batchedMatrixMultiply(A1,B1,A2,B2) は、行列 A1,B1 と A2,B2 のバッチについて、行列-行列乗算を実行します。関数 gpucoder.batchedMatrixMultiply は次の形式の行列-行列乗算を実行します。

D=αAB

ここで、α はスカラー乗算係数、A、B、および D はそれぞれ m 行 k 列、k 行 n 列、および m 行 n 列の次元をもつ行列です。A および B は、転置またはエルミート共役にすることもできます。既定の設定では、α は 1 に設定され、行列は転置されません。さまざまなスカラー乗算係数を指定し、入力行列の転置演算を実行するには、Name,Value ペアの引数を使用します。

関数 gpucoder.batchedMatrixMultiply に渡されるすべてのバッチは一様でなければなりません。つまり、すべてのインスタンスは同じ次元 m,n,k でなければなりません。

[D1,...,DN] = gpucoder.batchedMatrixMultiply(A1,B1,...,AN,BN) は、複数の A と B のペアについて、次の形式の行列-行列乗算を実行します。

Di=αAiBi                  i=1…N

___ = gpucoder.batchedMatrixMultiply(___,Name,Value) は、1 つ以上の Name,Value ペアの引数で指定されたオプションを使用することにより、バッチ行列乗算演算を実行します。

例

例

すべて折りたたむ

シンプルなバッチ行列-行列乗算を実行し、関数 gpucoder.batchedMatrixMultiply を使用して、対応する cublas<t>gemmBatched API を呼び出す CUDA® コードを生成します。

1 つのファイル内に、行列入力 A1、B1、A2、および B2 を受け入れるエントリポイント関数 myBatchMatMul を記述します。入力行列は転置しないため、'nn' オプションを使用します。 

function [D1,D2] = myBatchMatMul(A1,B1,A2,B2,alpha)

[D1,D2] = gpucoder.batchedMatrixMultiply(A1,B1,A2,B2, ...
    'alpha',alpha,'transpose','nn');

end

コード生成で使用する double の行列の型を作成するには、関数 coder.newtype を使用します。

A1 = coder.newtype('double',[15,42],[0 0]);
A2 = coder.newtype('double',[15,42],[0 0]);
B1 = coder.newtype('double',[42,30],[0 0]);
B2 = coder.newtype('double',[42,30],[0 0]);
alpha = 0.3;
inputs = {A1,B1,A2,B2,alpha};

CUDA ライブラリを生成するには、関数 codegen を使用します。

cfg = coder.gpuConfig('lib');
cfg.GpuConfig.EnableCUBLAS = true;
cfg.GpuConfig.EnableCUSOLVER = true;
cfg.GenerateReport = true;
codegen -config cfg-args inputs myBatchMatMul

生成された CUDA コードには、入力および出力行列を初期化するためのカーネル myBatchMatMul_kernelNN が含まれます。さらに、コードには cuBLAS ライブラリに対する cublasDgemmBatched API 呼び出しも含まれます。以下のコードは、生成されたコードのスニペットです。

//
// File: myBatchMatMul.cu
//
...
void myBatchMatMul(const double A1[630], const double B1[1260], const double A2
                   [630], const double B2[1260], double alpha, double D1[450],
                   double D2[450])
{
  double alpha1;
...

  myBatchMatMul_kernel1<<<dim3(2U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_A2,
    *gpu_A1, *gpu_input_cell_f2, *gpu_input_cell_f1);
  cudaMemcpy(gpu_B2, (void *)&B2[0], 10080UL, cudaMemcpyHostToDevice);
  cudaMemcpy(gpu_B1, (void *)&B1[0], 10080UL, cudaMemcpyHostToDevice);
  myBatchMatMul_kernel2<<<dim3(3U, 1U, 1U), dim3(512U, 1U, 1U)>>>(*gpu_B2,
    *gpu_B1, *gpu_input_cell_f4, *gpu_input_cell_f3);
  myBatchMatMul_kernel3<<<dim3(1U, 1U, 1U), dim3(480U, 1U, 1U)>>>(gpu_r3, gpu_r2);
  myBatchMatMul_kernel4<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(gpu_r2,
    *gpu_out_cell);
  myBatchMatMul_kernel5<<<dim3(1U, 1U, 1U), dim3(32U, 1U, 1U)>>>(gpu_r3,
    *gpu_out_cell);
...

  cublasDgemmBatched(getCublasGlobalHandle(), CUBLAS_OP_N, CUBLAS_OP_N, 15, 30,
                     42, (double *)gpu_alpha1, (double **)gpu_Aarray, 15,
                     (double **)gpu_Barray, 42, (double *)gpu_beta1, (double **)
                     gpu_Carray, 15, 2);
  myBatchMatMul_kernel6<<<dim3(1U, 1U, 1U), dim3(480U, 1U, 1U)>>>(*gpu_D2,
    *gpu_out_cell, *gpu_D1);
...
}

入力引数

すべて折りたたむ

オペランド。ベクトルまたは行列として指定します。A と B は 2 次元配列でなければなりません。A の列数は B の行数と等しくなければなりません。

データ型: double | single | int8 | int16 | int32 | int64 | uint8 | uint16 | uint32 | uint64
複素数のサポート: あり

名前と値の引数

すべて折りたたむ

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

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

例: [D1,D2] = gpucoder.batchedMatrixMultiply(A1,B1,A2,B2,'alpha',0.3,'transpose','CC');

A との乗算に使用されるスカラーの値。既定値は 1 です。

2 つの文字からなる文字ベクトルまたは string で、行列乗算の前に行列 A および B で実行される演算を示します。可能な値は、標準 ('N')、転置 ('T')、または複素共役転置 ('C') です。

出力引数

すべて折りたたむ

積。スカラー、ベクトル、または行列として返されます。配列 D には、入力 A と同じ数の行と、入力 B と同じ数の列があります。

バージョン履歴

R2020a で導入

参考

アプリ

関数

オブジェクト

トピック