Main Content

生成コードからのカスタム CUDA カーネルの呼び出し

MATLAB® コード内から、カスタム コードまたはレガシ コードとも呼ばれる外部 CUDA® カーネルを直接呼び出すことができます。CUDA カーネルを呼び出すには、coder.ceval を使用します。コード ジェネレーターは、CUDA カーネルを MATLAB から生成された CUDA コードに統合します。生成されたコードで使用する外部ライブラリ、最適化されたコード、または CUDA を使用して開発されたオブジェクト ファイルがある場合は、コードを統合します。

外部 CUDA カーネルは、__global__ 修飾子を使用して、GPU デバイスで関数 (カーネル) を実行し、ホストまたはデバイスから関数を呼び出さなければなりません。__device__ 修飾子をもつ関数は、デバイス関数と呼ばれます。デバイス関数は、他のデバイス関数またはグローバル関数からのみ呼び出すことができるという点で、グローバル関数とは異なります。カスタム デバイス関数の統合の詳細については、生成コードからのカスタム CUDA デバイス関数の呼び出しを参照してください。

メモ

coder.ceval は、コード生成用の MATLAB コードでのみ使用します。coder.ceval は、コンパイルされていない MATLAB コードでエラーを生成します。MATLAB 関数が MATLAB で実行されているかどうかを判定するには、coder.target を使用します。関数が MATLAB で実行されている場合は、CUDA カーネルの MATLAB バージョンを呼び出します。

カスタム CUDA カーネルの呼び出し

ライブ スクリプトを開く

この例では、coder.cevalを使用して単純な CUDA カーネルを MATLAB コードと統合する方法を示します。次の MATLAB 関数 saxpy について考えます。

type saxpy.m
function y = saxpy(a,x,y)
    y = a*x + y;
end

この例では、外部 CUDA カーネルを使用して x + y の演算を実装することにします。ファイル saxpy.cu に実装されている次の CUDA カーネル saxpy_kernel について考えます。

type saxpy.cu
#include "saxpy.h"

__global__
void saxpy_kernel(uint32_T n, real32_T a, real32_T *x, real32_T *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

saxpy_kernel を MATLAB コードに統合するには、関数プロトタイプを含むヘッダー ファイルが必要です。ファイル saxpy.h を確認します。

type saxpy.h
#ifndef real32_T 
#define real32_T float
#define uint32_T unsigned int
#endif

#define saxpy(grid,block,n,a,x,y) saxpy_kernel<<<grid,block>>>(n,a,x,y)

__global__ void saxpy_kernel(uint32_T n, real32_T a, real32_T *x, real32_T *y);

この例では、CUDA MEX を生成します。uint32_T と real32_T は、生成された MEX コードで使用されるカスタム型です。コード ジェネレーターは、MATLAB コードで使用するデータ型に対応するデータ型を CUDA コードで生成します。生成されるデータ型は、ターゲット プラットフォームとコンパイラによって異なります。コード ジェネレーターは、組み込みの CUDA/C++ データ型 (short、long、int など) と、typedef ステートメントを使用して定義されたカスタム データ型のいずれも生成できます。既定では、コード ジェネレーターは、スタンドアロン コード (lib、dll、または exe) では組み込み型を、MEX コードではカスタム型を生成します。詳細については、生成コードの型への MATLAB 型のマッピングを参照してください。

エントリポイント関数

coder.ceval コマンドを使用して、エントリポイント関数 saxpyRef.m で CUDA カーネルを呼び出します。coder.ceval は、コード生成用の MATLAB コードでのみ使用します。coder.rref および coder.ref コマンドは、ポインターをコピーするのではなく、配列に渡すようにコード ジェネレーターに指示します。

type saxpyRef.m
function y = saxpyRef(a,x,y)
%   saxpyRef Entry-point function for computing single-precision 
%   (A*X) Plus Y

%   Copyright 2022 The MathWorks, Inc.
coder.gpu.kernelfun;

if coder.target('MATLAB')
    y = a*x + y;
else
    coder.ceval('saxpy', uint32(floor((numel(x)+255)/256)), ...
        uint32(256),uint32(numel(x)), single(a), ...
        coder.rref(x,'gpu'),coder.ref(y,'gpu'));
end
end

CUDA コードの生成

CUDA コードを生成するには、GPU コード構成オブジェクトを作成します。構成オブジェクトのカスタム コード プロパティを設定して、カスタム CUDA ファイルの場所を指定します。詳細については、外部 C/C++ コードのビルドの構成を参照してください。

cfg = coder.gpuConfig("mex");
cfg.GenerateReport = true;
cfg.CustomSource = "saxpy.cu";
cfg.CustomInclude = pwd;
cfg.CustomSourceCode = '#include "saxpy.h"';

aType = coder.newtype('single', [1 1], [0 0]);
xType = coder.newtype('single', [4096 256], [0 0]);
yType = coder.newtype('single', [4096 256], [0 0]);
inputArgs = {aType,xType,yType};

codegen -config cfg saxpyRef -args inputArgs
Code generation successful: View report

生成コード

生成された CUDA コードと元の MATLAB コードを比較するには、work\codegen\mex\saxpyRef フォルダーにある CUDA ファイル saxyRef.cu を開きます。

#include "saxpy.h"
// Function Definitions
void saxpyRef(real32_T a, const real32_T x[1048576], real32_T y[1048576])
{
  real32_T(*gpu_x)[1048576];
  real32_T(*gpu_y)[1048576];
  cudaMalloc(&gpu_y, 4194304UL);
  cudaMalloc(&gpu_x, 4194304UL);
  //    saxpyRef Entry-point function for computing single-precision (A*X) Plus
  //    Y
  //    Copyright 2022 The MathWorks, Inc.
  cudaMemcpy(*gpu_x, x, 4194304UL, cudaMemcpyHostToDevice);
  cudaMemcpy(*gpu_y, y, 4194304UL, cudaMemcpyHostToDevice);
  saxpy(4096U, 256U, 1048576U, a, &(*gpu_x)[0], &(*gpu_y)[0]);
  cudaMemcpy(y, *gpu_y, 4194304UL, cudaMemcpyDeviceToHost);
  cudaFree(*gpu_x);
  cudaFree(*gpu_y);
}

生成された MEX の実行

生成された MEX をランダムな入力で実行し、その結果と MATLAB シミュレーションを比較します。

a = single(15);
x = randi(10,4096,256,'single');
y = zeros(4096,256,'single');

yMATLAB = saxpyRef(a,x,y);
yGPU = saxpyRef_mex(a,x,y);

if (yGPU - yMATLAB == 0)
    fprintf('\nMATLAB simulation matches GPU execution.\n');
end
MATLAB simulation matches GPU execution.

参考

関数

オブジェクト

関連するトピック