OpenACC ディレクティブによるプログラミング

15章 CUDA Unified Memory with OpenACC

CUDA Unified Memory について

 PGI 17.7リリースでは、allocatable として割り付けられたデータ用の CUDA Unified Memory の使用がベータ版からプロダクション版に移行しました。 この機能は、 OpenACC および CUDA Unified Memory PGInsider の記事で詳しく説明されており、Linux/x86-64 および Linux/OpenPOWER コンパイラで利用できます。 Linux/x86-64 では、デフォルトの PGI コードジェネレータと Beta版 LLVMベースのコードジェネレータの両方を使用してサポートされています。 この機能を有効にするには、オプション -ta=tesla:managed をコンパイラおよびリンカのコマンドラインに追加します。

 -ta=tesla:managed の場合、プログラム単位内のすべての C/C++、Fortran の明示的な allocate 文は、CUDA Unified Memory 上にデータを配置する同等の「管理された」データ割り当て呼び出しに置き換えられます。 管理対象データは CPU/GPU の単一アドレスを共有し、CPU と GPU メモリ間のデータ移動は CUDAドライバによって暗黙的に処理されます。 したがって、OpenACC のデータ句とディレクティブは、「管理された」データには必要ありません。 それらは本質的に無視され、実際には省略することができます。

 プログラムが管理されたメモリを割り当てると、デバイスのメモリだけでなくホストの固定メモリも割り当てられます。したがって、割り当てと解放の操作がやや高価になり、データの転送がやや高速になります。 メモリプールアロケータは、アロケートおよびフリーオペレーションのオーバーヘッドを軽減するために使用されます。プールアロケータは -ta=tesla:managed または -ta=tesla:pinned に対してデフォルトで有効になっています。 PGI 17.7リリースでは、 -Mcuda が存在するとプールアロケータが無効になります。 今後のリリースでその制限を解除する作業を進めています。

 管理データのデータ移動は、NVIDIA CUDA GPUドライバによって制御されます。 CPU または GPU 上でデータにアクセスするたびに、最後にアクセスされたときに同じデバイス上にデータが転送される可能性があります。 場合によっては、ページのスラッシングが発生し、パフォーマンスに影響を与えることがあります。 Parallel Forall では、CUDA Unified Memory の概要を紹介しています。

 この機能には次の制限があります。

  • 管理対象メモリの使用は、動的に割り当てられたデータにのみ適用されます。静的データ(C静的および外部変数、Fortran モ ジュール、共通ブロックおよび保存変数)および関数ローカルデータは、OpenACC ランタイムによって処理されるため、従来通り、ユーザによるデータ移動管理が必要です。 動的に割り当てられた Fortran ローカル変数と Fortran 割り当て可能配列は暗黙的に管理されますが、Fortran 配列ポインタは管理されません。
  • ローカル、グローバル、または静的データを指すメンバを持つ allocatable な集合型データを使用した際、 -ta=tesla:managed でコンパイルし、計算カーネルからそのポインタを介してメモリにアクセスしようとすると、実行時にエラーが発生します。
  • C++ 仮想関数はサポートされていません。
  • ファイルに OpenACC コードがない場合でも、変数が割り当てられたファイルをコンパイルするには、 -ta=tesla:managed コンパイラ・オプションを使用する必要があります。

この機能には、NVIDIA Kepler GPU で 使用するときに次の追加制限があります。

  • Kepler GPU 上のデータの動きは、高速の固定非同期データ転送によって実現されます。 しかし、プログラムの観点からは、転送は同期的です。
  • PGI runtime は、Keplar GPU を搭載したシステムで -ta=tesla:managed が使用されている場合、カーネルの同期実行を強制します。 このような状況は、余分な同期と CPU とGPU のオーバーラップの減少により、パフォーマンスが低下する可能性があります。
  • 管理されるメモリの総量は、Kepler GPU で利用可能なデバイスメモリの量に制限されています。

 この機能は、NVIDIA Fermi GPU ではサポートされていません。

CUDA Unified Memory Pool Allocator

 動的メモリ割り当ては、cudaMallocManage() を使用して行われます。このルーチンは、cudaMalloc()を使用して非ユニファイドメモリを割り当てるよりもオーバーヘッドが高いルーチンです。 cudaMallocManaged() の呼び出し回数が増えるほど、パフォーマンスへの影響が大きくなります。

 cudaMallocManaged() 呼び出しのオーバーヘッドを軽減するために、-ta=tesla:managed および -ta=tesla:pinned は CUDA Unified Memory プールアロケータを使用して、cudaMallocManaged() への呼び出し回数を最小限に抑えます。 プールアロケータは、デフォルトで有効になっています。次の環境変数を使用して無効にするか、動作を変更できます。

プールアロケータ環境変数

PGI_ACC_POOL_ALLOC プールアロケータを無効にします。 プールアロケータはデフォルトで有効になっています。 無効にするには、PGI_ACC_POOL_ALLOCを0に設定します。
PGI_ACC_POOL_SIZE プールのサイズを設定します。 デフォルトのサイズは1GBですが、他のサイズ(2GB、100MB、500KBなど)を使用できます。 実際のプール・サイズは、フィボナッチ・シリーズのサイズが、提供されたサイズまたはデフォルトサイズと比較して、最も近い、より小さい数になるように設定されます。 必要であれば、プールアロケータはプールを追加しますが、PGI_ACC_POOL_THRESHOLD値までしか追加しません。
PGI_ACC_POOL_ALLOC_MAXSIZE 割り当ての最大サイズを設定します。 割り当てのデフォルトの最大サイズは64Bですが、16B以上であれば別のサイズ(つまり100KB、10MB、250MBなど)を使用できます。
PGI_ACC_POOL_ALLOC_MINSIZE 割り当てブロックの最小サイズを設定します。 デフォルトサイズは16Bですが、他のサイズを使用することもできます。 サイズは16B以上でなければなりません。
PGI_ACC_POOL_THRESHOLD プールアロケータが占めることができるデバイスメモリの合計の割合を設定します。 デフォルトは50%に設定されていますが、他のパーセンテージを使用することもできます。

前章へ

次章へ

OpenACCプログラミングのインデックスへ

PGI アクセラレータコンパイラサイトへ