PGIアクセラレータ機能の概要

キーワード PGI アクセラレータコンパイラ　GPU用コンパイラ

QUESTION PGIアクセラレータ™ x64+GPU Fortran/C99 コンパイラサポート：　NVIDIA の GPU搭載の Linux / Windows / OS X システム上で、コンパイラ・ディレクティブ（指示行）ベースでプログラミング可能なアクセラレータ・コンパイラを業界で初めてリリースしました。PGF95並びにPGCC(R) アクセラレータ™ コンパイラは、CUDAソフトウェア環境を有するNVIDIA GPUが搭載されている、全てのインテル(R)並びにAMDの64ビットプロセッサベースのシステム上でサポートされます。この機能を使用することによって、Fortran/C99プログラムをマルチコアベースから x64+GPU ベースのプラットフォームへの移行が容易となり、さらなるチューニングを促します。　
2009年12月11日 Copyright © 株式会社ソフテック

PGIアクセラレータ™ コンパイラ機能 (NVIDIA GPUサポート）

「アクセラレータ」とは、特別の目的で CPU にアタッチして使用する協調プロセッサであり、時間の掛かる計算部分を CPU の演算機構からデータと実行部分のカーネルをオフロードするために使用されます。ここでは、新しい PGI アクセラレータ™ コンパイラについて説明します。この内容には、ホスト CPU からアタッチされているアクセラレータへオフロード可能な Fortran、C におけるコードの領域 (region) を指定するために使われるディレクティブに関する事項も含みます。 　PGIアクセラレータ、プログラミングモデル、ディレクティブに関する詳細な情報は、PGI User's Guide における7章「Using an Accelerator」と18章「PGI Accelerator Compilers Reference」をご参照下さい。

▶ コンポーネント

　PGIアクセラレータ™ コンパイラのテクノロジーは、次のコンポーネントを含みます。なお、アクセラレータ対応のデバッガあるいはプロファイラは、このリリースでは含まれておりません。

  • PGF95アクセラレータ付きFortran 90/95 自動並列化コンパイラ
  • PGCCアクセラレータ付き ANSI C99 and K&R C 自動並列化コンパイラ (PGC++には本機能無し）
  • NVIDIA CUDAツールキット・コンポーネント
  • システムが適切なGPUあるいは、アクセラレータカードを有するかどうかを検出するための単純なコマンドライン・ツール

▶ 可用性

　PGI 2010 (PGI 10.0) Fortran & Cアクセラレータ・コンパイラは、x86プロセッサベースのシステム上のワークステーションやサーバでのみ使用できます。もちろん、このシステム上にはNVIDIA CUDA-enable GPUあるいは、TESLAカードが実装されている必要があります。これらのコンパイラはPGIがサポートする、OS X 64bitを除く全てのプラットフォームで利用できます。

▶ ユーザ指示のアクセラレータ・プログラミング

　ユーザが指示するタイプのアクセラレータ・プログラミングにおいては、ユーザはアクセラレータ・デバイスへオフロードするために、その対象とすべきホスト・プログラムの領域を指定します。また、ターゲットとなるアクセラレータでサポートされていないコンストラクトを含む領域やユーザ・プログラムの大半は、ホスト上で実行されます。ここでの説明は、アクセラレータにオフロードされるコードの領域やループの指定方法に関してのみ、説明しています。

▶ カバーしていない、あるいは実装していない機能

　現在、PGIアクセラレータ・コンパイラは、コンパイラや他のツール等による、アクセラレータにオフロード化できるコード領域の自動検出やそのオフロード化のための機能は含んでおりません。PGIコンパイラの将来のバージョンでは、自動的なオフロード化や同一ホスト上のマルチ・アクセラレータの使用、あるいは、異なるタイプのマルチ・アクセラレータの使用もできるように計画していますが、現在のバージョンでは、これらはサポートされておりません。

▶ 必要とされるシステム仕様

　PGI アクセラレータの機能を使用するためには、システム上でNVIDIAドライバーコンポーネントを必ずインストールしておく必要があります。CUDA ToolkitとCUDA SDKのインストールは必須ではありませんが、PGIとは別に、完全な CUDA環境を構築しておきたい場合は、以下の全てのソフトウェア・コンポーネントをインストールすることを推奨します。 これらは、NVIDIA の Webサイト( www.nvidia.com/cuda/) からダウンロード可能です。これらは、PGIコンポーネントではありませんので、NVIDIA 社によってライセンス並びにサポートされます。

  • NVIDIA Driver
  • CUDA Toolkit
  • CUDA SDK

▶ サポートするプロセッサとGPU

　PGI アクセラレータ・コンパイラのリリースは、PGI コンパイラ＆ツールの2010リリース以降がサポートする全てのAMD64並びにIntel 64のホストプロセッサ上で利用できます。また、コンパイラのオプション・フラグである -tp <target>フラグを使用することもできます。
　NVIDIA GPU を対象としたアクセラレータのディレクティブの認識を有効にするための -ta=nvidia フラグも使用できます。そして、CUDA-enable な GeForce、Quadro、Tesla カードを有した CUDA がインストールされた任意のシステム上で、生成されたコードを使用することができます。 アクセラレータのテクノロジーに関するこれらのフラグの詳細な情報は、PGI User's Guide をご参照下さい。また、サポートされるGPUの詳細なリストは、NVIDIA のウェブサイトでご覧下さい。www.nvidia.com/object/cuda_learn_products.html

▶ インストールとライセンシング

　PGI アクセラレータ・コンパイラは、標準的な PGI Workstation、Server、CDKのライセンスキーの他に、別のライセンスキーが必要となります。 PGI アクセラレータ・コンパイラは、別のコンパイラ・エディション製品として発売しますので、アクセラレータ機能を使用したいお客様は、新規のご購入、あるいは、現製品のアップグレードを行っていただく必要があります。

▶ アクセラレータ・プログラムの実行

　アクセラレータ用のディレクティブを有し、-ta=nvidia フラグを使ってコンパイル・リンクしたプログラムを実行させる方法は、-ta=nvidia フラグを付けないでコンパイル・リンクした場合と同じである。プログラムは CUDA ライブラリを探し、動的にロードします。もし、ライブラリがない場合、あるいは、プログラムをコンパイルした際に存在した場所とは異なるディレクトリ上にある場合は、LD_LIBRARY_PATH 環境変数に CUDA ライブラリの所在を追加する必要があります。
　プログラムが最初のアクセラレータ領域に到達した際、約 0.5～1.5秒ほどのポーズ時間が存在します。これは、GPU への操作権を得て、静的なリソースを割り付けるための時間オーバーヘッドです。この後は、引き続きオーバーヘッドなしに、アクセラレータ領域を実行することができます。
　アクセラレータ・プログラムを CUDA-enable の NVIDIA GPU を有しないシステムで動作させた場合、あるいは、ランタイム・ライブラリの検索範囲のディレクトリの場所に CUDA ライブラリが存在しない場合、プログラムは、その実行時にエラーメッセージ無しに終了してしまします。 環境変数 ACC_NOTIFY にゼロ以外の整数をセットした場合、ランタイム・ライブラリは GPU のカーネル(kernel) が開始される度に、その履歴を標準エラーにプリントします。

▶ PGIアクセラレータ™ コンパイラのランタイム・ライブラリ

　PGI アクセラレータ・コンパイラは、アクセラレータの機能に関する問い合わせや実行時にアクセラレータ用のプログラムの挙動を制御するような、ユーザによる呼び出し可能な関数やライブラリ・ルーチンを提供します。Fortranにおいて、PGI アクセラレータ・コンパイラのランタイム・ルーチンは、PUREあるいはELEMENTAL手続きからコールされません。 アクセラレータ・ライブラリにアクセスするためには、-ta フラグ（コンパイル・オプション）と共にアクセラレータ・プログラムをリンクする必要があります。 C と Fortran では、各々、別のランタイム・ライブラリが存在します。

  • C ランタイム・ライブラリ・ファイル

C では、ランタイム・ライブラリ・ルーチンのプロトタイプが "accel.h"と言う名前のヘッダーファイルで用意されています。全てのランタイム・ライブラリ・ルーチンは、"C"リンケージを有する外部関数です。このファイルは、以下を定義します。

  1. このセクションの述べられている全てのルーチンのプロトタイプ
  2. アクセラレータの型を記述するenumulation型を含むこれらのプロトタイプで使用されている任意のデータ型
  • Fortranランタイム・ライブラリ・ファイル　

Fortranでは、interfaceの宣言は、accel_lib.h と言う名前の Fortran インクルード・ファイルの中、あるいは、accel_lib と言う Fortran モジュールの中で提供されます。これらのファイルでは、以下を定義します。

  1. このセクションの述べられている全てのルーチンのプロトタイプ
  2. これらのルーチンへの引数のための整数kind値を定義するための整数パラメータ
  3. アクセラレータの型を記述するための整数パラメータ

　yyyymm と言う値を持った整数パラメータ accel_version は、サポートしているアクセラレータ・プログラミングモデルのバージョンを示すもので、yyyy は年、mm は月を意味します。この値は、プリプロセッサ変数 _ACCEL の値に合致します。
　次のリストは、PGIが現時点でサポートしているPGIアクセラレータ・コンパイラのランタイム・ライブラリを簡単に述べたものです。これらのルーチンの完全な説明は、PGI User's GuideのPGI Accelerator Runtime Routinesの章をご覧下さい。

  • acc_get_device ： 使用中のアクセラレータ・デバイスのタイプを返す。
  • acc_get_num_devices ：ホスト側にアタッチしているアクセラレータ・デバイスの数を返す。
  • acc_init ：アクセラレータ・デバイスに接続し、初期化を行い、アクセラレータ・ライブラリの中の制御構造を割り付ける。
  • acc_on_device ：特定のデバイス上で実行しているかどうかをプログラムに伝える。
  • acc_set_device ：アクセラレータ計算領域を実行するときに使用するデバイスのタイプをランタイムへ伝える。
  • acc_set_device_num ：アクセラレータの領域を実行するときに使用されるデバイスに伝えるルーチン。
  • acc_shutdown ：アクセラレータ・デバイスへの接続を切るようにランタイムに伝える。そして、任意のランタイムリソースを自由にする。

▶ 環境変数

　PGIは、アクセラレータ領域の挙動を修正するための環境変数を提供します。 このセクションでは、実行時のアクセラレータを使用するプログラムの挙動を制御するために使用する、ユーザが指定可能な環境変数について説明します。これらの環境変数は、以下のルールに準拠する必要があります。

  • 環境変数の名前は、大文字でなければなりません。
  • 環境変数にセットする値は、大文字・小文字の区別はありません。また、最初と最後の余白指定も可能です。
  • もし、プログラムが開始した後で、環境変数の値が変更された場合、その挙動は実装依存です。また、例えプログラム自身が値を変えたとしてもこれは同様です。

　次のリストは、PGIが現時点でサポートしているPGIアクセラレータ環境変数を簡単に述べたものです。これらのルーチンの完全な説明は、PGI User's GuideのPGI Accelerator Runtime Routinesの章をご覧下さい。

  • ACC_DEVICE：アクセラレータのためのPGI Unified Binaryを実行する時に使用するアクセラレータのデフォルト・デバイスを制御する。NVIDIA GPU上で実装する場合は、この環境変数は NVIDIAにセットしなければならない。それ以外の場合は、これは HOSTとなる。
  • ACC_DEVICE_NUM：アクセラレータ領域を実行する際に使用するデフォルトのデバイス番号を制御する。この値は、負ではない整数で0からホストにアタッチしているデバイスの数までの値となる。
  • ACC_NOTIFY：負ではない整数値をセットしたとき、デバイス上のカーネル(kernel) が開始される度に、その事実を標準エラーにプリントします。

▶ コンパイラ・コマンド・オプション

　アクセラレータを動作させるために特に使用するコマンド・オプションは以下の通りです。

  • -tp　：　対象となるホストプロセッサのアーキテクチャを指定するオプション
  • -Minfo or -Minfo=accel　：　このオプションを指定すると、コンパイラがアクセラレータ領域をGPUカーネルに翻訳できたかどうかについて、コンパイラのメッセージとして出力します。
  • -ta=nvidia(,nvidia_suboptions), host　： PGI アクセラレータ・コンパイラに伴う新しいオプションです。-ta は、ターゲット・アーキテクチャを意味します。Fortranにおける !$ACC ディレクティブ、C における #pragma acc ディレクティブをコンパイラに認識させるために、このオプションを使用します。これは、Fortran 並びに C コンパイラのみで有効です。このオプションは、以下のサブオプションを有します。
    * nvidia - NVIDIAアクセラレータをターゲットとして選択します。さらに、以下の nvidia 用のサブオプションがあります。
  • analysis ：ループの解析のみ行い、コードの生成を行いません。
  • cc10 ：compute capability 1.0 のコードを生成
  • cc11 ：compute capability 1.1 のコードを生成
  • cc13 ：compute capability 1.3 のコードを生成
  • fastmath ：fast mathライブラリを使用
  • keepgpu：kernelソースファイルを保持する
  • keepptx ：GPUコードのためのportable assembly(.ptx)ファイルを保持する
  • maxregcount:n ：GPU上で使用するレジスタの最大数を指定。ブランクの場合は、 制約が無いと解釈する
  • mul24 ：添字計算のために 24ビット乗算を使用
  • nofma ：fused-multiply-add命令を生成しない
  • time：アクセラレータ領域の単純な時間情報を集積するためにプロファイル・ライブラリをリンクする
  • * host - ターゲットとして、hostを選択する。nvidiaオプションとの組み合わせで使用されます。アクセラレータ領域をホスト側で実行するようにコンパイルする。このオプションは、PGI Unified Binaryコードを生成します。

　コンパイラは、自動的に必要とする CUDA ソフトウェアのツールを発動し、GPU カーネルコードを生成し、Linux オブジェクト・ファイルの中にカーネルを埋め込みます。
　コマンドライン上でリンク時に、アクセラレータのライブラリにアクセスするためには、必ず、コンパイル時に指定したものと同じ、-taフラグ・オプションを指定することが必要です。

▶ アクセラレータ用のPGI Unified Binary

　PGI ンパイラは、異なるホストプロセッサ用に最適化がなされた関数を備え、一つの実行モジュール形態として実行モジュールを生成するための機能である PGI Unified Binary をサポートします。このリリースは、PGI Unified Binary をアクセラレータ用まで拡張します。特に、以下のような関数の二つのバージョンを含む単一バイナリを生成できます。

  • 一つは、アクセラレータ用に最適化したバージョン
  • 他は、アクセラレータが有効ではないとき、あるいは、アクセラレータ上での実行とホスト上での実行を比較したいときに、ホストプロセッサ上で実行するバージョン

この機能を有効にするには、拡張された-taフラグ・オプション-ta=nvidia,hostを使用します。このフラグは、コンパイラに対して、有効なアクセラレータ領域を有する、以下のような二つのバージョンの関数を生成するように指示する。

  • アクセラレータを対象としたコンパイル・バージョン
  • アクセラレータ・ディレクティブを無視し、ホストプロセッサを対象としたコンパイル・バージョン

　-Minfo フラグをコンパイル時に指定したときは、以下のコンパイル・メッセージと同じようなメッセージを得るでしょう。 

　s1: （ホスト上の executableバージョン）
      12, PGI Unified Binary version for -tp=barcelona-64 -ta=host
      18, Generated an alternate loop for the inner loop
          Generated vector sse code for inner loop
          Generated 1 prefetch instructions for this loop
  s1:（nvidia上の executableバージョン）
      12, PGI Unified Binary version for -tp=barcelona-64 -ta=nvidia
      15, Generating copy(b(:,2:90))
          Generating copyin(a(:,2:90))
      16, Loop is parallelizable
      18, Loop is parallelizable
           Parallelization requires privatization of array t(2:90)
          Accelerator kernel generated
            16, !$acc do parallel
            18, !$acc do parallel, vector(256)
    　Using register for t
　

上記では、PGI Unified Binaryのメッセージが、サブルーチン s1に関して、二つのバージョンを生成したことを示しています。

  • 一つは、アクセラレータのないバージョン（-ta=host）
  • もう一方は、NVIDIA GPUのためのバージョン（-ta=nvidia）

　実行時に、プログラムは NVIDIA CUDA 動的ライブラリをロードしようとします。そして、GPU の存在を確かめます。もし、ライブラリが有効でないか、あるいは、GPUが見つからない場合、プログラムはホストバージョンで実行されます。
　また、NVIDIA GPU 上で実行するようにプログラムに指示するために、環境変数をセットすることができます。これを行うためには、ACC_DEVICE を NVIDIA あるいは nvidia セットします。一方、この環境変数にこれ以外の任意の値をセットすると、ホストバージョンを使うようになります。
　なお、今回のリリースにおける -ta ターゲットは、"nvidia" と"host" の二つに限られます。

▶ 複数のプロセッサターゲット

　ta フラグと共に、複数のプロセッサターゲットを指定する形で、-tp フラグ・オプションも使用できます。この場合、次に述べるような挙動となります。

  • 一つの -tp　ターゲット値の指定と、一つの -ta 値を指定した場合、その指定された「プロセッサターゲット」と「アクセラレータターゲット」向けに生成された各サブプログラムを有する「一つの実行モジュール・バージョン」が生成される。
  • 一つの -tp　ターゲット値の指定と、複数の -ta ターゲット値を指定した場合、コンパイラは、その指定された「プロセッサターゲット」と各々の「アクセラレータターゲット」向けのアクセラレータ領域を含むサブプログラムを有する「二つの実行モジュール・バージョン」が生成される。
  • 複数の -tp　ターゲット値の指定と、一つの -ta の指定をした場合、二つあるいはそれ以上の「プロセッサターゲット」向けのそれぞれのサブプログラムのバージョンを一つの実行モジュール内に生成する。また、その各バージョンは、選択したアクセラレータ用のバイナリコードを含む。
  • 複数の -tp　ターゲット値の指定と、複数の -ta ターゲット値を指定した場合、例えば、N個の -tp の値と二つの -ta ターゲット値を指定すると、コンパイラは、N+1個のサブプログラムのバージョンを生成する。それは、始めに各 -tp ターゲット用にN個のバージョンを-ta=hostの指定と等価なアクセラレータ領域を無視したバイナリで生成します。そしてさらに、アクセラレータ向けの一つのバージョンを生成します。

▶ 複数のアクセラレータ

　並列のMPIあるいは、OpenMPプログラムを書くことにより、複数枚のNVIDIA GPUを使用することができます。

MPI
同じノード上で並列に各MPIランクが実行出来るMPIプログラムにおいては、acc_device_num手続きを使用して異なるGPUを選択するためにMPIランクの値を使用できます。
OpenMP
OpenMPスレッド並列プログラムにおいては、omp_num_thread_num関数を使用して、各スレッド用に異なるGPUを選択することができます。
OpenMP並列領域において、各スレッドがイテレーションの異なるサブセットを計算するようなループを持つアクセラレータ・リージョンを含むことができます。

▶ アクセラレータ・カーネルのプロファイリング

　本リリースでは、以下のコマンドライン・オプションをサポートします。

　　　-ta=nvidia,time

　Timeサブオプションは、タイマーライブラリをリンクします。このサブオプションは、アクセラレータ領域と生成されたカーネルに関するタイミング情報を集積し、印字します。

       Sample Accelerator Kernel Timing Data
       Accelerator Kernel Timing data 
       s1
          15: region entered 1 times
               time(us): total=1490738
                   init=1489138 region=1600
                   kernels=155 data=1445
               w/o init: total=1600 max=1600
                         min=1600 avg=1600
          18: kernel launched 1 times
              time(us): total=155 max=155 min=155 avg=155

上記の例において、いくつかの事象を見て取れます。

  • 各アクセラレータ領域において、ファイル名 /proj/qa/tests/accel/bb04.f90 とサブルーチンあるいは関数名 s1 とそのアクセラレータ領域の行番号が印字されています。上記例では、15行目という風に。
  • ライブラリは何回、領域に入ったかをカウントしています（上記例では、1）。そして、その領域内で消費した時間をマイクロ秒単位で表示します（上記例では、1490738）。この内訳は、初期化に係わる時間（例では、1489138）と実行時間（例では、1600）の二つに分けて表示します。
  • 実行時間は、「カーネル実行時間」と「ホストとGPU間のデータ転送時間」に分離して表示されます。
  • 各カーネルは、行番号が表示されます（上記例では18）。カーネルの開始数に沿って、カーネルで消費したtotal、max、min、平均の各時間が表示されます。この例では、1回のカーネル実行ですので、全て155 です。

▶ サポートする組込関数

　PGIアクセラレータ・コンパイラは、Fortran と C の組込関数、サブプログラムをサポートします。PGI組込関数に関する詳細な情報は、PGI User's Guide における7章「Using an Accelerator」の「Supported Intrinsics」をご参照下さい。さらに、組込関数については、今後のリリースで追加される予定です

以上

ページの先頭へトップへ