gpucoder.matrixMatrixKernel

行列-行列演算を含む関数の最適化された GPU 実装

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

構文

C = gpucoder.matrixMatrixKernel(FUN,A,B)

C = gpucoder.matrixMatrixKernel(FUN,A,B,orientation)

説明

例

C = gpucoder.matrixMatrixKernel(FUN,A,B) は、GEMM のような演算を含む関数からカーネルを生成します。たとえば、以下を使用して 2 つのイメージ間の特徴点をマッチさせます。

差の絶対値の和 (SAD) — F() = @(a,b)abs(a-b)
差の二乗の和 (SSD) — F() = @(a,b)(a-b).*(a-b)

FUN はユーザー定義関数へのハンドルです。行列 A から 1 つの行および列を、行列 B から 1 つの行および列を受け取り、入力と同じ型のベクトルを出力します。その後、出力ベクトルが合計され、C の単一のスカラー値が計算されます。数値入力 A および B は、同じサイズか、互換性のあるサイズでなければなりません。たとえば、A が M 行 K 列の行列、B が K 行 N 列の行列の場合、C は M 行 N 列の行列になります。

C = gpucoder.matrixMatrixKernel(FUN,A,B,orientation) のオプション引数 orientation は、A および B 行列の向きを指定します。これは次の 4 つの値のいずれかを取ることができます。

'nn' - 行列 A および B は正規です。
'nt' - 行列 B は転置されます。
'tn' - 行列 A は転置されます。
'tt' - 行列 A および B の両方が転置されます。

例

すべて折りたたむ

行列-行列乗算

この例では、シンプルな行列-行列乗算を実行し、matrixMatrixKernel 設計パターンを使用して、CUDA^® コードを生成します。

1 つのファイル内に、2 つの行列入力 f1 および f2 を受け入れるエントリポイント関数 matMul_nn を記述します。MATLAB^® 関数 @times を使用して、f1 および f2 を要素ごとに乗算します。@ 記号は、関数 times のハンドルを作成します。gpucoder.matrixMatrixKernel() ステートメントを挿入します。入力行列は転置しないため、'nn' オプションを使用します。

function scores = matMul_nn(f1, f2)
    scores = gpucoder.matrixMatrixKernel(@times, f1, f2, 'nn');
end

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

codegen -config coder.gpuConfig('mex') ...
    -args {ones(1024,1024,'double'),ones(1024,1024,'double')} ...
    -report matMul_nn

生成された CUDA コードは 2 つのカーネルを含みます。出力行列 scores を初期化するための matMul_nn_kernel1 と、times 演算を実行する matMul_nn_kernel2 です。以下は、生成コードのスニペットです。

  matMul_nn_kernel1<<<dim3(2048U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_scores);
  cudaMemcpy(gpu_f2, f2, 8388608U, cudaMemcpyHostToDevice);
  cudaMemcpy(gpu_f1, f1, 8388608U, cudaMemcpyHostToDevice);
  matMul_nn_kernel2<<<dim3(16U, 16U, 1U), dim3(16U, 16U, 1U)>>>(gpu_f2, gpu_f1,
                   gpu_scores);
  cudaMemcpy(scores, gpu_scores, 8388608U, cudaMemcpyDeviceToHost);

matMul_nn_kernel2 には 2 次元ブロックから成る 2 次元グリッドがあります。このカーネルには、1 ブロックあたり 256 スレッドを含む 16 行 16 列のブロックが用意されています。