3434

チュートリアル

科学技術計算の分野で始まった GPU（Graphic Processing Unit）による汎用計算は、製造業、医療、金融など様々な分野へ広まり、モバイル、組み込みの分野でも使われ始めています。
当初は専用言語を使う必要があり GPU による汎用計算は気軽に始められるものではありませんでしたが、用途が広がると同時に、従来の言語で GPU 向けにプログラミングする環境が整ってきました。OpenACC はその中の一つで、既存のプログラムにディレクティブと呼ばれるヒント情報を追加するだけで、コンパイラが自動で GPU コードを生成する開発手法です。今回のチュートリアルでは、全４回のシリーズとして、この OpenACC を使ってアプ
リケーションを効率良く GPU で加速する方法を解説します。

OpenACC で始める GPU コンピューティング :OpenACC 概要

っ た G P U ( G r a p h i c P r o c e s s i n g U n i t ) に よ る 汎 用 計 算 は 、製 造 業 、医 療 、り 、モ バ イ ル 、組 み 込 み の 分 野 で も 使 わ れ 始 め て い ま す 。当 初 は 専

U による汎用計算は気軽に始められるものではありませんでした、 従 来 の 言 語 で GPU 向 け に プ ロ グ ラ ミ ン グ す る 環 境 が 整 っ て き ま

一つで、既存のプログラムにディレクティブと呼ばれるヒント情成瀬 彰

パ イ ラ が 自 動 で GPU コ ー ド を 生 成 す る 開 発 手 法 で す 。 今 回 の チ ュ

シ リ ー ズ と し て 、 こ の OpenACC を 使 っ て ア プ リ ケ ー シ ョ ン を 効 率

解説します。

１

はじめに

科学技術計算に使われる計算機システムは、多様化の時代を迎えている。2008年のTop500スパコンランキ
U コンピューティング:
OpenACC
概要[1]
ング
では、90%
以上のシステムがx86プロセッサを使
computing in OpenACC: Summary of OpenACC
用しており、HPC の世界は均質だった。その後、アク成瀬 彰 はメニコ
セラレータとして NVIDIA GPU が普及、Intelア プ ロ セ ッ サ の XeonPhi を リ リ ー ス、 米 CORAL プ ロ
ク並列性、細粒度な
SIMT 並列性[5]、やはり粒度の異[2]
ジェクトと供に IBM の Power プロセッサが HPC に復は、多様化なる
2 つの並列性を高めることで高性能を実現するア
帰、一部システムでは ARM プロセッサも使われ始め、
パコンラン
ーキテクチャである。近年の、そして、将来のプロセ
プロセッサ
ッサを有効に使うには、単にアプリケーションから大特定のアプリケーションに特化して FPGA や DSP も活。その後、
量の並列性を抽出するだけでなく、異なる粒度に合っ用されている。また、一つの計算ノードに、複数の種、Intel はメた並列性を抽出する必要があると言える。
類のプロセッサを搭載することも珍しくない。競争は
米 CORAL
もう一つの共通点は、メモリシステム階層の深化でイノベーションを加速する。多種多様なプロセッサのロセッサがある。CPUのメモリ階層は、主記憶（典型的には DDR
セッサも使メモリ)、同一チップ内のCPU コア間でシェアされる
存在は、高い性能を求めるユーザには朗報である。
して FPGA
共有キャッシュ、CPU
コア内のプライベートキャッシ
しかし、ソースコードのポータビリティの確保が問ノードに、ュで構成される。一方、GPU のメモリ階層は、CPU の題である。プログラム言語や開発手法も多様化してお珍しくない。 主記憶、GPU の主記憶（DDR メモリよりスループッ
り、事前の検討が不十分だと、特定のプロセッサ以外
様なプロセ
トの高い
GDDR メモリが典型的）、複数階層のキャッでは、開発したプログラムが動作しない、もしくは、は朗報であシュ、スクラッチパッドメモリで構成される。
まともな性能が出ないという状況になりかねない。後の確保が問続きは Web で者は、性能ポータビリティと呼ばれる問題だが、多様様化してお
日本計算工学会誌「計算工学(Vol.21, No.2)」HP:
なプロセッサが存在する中で、どうすれば性能ポータ
セッサ以外
http://www.jsces.org/Issue/Journal/
ビリティを確保できるのかは、大きな課題となってい
もしくは、

ねない。後
だが、多様
性能ポータ
となってい
ロセッサ間
。
「並列性」
とである。
MD 幅を広ッサのコアになり[3]、て AVX-512
、粒度の異
方針を採っ
度のブロッ

筆者紹介

筆者紹介

る。多様化が進む中での救いは、これらプロセッサ間に少なからぬ共通点が存在することである。
一点目の共通点は、全てのプロセッサが「並列性」を高める方向で、高性能化を指向していることである。
CPU は、CPU コア数を増やすと当時に、SIMD 幅を広げる傾向にある。例えば Intel Xeon プロセッサのコア数
は10年前の2コアから18コアと9倍になり [3]、SIMD 幅も10年 前 の SSE2の128bit と 比 べ て AVX-512は512bit で 間もなく4倍になる [4]。つまり、粒度の異なる2つの並列性を上げて、性能を高める方針を採っている。GPUは、CPU とは異なるが、粗粒度のブロック並列性、細粒度な SIMT 並列性 [5]、やはり粒度の異なる2つの並列性を高めることで高性能を実現するアーキテクチャである。近年の、そして、将来のプロセッサを有効に使うには、単にアプリケーションから大量の並列性を抽出するだけでなく、異なる粒度に合った並列性を抽出する必要があると言える。
もう一つの共通点は、メモリシステム階層の深化である。CPU のメモリ階層は、主記憶（典型的には DDRメモリ）、同一チップ内の CPU コア間でシェアされる
共有キャッシュ、CPU コア内のプライベートキャッシュで構成される。一方、GPU のメモリ階層は、CPUの主記憶、GPU の主記憶（DDR メモリよりスループッ
トの高い GDDR メモリが典型的）、複数階層のキャッ
シュ、スクラッチパッドメモリで構成される。

なるせAkira NARUSE
なるせあきらあきら
1996
年、名古屋大学大学院工学系研究科修士課
1996年、名古屋大学大学院工学系研究科修士課程修了（情報工学専攻）。同年、富士通研究所に入社、程修了(情報工学専攻)。同年、富士通研究所に入社、大規模サーバ開発、HPC システム開発な

大規模サーバ開発、HPC システム開発など、様々

ど、様々なプロジェクトに参加、関連するハーなプロジェクトに参加、関連するハードウェアや続きは Web で
日本計算工学会誌「計算工学（Vol.21, No.2）」HP:
http://www.jsces.org/Issue/Journal/

ドウェアやソフトウェアの研究開発に従事。

ソフトウェアの研究開発に従事。2009年、山下記2009 年、山下記念研究賞受賞。2010 年、念研究賞受賞。2010年、SACSIS2010最優秀論文賞

SACSIS2010 最優秀論文賞受賞。2013 年、NVIDIA に参加、シニアデベロッ受 賞。2013年、NVIDIA に 参 加、 シ ニ ア デ ベ ロ ッ

パーテクノロジーエンジニアとして様々なアルゴリズムとアプリケーショ

パーテクノロジーエンジニアとして様々なアルゴ

ンの GPU 向け並列化に従事、GPU コンピューティングの普及に努めている。

リズムとアプリケーションの GPU 向け並列化に従事、GPU コンピューティングの普及に努めている。

（28）

計算工学

チュートリアル

OpenACC で始める GPU コンピューティング：OpenACC 概要

近年のプロセッサには、このような並列化・階層化という共通性があるのに加え、内在するアーキテク
チャの複雑性をハードウェアレベルで隠蔽せずに、ソフトウェアに解放する傾向にある。つまり、プロセッ
サの機能を最大限に活用するには、これまで以上に、その機能をソフトウェアから使いこなす必要があり、そのためプロセッサベンダーは各種 API やプログラミ
ングモデルを開発者に提供している。ベンダー提供の
API やプログラミングモデルを使うと、そのベンダーのプロセッサ上でのアプリケーション性能は向上する
が、コードのポータビリティが問題として残る。高いレベルで性能とポータビリティを両立するプログラミ
ングモデルが存在するのが理想であるが、一般的には、性能とポータビリティはトレードオフの関係にあり、単一のプログラミングモデルでそれを両立させるのは容易ではない。従って、大規模なアプリケーショ
ンになるほど、複数のプログラミングモデルを組み合わせて使用するケースが増えている。

2

高並列プロセッサ向けプログラミングモデル

本稿は OpenACC による高並列プロセッサ向けのプログラム開発手法を説明するものだが、だからといってOpenACC のみを使用してプログラム並列化を進めよと推奨するものではない。OpenACC はシンプルで強力なプログラミングモデルだが、決して万能ではない。世の中には様々なプログラミングモデルが存在し、それぞれ長所と短所がある。我々が推奨しているのは、
OpenACC を中核として、様々なプログラミングモデルのいいところを利用して、効率良くプログラム開発を進める方法である。
HPC 分野で使用されているプログラミングモデル
は、大別すると次の4種類である。
1. 標準言語
2. ライブラリ
3. ディレクティブ
4. 拡張言語
以下では、これらプログラミングモデルのそれぞれの特徴に触れた後、我々の推奨方法を説明する。
2.1 標準言語
C/C++ や Fortran などの標準言語で開発されたソース
コードは、当然ながら、最もポータビリティが高い。
しかし、現段階で標準言語が提供している機能は、高並列プロセッサにとっては十分とは言えない。標準言語コミュニティ内でも高並列時代への対応が始まっているが、将来の規格で、高並列プロセッサ向け機能をどう標準化するか、議論されているところである。また、新しい規格がリリースされても、コンパイラベン
ダーから新機能が安定した状態で提供されるまでには数年かかるのが常である。従って、今後数年間は、高並列を指向するプロセッサ向けに、標準言語で高性能を実現するのは難しい状態が続くと考えられる。

計算工学

2.2 ライブラリ
BLAS などの標準ライブラリの使用は、ポータビリ
ティの高い方法である。API が標準化されているため、使用プロセッサを変更しても、ソースコードを変更することなく、ライブラリ実装を切り替えるだけで、プログラムを動作させることができる。また、プロセッサベン
ダーは、標準ライブラリに関しては、高度にチューニン
グされた実装を提供する傾向にあり、高性能も実現できることが多い。例えば、BLAS の中でも使用頻度の高
い 行 列 積 ル ー チ ン は、Intel プ ロ セ ッ サ 上 で は IntelMKL[6] が、NVIDIA GPU 上 で は NVIDIA cuBLAS[7] が それぞれ最速であり、どちらもユーザコードではほぼ実現不可能なレベルの性能を提供する。標準ライブラリの使用は、ポータビリティと性能を高いレベルで両立できる方法であるが、残念ながら、標準ライブラリだ
けで開発できるアプリケーションはほとんどない。
プロセッサベンダーは標準ライブラリ以外にも、ユーザのために独自ライブラリを開発・提供することがある。独自ライブラリ使用のメリットは高性能である が、API が 標 準 化 さ れ て い な い た め、 ポ ー タ ビ リ
ティが問題となる。ただ、ライブラリの適用箇所は局所的であることが多いので、ベンダー独自ライブラリを使っても、アプリケーション全体としてはポータビ
リティを高い状態に保つことができる。
2.3 ディレクティブ
ディレクティブは、標準言語で開発された既存のプ
ログラムに対して、コンパイラへのヒントを追加する方法である。一般的には、ディレクティブは、C/C++の場合は pragma 行として、Fortran の場合はコメント行
として既存コードに追加される。ディレクティブは、標準言語ではサポートされていない機能を提供する手段として使われることが多く、プログラム開発者はディレクティブを使ってコンパイラに対してどのようにコードをビルド・最適化すべきかを指示できる。
ディレクティブにはコンパイラベンダー独自のもの
も あ る が、 広 く 使 わ れ て い る の は、OpenMP [8] やOpenACC[9] のように、産学グループで標準化され、複数のコンパイラベンダーからサポートされるディレク
ティブである。標準化されたディレクティブを使ったプログラムであれば、使用するプロセッサやコンパイ
ラを変えても、プログラムをビルド・実行することができる。
一般的には、ディレクティブで最適化したプログラ
ムは、ハンドチューニングしたコードと比べると、１〜２割程度低い性能となることが多い。ディレクティ
ブ が 提 供 す る 機 能 を フ ル に 使 い こ な せ ば、 ハ ン ド
チューニングコードに近い性能を引き出すことも可能だが、そのようにディレクティブを追加すると、対象プロセッサに特化したコードとなりやすく、性能ポー
タビリティが低下する。また、ソースコードの可読性も低下する。それではディレクティブのメリットが損
なわれてしまう。１〜２割の性能ロスの代償はあるも（28-2）

Vol.21, No.2 2016

チュートリアルチュートリアルOpenACC
で始めるGPU
GPUコンピューティング:
コンピューティング：OpenACC
概要
OpenACC で始める
OpenACC 概要(1)てしまう。１～２割の性能ロスの代償はあるものの、のの、コードのポータビリティと可読性を確保できる
コードのポータビリティと可読性を確保できるため、ため、ディレクティブは多くのアプリケーション、特ディレクティブは多くのアプリケーション、特に大規に大規模なアプリケーションで採用されている。
模なアプリケーションで採用されている。
また、追加したディレクティブは、ディレクティブ
また、追加したディレクティブは、ディレクティブに未対応のコンパイラでは認識されないが、ビルド・に未対応のコンパイラでは認識されないが、ビルド・実行はできる。ソースコードを標準コンパイラでもビ実行はできる。ソースコードを標準コンパイラでもビルドできる状態に保てるという意味でも、ディレク
ルドできる状態に保てるという意味でも、ディレクテ
ティブはポータビリティの高い方法である。
ィブはポータビリティの高い方法である。
2.4 拡張言語2-4 拡張言語とは、例えば
拡張言語
CUDA[10] や OpenCL[11] のように、特定のプロセッサの機能を最大限に引き出すため
拡張言語とは、例えば CUDA[10]や OpenCL[11]のよに拡張された言語である。CUDA と OpenCL はベース
うに、特定のプロセッサの機能を最大限に引き出すた言語として C/C++ を採用しているが、GPU等の高並列めに拡張された言語である。CUDAと OpenCL
はベープロセッサの性能を活用しやすいよう拡張されていス言語としてC/C++を採用しているが、GPU 等の高並る。拡張言語でのプログラミングは低レベルの記述が列プロセッサの性能を活用しやすいよう拡張されてい必要になりがちであるが、拡張言語を使わないと活用る。拡張言語でのプログラミングは低レベルの記述ができないハードウェア機能が存在するため、非常に高必要になりがちであるが、拡張言語を使わないと活用い性能が求められるアプリケーションでは拡張言語ができないハードウェア機能が存在するため、非常に高使用される。
い性能が求められるアプリケーションでは拡張言語が一方で、拡張言語で開発されたコードは、例えば使用される。
CUDA
OpenCL
コードは、標準のコンパ一 方 でコードや、拡張言語で開発 さ れ た コ ー ドC/C++
は、例
えば
イラではビルドすることができない。また、拡張言語
CUDA
コードや OpenCL コードは、標準の C/C++コンはその特性から、特定プロセッサでしかサポートされパイラではビルドすることができない。また、拡張言ない傾向にあり、必然的にコードのポータビリティは
語はその特性から、特定プロセッサでしかサポートさ低くなる。拡張言語は、性能面では魅力的だが、ポーれない傾向にあり、必然的にコードのポータビリティ
タビリティに難点があるため、採用の前には十分な検は低くなる。拡張言語は、性能面では魅力的だが、ポ討が必要である。
ータビリティに難点があるため、採用の前には十分な検討が必要である。
2.5 プログラミングモデルの併用
2-5 一つのプログラミングモデルで、性能とポータビリ
プログラミングモデルの併用
ティを高いレベルで両立するのは難しい。しかし、単一つのプログラミングモデルで、性能とポータビリ
一のプログラミングモデルに依存する必要はない。ほティを高いレベルで両立するのは難しい。しかし、単
とんどのアプリケーションでは、複数のプログラミン
一のプログラミングモデルに依存する必要はない。ほグモデルを併用して、プログラムを開発することが可
とんどのアプリケーションでは、複数のプログラミン能である。従って、高並列プロセッサ向けに、高性
グモデルを併用して、プログラムを開発することが可能、かつ、メンテナンスの容易なコード開発を要求さ能である。
従って、高並列プロセッサ向けに、高性能、れる開発者への推奨方法は次の方法となる。
かつ、メンテナンスの容易なコード開発を要求される
1. ライブラリで可能な部分はライブラリに置換える
開発者への推奨方法は次の方法となる。
2. コードの並列化は、最も生産性の高いプログラミ
1. ライブラリで可能な部分はライブラリに置換える
ングモデルで開始する
2. コードの並列化は、最も生産性の高いプログラミ
3. 性能が必要なモジュールのみ、低レベルなプログ
ングモデルで開始する
ラミングモデルを適用する
3. 性能が必要なモジュールのみ、低レベルなプログ
ライブラリは、標準ライブラリの使用が望ましい
ラミングモデルを適用するが、ニーズによってはベンダーライブラリも OK であライブラリは、標準ライブラリの使用が望ましいが、る。2. と3. の部分は、言い換えると、ディレクティブ
ニーズによってはベンダーライブラリもOK である。
の使用から始めて、必要な部分だけ低レベルの拡張言
2.と 3.の部分は、言い換えると、ディレクティブの使語を使うということである。現時点では、GPU に対応
用から始めて、必要な部分だけ低レベルの拡張言語をし標準化されたディレクティブは OpenACC
のみであ使うということである。現時点では、GPUに対応し標り、GPU をターゲットプロセッサとするなら、高性能準化されたディレクティブはOpenACC のみであり、とポータビリティを両立させる現実的なアプローチはGPUをターゲットプロセッサとするなら、高性能とポ
OpenACCー タ ビ リ テでコード開発を始めることとなる。
ィを両立させる現実的なアプローチは
OpenACC でコード開発を始めることとなる。
計算工学
計算工学

3 OpenACC とは
3 HPC
OpenACCとは分野では、メニコアアーキテクチャやGPU などHPC 分野では、メニコアアーキテクチャや GPU なの高並列プロセッサの出現とともに、高級言語で性能
どの高並列プロセッサの出現とともに、高級言語で性を引き出したい、コードのポータビリティと実効性能能を引き出したい、コードのポータビリティと実効性を高いレベルで両立させたいというプログラマーから能を高いレベルで両立させたいというプログラマーかの ニ ー ズ が 高 ま っ て い た。 そ れ に 応 え る 形 で、らの ニ ー は、アクセレレータなど並列性の高いプロズが高まっていた。それに応える形で、OpenACC
OpenACCは、アクセレレータなど並列性の高いプロセ
セッサ向けのプログラミングモデルとして2011年に提ッサ向けのプログラミングモデルとして2011 年に提案された。OpenACC はディレクティブ方式を採用して案された。
OpenACC はディレクティブ方式を採用して
おり、コード開発者はコード内の並列化が可能な部分
おり、コード開発者はコード内の並列化が可能な部分
やその並列化の方法を、ディレクティブを追加することでコンパイラに指示することができる。そして、コやその並列化の方法を、ディレクティブを追加するこ
ンパイラはその情報を元に、様々な並列プロセッサ向とでコンパイラに指示することができる。そして、コけのコードを生成する。
ンパイラはその情報を元に、様々な並列プロセッサ向
けのコードを生成する。
3.1 アクセレレータ・モデル
で開発されたコードが、様々なプロセッサ
3-1OpenACC
アクセレレータ・モデルで実行できるようにするために、OpenACC
ではアクセ
OpenACC で開発されたコードが、様々なプロセッ
ラレータ・コンピューティングを抽象化したモデルがサで実行できるようにするために、
OpenACC ではアク定義されている。このモデルは、対象となるプロセッ
セラレータ・コンピューティングを抽象化したモデル
サでは様々な粒度の並列性が必要であること、メモリが定義されている。このモデルは、対象となるプロセ
システムは容量・速度・アドレスの異なる階層から構
ッサでは様々な粒度の並列性が必要であること、メモ成されることを想定している。アクセラレータ・モデ
リシステムは容量・速度・アドレスの異なる階層から
ルの目的は、現時点で広く使われているプロセッサに構成されることを想定している。アクセラレータ・モ
対してだけでなく、これから出現するであろう全てのデルの目的は、現時点で広く使われているプロセッサ
プロセッサに対して、OpenACC を適用可能にすることに対してだけでなく、これから出現するであろう全てである。
のプロセッサに対して、
OpenACC を適用可能にするこ
とである。
ホスト
デバイス

ホスト
メモリ

デバイス
メモリ

図 図1
1: OpenACC
OpenACCのアクセラレータ・モデルのアクセラレータ・モデル
OpenACCの実行モデルでは、プログラム実行を主導の実行モデルでは、プログラム実行を主OpenACC
導するのはホストであり、デバイス（アクセラレータ）
するのはホストであり、デバイス（アクセラレータ）
ははホストの指示に従って動作する。ホストは、プログ
ホストの指示に従って動作する。ホストは、プログラ
ラム内の計算負荷の重い部分をアクセラレータにオフ
ム内の計算負荷の重い部分をアクセラレータにオフロードし、必要な場合はホストメモリとデバイスメモ
ロードし、必要な場合はホストメモリとデバイスメモ
リ間でデータを転送する。図1 にアクセラレータ・モ
リ間でデータを転送する。図1にアクセラレータ・モデ
デルの概略を示す。この図は、近年の
GPU
搭載ノード
ルの概略を示す。この図は、近年の GPU
搭載ノードでで見られるような、ホストとしてCPUを使い、デバイ
見られるような、ホストとして CPUを使い、デバイススとしてGPUを使う形態を模している。この形態のシとして GPUを使う形態を模している。この形態のシスステムでは、ホストとデバイスでメモリ空間が分離しテムでは、ホストとデバイスでメモリ空間が分離してているので、アクセラレータにオフロードされる計算いるので、アクセラレータにオフロードされる計算で（28-3）(??-3)

Vol.21,
No.2
2016
Vol.21,
No.2
2016

チュートリアル

OpenACC で始める GPU コンピューティング：OpenACC 概要

使われるデータは、事前に、OpenACC ランタイムによりホストからデバイスへ転送される。同様に、デバイ
スメモリに書き込まれたアクセラレータの計算結果
は、 ホ ス ト プ ロ セ ッ サ が そ れ を 必 要 と す る 前 に、
OpenACC ランタイムによりデバイスからホストへと転送される。
図1より、ホストとデバイスが別個のプロセッサであることをこのモデルは前提としているように見えるかもしれない。実際には、ホストとデバイスは同じアー
キテクチャでも問題なく、同一チップでも構わない。
その場合のプログラムの実行方式は、マルチコア CPU
上での OpenMP プログラムの実行に近いものとなる。
メモリも同様であり、ホストメモリとデバイスメモリ
が同一のメモリでも問題ない。この場合は、単にホストとデバイス間のデータ転送が不要になるだけである。
3.2 メリットと制限
OpenACC のメリットは、高級言語でアクセラレータ
向けコード開発を行えること、そして、開発したコードの適用対象が特定プロセッサに限定されないことである。プログラマは、OpenACC 化した単一のソースコー
ドから、様々なプロセッサ上で動作するプログラムを開発できる。例えば、PGI の OpenACC コンパイラ [12] は、
コンパイルオプションを変えるだけで、単一のソースコードから NVIDIA GPU、AMD GPU、Intel マルチコア
CPU 向けのコードを生成できる。しかし、この高いポータビリティは性能面では制限となることもある。
OpenACC のアクセラレータモデルは、想定される
様々なアクセラレータの共通項に基づいて決定されている。これは、特定のアクセラレータに固有な機能をモデルに取り入れることにより、OpenACC で開発されたコードのポータビリティが低下するのを回避するためである。従って、CUDA や OpenCL のような、低レ
ベルなプログラミングモデルで容易に実装できる最適化手法の中には、OpenACC では実現するのが難しいものもある。例えば NVIDIA GPU の場合、チップ内の「共有 メ モ リ」と 呼 ば れ る ス ク ラ ッ チ パ ッ ド メ モ リ は、
CUDA な ら 比 較 的 容 易 に 活 用 で き る の に 対 し て、
OpenACC には「cache」ディレクティブが用意されているものの、OpenACC で共有メモリを活用するのは容易ではない。同様の事例は、他のプロセッサやアクセラ
レータにも存在する。特定のプロセッサにのみ有効な最適化手法ほど、そのプロセッサに密接した低レベルなプログラミングモデルなら容易に実現できるにもかかわらず、OpenACC では実現が難しい傾向にある。
性能が重視されるアプリケーションに OpenACC を使
うのか、それとも低レベルな言語を採用するかは、それぞれのメリットとデメリットを把握した上で、開発者が慎重に決定する必要がある。しかし、性能が非常に重要で、低レベルな言語を使うしかない場合であっても、OpenACC は低レベル言語と組み合わせて使うこと が で き る の で、 ア プ リ ケ ー シ ョ ン の 大 部 分 はOpenACC で高速化して、コード内の必要最小限の部分
計算工学

だけ低レベル言語を採用するアプローチも可能である。
3.3 ディレクティブの構文
本節では OpenACC の構文を説明する。OpenACC の様々なディレクティブの機能に関しては、後述するサン プ ル コ ー ド を 事 例 と し て、 ア プ リ ケ ー シ ョ ン をOpenACC で高速化するプロセスの各過程で、関連する
ディレクティブの機能を詳細に説明する予定である。
なお、OpenACC の構文は OpenMP の構文を踏襲しており、OpenMP の経験者には違和感が少ないだろう。
OpenACC ディレクティブの中身は、基本的には言語に関係なく共通であるが、コードへの追加方法は、以下に示す通り、C/C++ では pragma 形式、Fortran では特
殊なコメント行形式で追加する方法である。
[C/C++]#pragma acc directive-name [clause [clause […]]][Fortran]!$acc directive-name [clause [clause […]]]
キ ー ワ ー ド acc の 後 に 記 述 さ れ た 内 容、 上 記 例 の「directive-name [clause [clause […]]]」の部分がディレク
ティブの中身である。この中の directive-name の箇所に記載する内容がコンパイラへの指示内容のメインである。角括弧内はオプションであり、clause の箇所に記載する内容は前記 directive-name の指示内容をより詳細化するためのものである。一つの directive-name に対して、複数の clause を付けることもできる。
OpenACC ディレクティブの中には、コード領域に対し て 適 用 さ れ る デ ィ レ ク テ ィ ブ が あ る。 例 え ば、
kernels ディレクティブはコード内の並列化可能領域を
コンパイラに指示するものであるが、C/C++ の場合は、kernels ディレクティブを追加した行の次行から始まる
波括弧で囲まれたコード領域、いわゆる構造ブロック
が、その適用対象領域となる。
[C/C++]#pragma acc kernels{(computations …)}
Fortran では、C/C++ のように構造ブロックをコード内に記述するのが一般的ではないため、kernels ディレ
クティブを追加した行の次行から、それに対応する終了ディレクティブ（end kernels ディレクティブ）を追加
した行の前行までのコード領域が、適用対象領域となる。
[Fortran]!$acc kernels(computations …)!$acc end kernels（28-4）

Vol.21, No.2 2016

チュートリアル

4

OpenACC で始める GPU コンピューティング：OpenACC 概要

OpenACC によるアプリケーション並列化のチュートリアル
推奨方法

本章ではアプリケーションを OpenACC で並列化する
4 OpenACC に よ るア プ リケ ー ショ ン
並列 化 の推 奨
場合の推奨方法を説明する。OpenACCに限らず、アプ
方法
リケーションを並列化するときに、コード開発者が頻
本章ではアプリケーションを OpenACC で並列化す繁に直面する問題の一つは、プログラムの計算結果がる場合の推奨方法を説明する。OpenACC に限らず、ア
並列化前と並列化後で一致しないことである。並列化
プリケーションを並列化するときに、コード開発者がにより演算順序が変わり、浮動小数点演算の丸め誤差
頻繁に直面する問題の一つは、プログラムの計算結果が変わったことがその原因かもしれない。これは、数が並列化前と並列化後で一致しないことである。並列学的にはどちらも正しいので、許容すべき不一致と言化により演算順序が変わり、浮動小数点演算の丸め誤える。しかし、並列化の過程でバグが混入、もしくは差が変わったことがその原因かもしれない。これは、バグが表面化したのかもしれない。こちらは修正が必
数学的にはどちらも正しいので、許容すべき不一致と言える。しかし、並列化の過程でバグが混入、もしく
要な不一致である。結果が一致しない原因は、許容しはバグが表面化したのかもしれない。こちらは修正がて良いものと、修正が必要なバグが存在するので、そ必要な不一致である。結果が一致しない原因は、許容
れぞれ明確に切り分ける必要がある。しかし、開発者して良いものと、修正が必要なバグが存在するので、がプログラムの多数の箇所を並列化した後に、この結それぞれ明確に切り分ける必要がある。しかし、開発果が一致しない問題に直面すると、並列化したコード者がプログラムの多数の箇所を並列化した後に、このの中で、どの部分は OK で、どの部分は修正が必要な結果が一致しない問題に直面すると、並列化したコーのか、その切り分け作業に膨大な時間を費やすことにドの中で、どの部分は OK で、どの部分は修正が必要なる。場合によっては、並列化前のコードに戻って、なのか、その切り分け作業に膨大な時間を費やすこと全てをやり直すことになりかねない。
になる。場合によっては、並列化前のコードに戻って、全てをやり直すことになりかねない。
1) ホットスポ
ットの特定
4) ループの最適化

2) ループの並列化

2: OpenACC
によるアプリケーション並列化の
図2 図
OpenACCによるアプリケーション並列化のサイクル
サイクル

それ故に、我々はアプリケーションを一挙に並列化
それ故に、我々はアプリケーションを一挙に並列化するのではなく、重要な部分から一つずつ順番に並列するのではなく、重要な部分から一つずつ順番に並列化していく方法を推奨している。その具体的なプロセ化していく方法を推奨している。その具体的なプロセスは、図2の通りである。1）コード内の実行時間の長い
スは、図 2 の通りである。1) コード内の実行時間の長場所（ホットスポット）の特定、2）ホットスポットをい場所（ホットスポット）の特定、2) ホットスポット
OpenACC
で並列化、3）
ホストとアクセラレータ間の無を OpenACC
で並列化、3)
ホストとアクセラレータ間用なデータ転送を削減、4）
対象アクセラレータ向けにの無用なデータ転送を削減、
4) 対象アクセラレータ向
ループを最適化。そして、アプリケーション全体の性
けにループを最適化。そして、アプリケーション全体
能が所望の性能に達するまで、このプロセスを繰り返の性能が所望の性能に達するまで、このプロセスを繰すのである。
り返すのである。
4-1 ホットスポットの特定
ホットスポットの特定
4.1
アプリケーションの並列化に取り組む前に、まず、アプリケーションの並列化に取り組む前に、まず、アプリケーションのどのサブルーチン、どのループがアプリケーションのどのサブルーチン、どのループが全実行時間のどの程度の時間を要しているかを把握す
全実行時間のどの程度の時間を要しているかを把握することは重要である。アムダールの法則[13]によれば、ることは重要である。アムダールの法則 [13] によれば、並列化したアプリケーションの最終性能は、アプリケ
並列化したアプリケーションの最終性能は、アプリ
ケーションの逐次部分の実行時間に支配される。つま
計算工学

性能解析ツールは豊富に存在するが、効率良くホット
ホットスポットの並列化
ス4.2
ポットを特定できるのであれば、ツールは何でも高速化すべき場所を決定したら、次はその中のどのOK である。ユーザタイマーでも構わない。

部 分 が、 ど の ル ー プ が 並 列 化 で き る か を 調 査 し、

4-2
ホットスポットの並列化
OpenACC
ディレクティブを使って、どのループをどう
高速化すべき場所を決定したら、次はその中のどの並列化するかをコンパイラに指示する。コンパイラへ部の指示は、最初から完璧を目指す必要はない。多くの分が、どのループが並列化できるかを調査し、OpenACC
どのループをどう
OpenACCディレクティブを使って、コンパイラは、どう並列化したかのフィード並列化するかをコンパイラに指示する。コンパイラへ
バック・メッセージをコンパイル時に出力するので、の指示は、最初から完璧を目指す必要はない。多くのそれを参考にして並列化方法を詳細化していけばよい。
OpenACC コンパイラは、どう並列化したかのフィード
また、この段階では、ホストとデバイス間で冗長にバック・メッセージをコンパイル時に出力するので、データ転送が行われていることが多く、その結果としそれを参考にして並列化方法を詳細化していけばよい。
て、アクセラレータを使用する前より、性能が低下す
また、この段階では、ホストとデバイス間で冗長にデることも多い。この段階で無理にアクセラレータ上のータ転送が行われていることが多く、その結果として、計算性能を最適化する必要はない。
アクセラレータを使用する前より、性能が低下することも多い。この段階で無理にアクセラレータ上の計算性能を最適化する必要はない。
4.3 無用なデータ転送の削減

CPU+GPU のように、現在のアクセラレータ搭載ノー3) データ転
送の最適化

計算工学

り、並列性能を高めるには、逐次実行時間の削減が重OpenACC で始める GPU コンピューティング: OpenACC 概要 (1)
要ということである。逐次実行時間の長いところ、CPU 実行時のホットスポットから並列化に取り組むとーションの逐次部分の実行時間に支配される。
つまり、いうのは、効率良く並列化の効果が得られるという意
並列性能を高めるには、逐次実行時間の削減が重要と味でも理にかなっている。Gprof[14] などフリーのツールいうことである。逐次実行時間の長いところ、CPU 実
から、より詳細に解析できる有償のツールまで、CPU
行時のホットスポットから並列化に取り組むというの向けの性能解析ツールは豊富に存在するが、効率良くは、効率良く並列化の効果が得られるという意味でもホットスポットを特定できるのであれば、ツールは何理にかなっている。Gprof[14]などフリーのツールから、でも OK である。ユーザタイマーでも構わない。
より詳細に解析できる有償のツールまで、CPU向けの4-3
無用なデータ転送の削減
ドでは、ホストとアクセラレータでメモリが分離してCPU+GPUのように、現在のアクセラレータ搭載ノいることが多い。この条件で正しい結果を得るためにードでは、ホストとアクセラレータでメモリが分離しは、コンパイラは、ホストとアクセラレータ間で適切ていることが多い。この条件で正しい結果を得るためにデータを転送するコードを生成しなくてはいけなには、コンパイラは、ホストとアクセラレータ間で適い。しかし、ソースコードからそのアプリケーション
切にデータを転送するコードを生成しなくてはいけなに関してコンパイラが読み取れる情報は、その一部でい。しかし、ソースコードからそのアプリケーションしかない。その結果として、正しい計算結果が得られに関してコンパイラが読み取れる情報は、その一部でることを保証するため、コンパイラが生成するコードしかない。その結果として、正しい計算結果が得られでは、必要以上にホストとアクセラレータ間でデータることを保証するため、コンパイラが生成するコード転送が行われることが多い。コード開発者はそのアプでは、必要以上にホストとアクセラレータ間でデータ転送が行われることが多い。コード開発者はそのアプ
リケーションに関してコンパイラより多くの知識をリケーションに関してコンパイラより多くの知識を持持っている。加えて、OpenACC にはデータ転送をコンっている。加えて、OpenACCにはデータ転送をコント
トロールするディレクティブが用意されている。コー
ロールするディレクティブが用意されている。コードド開発者はこのディレクティブを使って、できる限り
開発者はこのディレクティブを使って、できる限り長長い間、データをアクセラレータメモリ上に滞在さい間、データをアクセラレータメモリ上に滞在させ、せ、ホストとアクセラレータ間のデータ転送を最小化
ホストとアクセラレータ間のデータ転送を最小化する
することができる。現行の GPU 搭載ノードは、CPU とことができる。現行の GPU 搭載ノードは、CPU と GPU
GPU 間が、主記憶のメモリバンド幅と比べると、かな
間が、主記憶のメモリバンド幅と比べると、かなり帯
り帯域の狭いPCI エクスプレスで接続されていること
域の狭い PCI エクスプレスで接続されていることもあもあり、このデータ転送の最適化がアプリケーション
り、このデータ転送の最適化がアプリケーション性能
性能に寄与する効果は大きい。
に寄与する効果は大きい。
4-4
4.4高並列向けにループを最適化
高並列向けにループを最適化
シンプ
ルに ル ープが 並列 化 可能で ある こ とをコ ン
シンプルにループが並列化可能であることをコンパ
パイラに指示するだけでも、コンパイラは自身の知識
イラに指示するだけでも、コンパイラは自身の知識に基づいて、様々な粒度の並列性をコードから抽出し、(??-5)（28-5）

Vol.21, No.2 2016

Vol.21, No.2 2016

チュートリアル

OpenACC で始める GPU コンピューティング：OpenACC 概要

それをターゲットプロセッサの並列特性にマップする。しかし、前述の通り、コンパイラが持つアプリ
ケーションの知識には制限があるので、コード開発者
がより詳細に並列化方式を指示することで、更に性能を向上させることも可能である。
それよりも性能向上の可能性が高いのは、コードの再構成である。計算負荷の重い部分は、複数の多重ループから構成されていることが多い。往々にして、
これらコードはマルチコア CPU をターゲットに開発されていることが多く、マルチコア CPU には十分な並列性、つまり数〜数十程度の並列性は抽出できている。
しかし、ループの構成を見直すことで、より多くの並列性、数千、数万、それ以上の並列性を抽出することも可能である。GPU は超並列プロセッサであり、その性能を最大限に活用するには、数万のオーダーの並列性が必要である。このようなコードの再構成は、アクセラレータ上で高い性能を得られるのはもちろんのこと、今後、全てのプロセッサが高並列へと向かっていることを考えると、将来的には全てのプロセッサにとって有意義と言える。
4.5 結果検証とツール使用
本章の最初に述べた通り、OpenACC でアプリケー
ションを並列化・高速化する過程では、継続的に計算結果が正しいことを確認することが非常に重要である。シミュレーション結果の一部、例えばある物理量が格納されている配列内の最小値・最大値・平均値な
ど、また収束計算なら残差・収束回数など、計算結果を検証できるデータを定期的に出力するコードを追加することを推奨する。定常的に計算結果が正しいことを検証できれば、コード開発の手戻りが減り、結果的にはコード開発を効率良く進めることができる。また、Git[15] などのソースコードのバージョン管理ツール
を使うことも合わせて推奨する。こまめにスナップ
ショットを保存しておけば、それまでの作業工数を大幅に無駄にすることなく、素早く正常なソースコードに戻ることができる。

5

サンプルコード : ヤコビ反復法

本チュートリアルの2回目と3回目では、比較的シン
プルなサンプルコードを題材として、前章で説明した、OpenACC によるアプリケーション並列化のプロセスに沿って、具体的な OpenACC ディレクティブの使い方とその機能を説明していく。最初に用いるサンプル
コードは、ヤコビ反復法による2次元ラプラス方程式の解法である。本サンプルコードのメインループ部分
は、C/C++ と Fortran で、それぞれ以下の通りである。
[C/C++]
1 while ( error > tol && iter < iter_max ) {2
error = 0.0
3
for (int j = 1; j < N-1; j++ ) {4for (int i = 1; i < M-1; i++ ) {計算工学

5Anew[j][i] = 0.25 *6(A[j][i-1] + A[j][i+1] + A[j-1][i] + A[j+1][i]);7error = fmax(error, fabs(Anew[j][i] – A[j][i] ));
8
}
9
}
10
for (int j = 1; j < N-1; j++ ) {11
for (int i = 1; i < M-1; i++ ) {12
A[j][i] = Anew[j][i];
13
}
14
}
15
if (iter % 100 == 0) printf( %d, %f￥n , iter, error)
16
iter++;
17 }[Fortran]
1 do while ( error > tol .and. iter < iter_max )2error = 0.0
3
do j=2, N-1
4
do i=2, M-1
5
Anew(i,j) = 0.25 * &6(A(i-1,j) + A(i+1,j) + A(i,j-1) + A(i,j+1))
7
error = max(error, abs(Anew(i,j) – A(i,j) ))8
end do9
end do10
do j=2, N-1
11
do i=2, M-1
12A(i,j) = Anew(i,j)
13
end do14
end do15
if (mod(iter,100)==0) write(*,*) iter, error
16
iter = iter + 1
17 end do
計算の内容は、隣接4点の平均を各点の次の値とし、その値と現時点の値との差分が許容値を下回るまで、同じ計算を繰り返すというシンプルなものである。このように収束的に解を求める計算方法は、科学技術計
算では典型的な方法であり、このサンプルコードを並列化する過程で習得できる OpenACC の知識・ノウハウ
は、大規模なアプリケーションを OpenACC で加速する
ときに大いに役に立つだろう。
■参考文献 :
[1]
[2]
[3]
[4]
[5]
[6]
[7]

（28-6）

Top500 Supercomputer Sites, www.top500.org
CORAL: Collaboration of Oak Ridge, Argonne, and Lawrence
Livermore Labs, asc.llnl.gov/CORAL/Intel Xeon, ja.wikipedia.org/wiki/Xeon
SSE (Streaming SIMD Extensions), ja.wikipedia.org/wiki/
Streaming_SIMD_Extensions
SIMT (Single Instruction Multiple Threads), en.wikipedia.org/
wiki/Single_instruction,_multiple_threads
Intel MKL (Math Kernel Library), software.intel.com/en-us/
intel-mkl
NVIDIA cuBLAS, developer.nvidia.com/cublas
Vol.21, No.2 2016

チュートリアル

OpenACC で始める GPU コンピューティング：OpenACC 概要

[8]
[9]
[10]
[11]
[12]

OpenMP, www.openmp.org
OpenACC, www.openacc.org
CUDA, developer.nvidia.com/cuda-toolkit
OpenCL, www.khronos.org/opencl/
PGI Accelerator Compilers, www.pgroup.com/resources/accel.
htm[13] Amdahl's law, en.wikipedia.org/wiki/Amdahl%27s_law[14] GNU gprof, sourceware.org/binutils/docs/gprof/[15] Git, git-scm.com/

計算工学（28-7）

Vol.21, No.2 2016