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

　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() への呼び出し回数を最小限に抑えます。 プールアロケータは、デフォルトで有効になっています。次の環境変数を使用して無効にするか、動作を変更できます。