Main Content

cuSOLVER の例

この例では、cuSOLVER ライブラリを使用して、x に対する線形方程式系 Ax = B を解きます。行列 A および B の行数は同じでなければなりません。A がスカラーの場合、A\B は A.\B と等価です。A が n 行 n 列の正方行列で、B が n 行の行列である場合、x = A\B は方程式 A*x = B の解になります (存在する場合)。backslash の MATLAB® 実装は次のようになります。

function [x] = backslash(A,b)
if (isscalar(A))
    x = coder.nullcopy(zeros(size(b)));
else
    x = coder.nullcopy(zeros(size(A,2),size(b,2)));
end

x = A\b;

end

カーネル作成用の backslash の準備

GPU Coder™ は、ライブラリの呼び出しを生成するために特殊なプラグマを必要としません。前と同じように、CUDA® カーネルを生成する方法は 2 つあります。coder.gpu.kernelfun および coder.gpu.kernel です。この例では、coder.gpu.kernelfun プラグマを利用して CUDA カーネルを生成します。変更した関数 backslash は次のようになります。

function [x] = backslash(A,b) %#codegen

if (isscalar(A))
    x = coder.nullcopy(zeros(size(b)));
else
    x = coder.nullcopy(zeros(size(A,2),size(b,2)));
end

coder.gpu.kernelfun()
x = A\b;

end

メモ

算術演算子および関数を cuSOLVER ライブラリ実装に置き換える場合、入力データには最小サイズが必要です。最小しきい値は 128 要素です。

生成された CUDA コード

CUDA コードを生成すると、GPU Coder は cuSOLVER ライブラリを初期化する関数呼び出しを作成し、mldivide 演算を実行して、cuSOLVER ライブラリで使用されたハードウェア リソースを解放します。生成された CUDA コードのスニペットは次のとおりです。

  cusolverEnsureInitialization();

  /*    Copyright 2017 The MathWorks, Inc. */
  cudaMemcpy(b_gpu_A, A, 1152UL, cudaMemcpyHostToDevice);  
  blackslash_kernel1<<<dim3(1U, 1U, 1U), dim3(160U, 1U, 1U)>>>(b_gpu_A,gpu_A);
  cudaMemcpy(b_A, gpu_A, 1152UL, cudaMemcpyDeviceToHost);
  cusolverDnDgetrf_bufferSize(cusolverGlobalHandle, 12, 12, &gpu_A[0], 12,
    &cusolverWorkspaceReq);
  cusolverWorkspaceTypeSize = 8;
  cusolverInitWorkspace();
  cudaMemcpy(gpu_A, b_A, 1152UL, cudaMemcpyHostToDevice);
  cusolverDnDgetrf(cusolverGlobalHandle, 12, 12, &gpu_A[0], 12, (real_T *)
                   cusolverWorkspaceBuff, &gpu_ipiv_t[0], gpu_info_t);
  A_dirtyOnGpu = true;
  cudaMemcpy(&info_t, gpu_info_t, 4UL, cudaMemcpyDeviceToHost);

cuSOLVER ライブラリを初期化し、cuSOLVER ライブラリ コンテキストのハンドルを作成するために、関数 cusolversEnsureInitialization() で cusolverDnCreate() cuSOLVER API を呼び出します。これにより、ホストとデバイスのハードウェア リソースが割り当てられます。

static void cusolverEnsureInitialization(void)
{
  if (cusolverGlobalHandle == NULL) {
    cusolverDnCreate(&cuSolverGlobalHandle);
  }
}

backslash_kernel1 は行列 A をゼロ パディングします。このカーネルが起動し、512 スレッドから成る 1 つのブロックが用意されます。 

static __global__ __launch_bounds__(160, 1) void backslash_kernel1(const real_T *
  A, real_T *b_A)
{
  int32_T threadId;
  ;
  ;
  threadId = (int32_T)(((gridDim.x * gridDim.y * blockIdx.z + gridDim.x *
    blockIdx.y) + blockIdx.x) * (blockDim.x * blockDim.y * blockDim.z) +
                       (int32_T)((threadIdx.z * blockDim.x * blockDim.y +
    threadIdx.y * blockDim.x) + threadIdx.x));
  if (!(threadId >= 144)) {
    /*    Copyright 2017 The MathWorks, Inc. */
    b_A[threadId] = A[threadId];
  }
}

cudaMemcpy の呼び出しは、行列 A をホストからデバイスに転送します。関数 cusolverDnDgetrf は m 行 n 列の行列の LU 分解を計算します。

P*A = L*U

ここで、A は m 行 n 列の行列、P は置換行列、L は単位対角要素を持つ下三角行列、U は上三角行列です。

cuSOLVER スタンドアロン コード

cuSOLVER の部分的なサポートしかない qr などの関数の場合、GPU Coder は必要に応じて LAPACK ライブラリを使用します。MEX 関数の場合、コード ジェネレーターは MATLAB に含まれている LAPACK ライブラリを使用します。スタンドアロン コードの場合、コード ジェネレーターは指定された LAPACK ライブラリを使用します。LAPACK ライブラリを指定するには、以下を行います。

  • コマンド ラインで、LAPACK ライブラリ情報を含む独自の coder.LAPACKCallback クラスを定義し、それをコード構成オブジェクトの CustomLAPACKCallback プロパティに割り当てます。

  • GPU Coder アプリで、カスタム LAPACK ライブラリのコールバックを LAPACK ライブラリに設定します。

たとえば、スタンドアロン実行可能ファイルを生成するには、次のコード生成スクリプトを使用できます。ここで、myLAPACK は LAPACK ライブラリ情報を含むカスタム coder.LAPACKCallback クラスの名前です。

cfg = coder.gpuConfig('exe');
cfg.CustomLAPACKCallback = 'myLAPACK';
cfg.GenerateExampleMain = 'GenerateCodeAndCompile';

classdef myLAPACK < coder.LAPACKCallback
    methods (Static)
        function hn = getHeaderFilename()
            hn = 'lapacke.h';
        end
        function updateBuildInfo(buildInfo, buildctx)
            [~,linkLibExt] = buildctx.getStdLibInfo();
            cudaPath = getenv('CUDA_PATH'); 
            libPath = 'lib\x64';            
            
            buildInfo.addIncludePaths(fullfile(cudaPath,'include'));
            libName = 'cusolver';
            libPath = fullfile(cudaPath,libPath);
            buildInfo.addLinkObjects([libName linkLibExt], libPath, ...
                '', true, true);
            
            lapackLocation = 'C:\LAPACK\win64'; % specify path to LAPACK libraries 
            
            includePath = fullfile(lapackLocation,'include');
            buildInfo.addIncludePaths(includePath);
            libPath = fullfile(lapackLocation,'lib');
            libName = 'mllapack';
            
            buildInfo.addLinkObjects([libName linkLibExt], libPath, ...
                '', true, true);
            buildInfo.addDefines('HAVE_LAPACK_CONFIG_H');
            buildInfo.addDefines('LAPACK_COMPLEX_STRUCTURE');
        end
    end
end
詳細については、生成されたスタンドアロン コードでの LAPACK 呼び出しを使用した線形代数の高速化を参照してください。