GPU メモリの割り当てと最小化

離散および管理モード

GPU Coder™ では、CUDA^® プログラミングモデルで利用可能な 2 つの異なるメモリ割り当て (malloc) モード cudaMalloc および cudaMallocManaged にアクセスできます。cudaMalloc API は従来の個別の CPU と GPU のグローバルメモリに適用できます。cudaMallocManaged は "ユニファイドメモリ" に適用できます。

プログラマの観点から見ると、従来のコンピューターアーキテクチャでは、データを CPU および GPU のメモリ空間に割り当てて、その間で共有する必要があります。アプリケーションでこれら 2 つのメモリ空間の間でのデータ転送を管理する必要があると複雑さが増します。ユニファイドメモリでは CPU と GPU 間で共有される管理メモリのプールが作成されます。この管理メモリは単一のポインターを介して CPU および GPU の両方からアクセスできます。ユニファイドメモリは、データを必要とするデバイスにそのデータを移動することでメモリパフォーマンスを最適化しようとします。同時に、その移動の詳細はプログラムから見えないようにします。ユニファイドメモリによってプログラミングモデルは簡略化されますが、GPU に書き込まれたデータに CPU からアクセスするときにデバイス同期呼び出しが必要になります。GPU Coder によりこれらの同期呼び出しが挿入されます。NVIDIA^® によれば、ユニファイドメモリは、NVIDIA Tegra^® のような組み込みハードウェアをターゲットにしている場合に、大きなパフォーマンス上のメリットをもたらすことができます。

GPU Coder アプリでメモリ割り当てモードを変更するには、[詳細設定]、[GPU Coder] の下にある [Malloc モード] ドロップダウンボックスを使用します。コマンドラインインターフェイスを使用する場合、MallocMode ビルド構成プロパティを使用し、それを 'discrete' または 'unified' に設定します。

メモ

ホスト開発コンピューター上の NVIDIA GPU デバイスをターゲットとする場合のユニファイドメモリ割り当て (cudaMallocManaged) モードは、将来のリリースで削除される予定です。NVIDIA 組み込みプラットフォームをターゲットとする場合は、引き続きユニファイドメモリ割り当てモードを使用できます。

GPU メモリマネージャー

効率的なメモリ割り当て、メモリ管理、および実行時のパフォーマンス向上のために、GPU メモリマネージャーを使用できます。GPU メモリマネージャーは、大きな GPU メモリプールのコレクションを作成し、これらのプール内のメモリブロックのチャンクの割り当ておよび割り当て解除を管理します。大きなメモリプールを作成することにより、メモリマネージャは CUDA のメモリ API に対する呼び出し回数を減らして、実行時のパフォーマンスを向上させます。GPU メモリマネージャーは、MEX およびスタンドアロンの CUDA コード生成に使用できます。

GPU メモリマネージャーを有効にするには、次のいずれかの方法を使用します。

GPU コード構成オブジェクト (coder.gpuConfig) で、EnableMemoryManager プロパティを true に設定します。
GPU Coder アプリの [GPU コード] タブで、[GPU メモリマネージャー] を選択する。
Simulink^® コンフィギュレーションパラメーターダイアログボックスの [コード生成]、[GPU コード] ペインで、[メモリマネージャー] パラメーターを選択する。

cuFFT、cuBLAS、cuSOLVER などの NVIDIA CUDA ライブラリを使用する CUDA コードでは、効率的なメモリ割り当てと管理のために GPU メモリマネージャーの使用を有効にできます。

CUDA ライブラリでメモリプールを使用するには、上記の方法のいずれかを使用してメモリマネージャーを有効にし、次のようにします。

GPU コード構成オブジェクト (coder.gpuConfig) で、EnableCUFFT、EnableCUBLAS、または EnableCUSOLVER のプロパティを有効にする。
GPU Coder アプリの [GPU コード] タブで、[cuFFT の有効化]、[cuBLAS の有効化]、または [cuSOLVER の有効化] を選択する。
Simulink コンフィギュレーションパラメーターダイアログボックスの [コード生成]、[GPU コード] ペインで、[cuFFT]、[cuBLAS]、または [cuSOLVER] のパラメーターを選択する。

メモリ最小化

GPU Coder は CPU および GPU 区画間のデータの依存関係を解析し、生成コード内の関数 cudaMemcpy の呼び出しの数を最小化するように最適化を実行します。解析では、cudaMemcpy を使用して CPU と GPU 間でデータをコピーしなければならない位置の最小セットも特定されます。

たとえば、関数 foo には、データを CPU で逐次的に処理するコードのセクションと、GPU で並列に処理するコードのセクションがあります。

function [out] = foo(input1,input2)
	   …
     % CPU work
			input1 = …
			input2 = …
			tmp1 = …
			tmp2 = …
   	…
     % GPU work
			kernel1(gpuInput1, gpuTmp1);
       kernel2(gpuInput2, gpuTmp1, gpuTmp2);
       kernel3(gpuTmp1, gpuTmp2, gpuOut);

   	…
     % CPU work
       … = out

end

最適化されていない CUDA 実装には、入力 gpuInput1,gpuInput2 と一時的な結果 gpuTmp1,gpuTmp2 をカーネル呼び出し間で転送するための関数 cudaMemcpy の呼び出しが複数ある可能性があります。中間結果 gpuTmp1,gpuTmp2 は GPU 以外では使用されないため、これらは GPU メモリ内部に保存でき、その結果として関数 cudaMemcpy の呼び出しは少なくなります。これらの最適化により、生成コードの全体的なパフォーマンスが改善します。最適化された実装は以下のようになります。

gpuInput1 = input1;
gpuInput2 = input2;

kernel1<<< >>>(gpuInput1, gpuTmp1);
kernel2<<< >>>(gpuInput2, gpuTmp1, gpuTmp2);
kernel3<<< >>>(gpuTmp1, gpuTmp2, gpuOut);

out = gpuOut;

冗長な cudaMemcpy の呼び出しを削除するために、GPU Coder は与えられた変数のすべての用途と定義を解析し、ステータスフラグを使用して最小化を実行します。元のコードと生成コードの例を次の表に示します。

元のコード	最適化された生成コード
A(:) = … … for i = 1:N gB = kernel1(gA); gA = kernel2(gB); if (somecondition) gC = kernel3(gA, gB); end … end … … = C;	A(:) = … A_isDirtyOnCpu = true; … for i = 1:N if (A_isDirtyOnCpu) gA = A; A_isDirtyOnCpu = false; end gB = kernel1(gA); gA = kernel2(gB); if (somecondition) gC = kernel3(gA, gB); C_isDirtyOnGpu = true; end … end … if (C_isDirtyOnGpu) C = gC; C_isDirtyOnGpu = false; end … = C;

元のコード

最適化された生成コード

A(:) = …
…
for i = 1:N
   gB = kernel1(gA);
   gA = kernel2(gB);

   if (somecondition)
      gC = kernel3(gA, gB);
   end
   …
end
…
… = C;

A(:) = …
A_isDirtyOnCpu = true;
…
for i = 1:N
   if (A_isDirtyOnCpu)
      gA = A;
      A_isDirtyOnCpu = false;
   end
   gB = kernel1(gA);
   gA = kernel2(gB);
   if (somecondition)
      gC = kernel3(gA, gB);
      C_isDirtyOnGpu = true;
   end
   …
end
…
if (C_isDirtyOnGpu)
   C = gC;
   C_isDirtyOnGpu = false;
end
… = C;

_isDirtyOnCpu フラグは、ルーチンに関する GPU Coder のメモリ最適化に対し、与えられた変数が CPU と GPU のどちらで宣言および使用されているかを示します。

CUDA MEX 用の共有メモリマネージャー

特に、CUDA MEX コードを生成すると、実行中のすべての CUDA MEX 関数のメモリ管理を処理する単一のユニバーサルメモリマネージャーが GPU Coder によって作成されるため、MEX 関数のパフォーマンスがさらに向上します。共有 MEX メモリマネージャーのプロパティを表示し、割り当てを管理するには、関数 cudaMemoryManager を使用して gpucoder.MemoryManager オブジェクトを作成します。使用されていない GPU メモリを解放するには、関数 freeUnusedMemory を呼び出します。How Shared GPU Memory Manager Improves Performance of Generated MEXを参照してください。

参考