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' に設定します。

メモリ最小化

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