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^® によれば、ユニファイドメモリは、CUDA 8.0 を使用している場合や、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) で、MemoryManager プロパティを有効にする。
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 メモリプールのカスタマイズオプション

GPU メモリマネージャーでは、GPU メモリプール内のメモリブロックの割り当ておよび割り当て解除を管理するために、表にリストされている追加のコード構成パラメーターを利用できます。

コード構成パラメーター	説明	値
GPU コード構成オブジェクト (`coder.gpuConfig`): `BlockAlignment` GPU Coder アプリ: [GPU コード] タブの [ブロックの配置]	ブロックの配置を制御します。プール内のブロックサイズ (バイト) は、指定した値の倍数になります。	2 のべき乗である正の整数。既定値は 256 です。
GPU コード構成オブジェクト: `FreeMode` GPU Coder アプリ: [GPU コード] タブの [解放モード]	メモリマネージャーが GPU デバイスのメモリを解放するタイミングを制御します。 `'Never'` に設定すると、メモリマネージャーが破棄されたときにのみメモリを解放します。生成コードで関数 `terminate` が呼び出されたときに、空の GPU プールを解放するには、`'AtTerminate'` を使用します。MEX ターゲットの場合、生成された MEX 関数が呼び出されるたびにメモリを解放します。他のターゲットの場合、関数 terminate が呼び出されたときにメモリを解放します。 `'AfterAllocate'` に設定すると、CUDA のメモリ割り当てが呼び出されるたびに空のプールを解放します。	`'Never'` (既定値) \| `'AtTerminate'` \| `'AfterAllocate'`
GPU コード構成オブジェクト: `MinPoolSize` GPU Coder アプリ: [GPU コード] タブの [最小プールサイズ]	最小プールサイズをメガバイト (MB) 単位で指定します。	2 のべき乗である正の整数。既定値は 8 です。
GPU コード構成オブジェクト: `MaxPoolSize` GPU Coder アプリ: [GPU コード] タブの [最大プールサイズ]	最大プールサイズをメガバイト (MB) 単位で指定します。メモリマネージャーは、`MinPoolSize` パラメーターと `MaxPoolSize` パラメーターを使用して、2 つの値の間を 2 のべき乗で内挿することにより、サイズレベルを計算します。たとえば、`MinPoolSize` が 4 で、`MaxPoolSize` が 1024 の場合、サイズレベルは {4, 8, 16, 32, 64, 128, 256, 512, 1024} になります。	2 のべき乗である正の整数。既定値は 2048 です。

メモリ最小化

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 のどちらで宣言および使用されているかを示します。

参考