Main Content

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

gpucoder.matrixMatrixKernel

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

説明

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 は、同じサイズか、互換性のあるサイズでなければなりません。たとえば、AMK 列の行列、BKN 列の行列の場合、CMN 列の行列になります。

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 列のブロックが用意されています。

R2017b で導入