アクセラレータ技術 WG 成果報告書(WG 活動期間:2009 年 10 月-2012 年 3 月)

2012 年 3 月 26 日
サイエンティフィック・システム研究会
アクセラレータ技術ＷＧ

目次
1. はじめに ........................................................................................................................................ 1
1.1 背景と目的............................................................................................................................... 1
1.2 メンバー .................................................................................................................................. 2
1.3 活動経緯 .................................................................................................................................. 3
2. ディレクティブ方式による GPU 利用.......................................................................................... 4
2.1 PGI Accelerator 概要................................................................................................................. 4
2.2 PGI Accelerator によるプログラム GPU 化事例 ....................................................................... 6
2.2.1 事例 1：姫野 BMT ............................................................................................................. 6
2.2.2 事例 2：2D-FDTD.............................................................................................................. 9
2.2.3 事例 3：分子軌道法（ERI-MO）..................................................................................... 10
2.3 まとめ .................................................................................................................................... 12
3 アクセラレータ向け拡張言語による GPU 利用 ........................................................................... 14
3.1 CUDA 概要 .............................................................................................................................. 14
3.2 CUDA によるプログラム最適化事例 ...................................................................................... 15
3.2.1 事例 1:分子軌道法（ERI-MO）....................................................................................... 15
3.2.2 事例 2:Linpack................................................................................................................. 18
3.3 OpenCL 概要 (CUDA との比較).............................................................................................. 23
3.3.1 OpenCL と CUDA の性能比較 ........................................................................................ 24
3.3.2 OpenCL のポータビリティ .............................................................................................. 26
3.4 まとめ .................................................................................................................................... 27
4. ハードウェア動向 ....................................................................................................................... 29
4.1 新旧 nVidia GPU アーキテクチャ比較 ................................................................................... 29
5. おわりに ...................................................................................................................................... 35

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

1. はじめに1.1 背景と目的
スパコンシステムの計算性能は、これまで概ね 3 年で 10 倍のペースで向上してきた。近年はプロセッサの動作周波数を向上させることが技術的に難しくなったため、プロセッサコアの単体性能を向上させるのではなく、スパコンを構成するノード数を増やす、プロセッサのコア数を増やす（マルチコア化）など、数を増やすアプローチで性能向上のペースを維持してきた。その中で、近年注目されているのがアクセラレータである。
汎用プロセッサは、HPC 以外の様々な用途での使用を想定しており、様々なプログラムで簡単に安定した性能が得られる様に実装面積の半分以上をキャッシュメモリとして使用しており、演算ユニットの占める面積比率は低い。それに対して、アクセラレータは HPC 用途にフォーカスしたプロセッサで、キャッシュメモリなどの非演算ユニットのリソース量を最小限に抑え、多数の演算ユニットを搭載して、演算ユニットの面積比率を高めることで、演算性能を高めるアプローチのプロセッサである。アクセラレータは高い演算性能を持つ反面、プログラム開発が難しいなど、その性能を引き出すのは簡単でない。実際に、特定アプリでは高い性能が得られるが、それ以外のアプリでは高い性能が得られないという問題がある。しかし、エクサスケール世代では飛躍的に電力あたり演算性能を高める必要があるため、汎用プロセッサと比べて電力あたり演算性能の優れたアクセラレータ技術は、エクサスケール世代の必須技術であると期待されている。
このような背景の下、SS 研アクセラレータ技術 WG を立ち上げ、アクセラレータとして近年普及しつつある GPU を事例として、将来のアクセラレータのあるべき姿を模索・議論した。特に、GPU に代表されるアクセラレータ上でのプログラム開発の現状とその課題、また、どうすればアクセラレータで高い性能を得られるのか、具体的な事例を用いて調査した結果を報告する。本報告は2 年にわたる WG 活動での議論・調査結果をまとめたものである。

-1-

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

1.2 メンバー
氏名

所属

担当幹事

村上 和彰

九州大学

推進委員

井上 弘士 ※

九州大学

青木 尊之

東京工業大学

遠藤 敏夫

東京工業大学黒川 原佳

理化学研究所

滝沢 寛之

東北大学

伊野 文彦

大阪大学

本田 宏明

九州先端科学技術研究所

久門 耕一 ※(株)富士通研究所

堀田 耕一郎富士通(株) 次世代テクニカルコンピューティング開発本部

丸山 拓巳

富士通(株) エンタプライズサーバ事業本部

坂口 吉生

富士通(株) TC ソリューション事業本部

佐々木 啓

富士通(株) TC ソリューション事業本部

成瀬 彰(株)富士通研究所

※まとめ役

-2-

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

1.3 活動経緯
会合

日時

場所

活動内容

第一回

2009 年

富士通本社

 WG 全体スケジュールの議論
 今後の議題検討

10 月 29 日(木)14:30～16:30
第二回

 会員報告 (PGI アクセラレータコンパイラ、2010 年

蒲田

1 月 19 日(火)ソリューション

コーンビーム再構成の高速化、初心者による14:30～17:00

スクエア

GPU プログラミング事例、複合型計算機向け
ソフトウェア開発環境)
 富士通報告 (GPGPU プログラミングの状況、
GPUSIM から見た GPU) 添付資料：1～7

第三回

2010 年

 会員報告 (Linpack on GPU 搭載スパコン

富士通本社

4 月 7 日(水)Tsubame) 富士通報告 (OpenCL の評価)

14:30～17:10

 OpenCL に関する議論
 添付資料：8～9
第四回

第五回

2010 年

 会員報告 (PGI コンパイラ評価)

富士通本社

7 月 14 日(水) 富士通報告 (nVidia 新 GPU 評価)

14:30～17:10

 添付資料：10～11

2010 年

 会員報告 (アクセラレータの大規模システム

富士通本社

12 月 3 日(金)導入課題、Linpack on Tsubame2)
 添付資料：12～13

14:30～17:00
第六回

2011 年

 会員報告 (PGI コンパイラ評価)

富士通本社

 富士通報告 (分子軌道法プログラムの GPU

3 月 25 日(金)14:30～17:00

化検討)
 添付資料：14～15

第七回

2011 年

 会員報告 (CUDA Fortran 評価)

富士通本社

 富士通報告 (分子軌道法プログラムの GPU

8 月 19 日(金)14:30～17:00

化結果)
 添付資料：16～17

第八回

2012 年

トラステッド1 月 19 日(木)クラウド

14:30～17:00

スクエア

 まとめ

-3-

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

2. ディレクティブ方式による GPU 利用
ここでは、汎用言語で記述されたプログラムに、ディレクティブを挿入して GPU を利用する形
態の評価結果を報告する。
汎用言語を用いて記述されたプログラムを GPU で動作させる一般的な方法は、既存の C や
Fortran で記述されたソースコードに対して、開発者が GPU 向けのディレクティブを挿入する方法である。共有メモリ並列環境向けの並列プログラミング言語・環境である OpenMP との類似性が高いプログラミングモデルと言える。コンパイラは、挿入されたディレクティブに基づいて、ソース
コード内のホスト実行部分とアクセラレータ実行部分を識別し、ホスト用オブジェクトとアクセラ
レータ用オブジェクトを生成する。OpenMP と同様、ディレクティブ方式の言語・開発環境では、GPU 向けにソースコードを大幅に変更する必要がないので、ソースコードが表現しているアルゴリ
ズムの可読性を保つことができ、アルゴリズムを変更するときの開発者の手間が少ないメリットがある。その反面、後節で説明する CUDA による GPU 化と比べ、性能は低くなるといったデメリットがある。さらに、汎用言語によるオリジナルコードの記述方法によっては、ディレクティブ方式の言語面
からの制約により、アクセラレータ実行用のオブジェクト生成が行われない場合もある。

2.1 PGI Accelerator 概要
ここでは、nVidia GPU 向けのディレクティブ挿入方式のプログラム開発言語・環境として、本
WG 期間に利用可能であった PGI Accelerator の評価結果を報告する。
PGI Accelerator は、PGI 社が販売している nVidia GPU に対応したディレクティブベースのコン
パイラであり、下記のディレクティブを挿入して GPU 実行領域を指定することで、汎用 CPU 向けに開発されたプログラムを GPU 上で実行することができる。また、GPU コンピューティングで問題となる Host・GPU 間のデータ転送も、適切にディレクティブを挿入することで制御できる。以
下に PGI Accelerator の主要ディレクティブを示す。
Accelerator Compute Region Directive:
GPU 実行部分を指定するディレクティブ。
このディレクティブだけで既存のプログラムを GPU 上で実行させることができる。コン
パイラの判定により、ある程度のデータ転送は自動で行われるが、性能を求めるのであれ
ば、人手で細かく制御する必要がある。
!$acc region [clause [, clause] ..]structured block!$acc end region-4-サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

Clause:copy(list):region の入口で CPU→GPU のコピー、
出口で GPU→CPU のコピーを行なうデータを指定

copyin(list):

入口で CPU→GPU のコピーを行なうデータを指定

copyout(list):

出口で GPU→CPU のコピーを行なうデータを指定

local(list):region 内で local なデータを指定

Accelerator Loop Mapping Directive:
ループをどのように並列化するかを指定するディレクティブ。
!$acc do [ clause [,clause] … ]do loopclause: (指定しない場合はコンパイラが自動選択)



host：

host 上で逐次に実行する

seq：

accelarator 上で逐次に実行する

unroll：unroll を指示，accelerator 上で逐次で実行する

parallel：parallel mode で並列化する

vector：vector mode で並列化する

cache：

メモリ階層の上位に移動することを試す

independent：

ループ内のデータに依存関係がないことを示す

kernel：ループボディが計算カーネルであることを示す

private：

デバイス上のループの 1 反復で 1 回しかコピーされない

shortloop：

accelarator に十分なループかどうかを指示

Accelerator Data Region Directive:
GPU でデータを保持する領域を設定するディレクティブ。
!$accdata region [clause [, clause] … ]structured block!$acc

end data region

clause:copy：

region の入口で CPU→GPU のコピー、
出口で GPU→CPU のコピーを行なうデータを指定-5-

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

copyin：

入口で CPU→GPU のコピーを行なうデータを指定

copyout：

出口で GPU→CPU のコピーを行なうデータを指定

local：

region 内で local なデータを指定

mirror：

Fortran だけで有効．領域内で動的に領域の allocation，deallocation が可能

update device/host：region，data region の外部でコピー

update device：

計算実行前か入口で CPU→GPU にコピー

update host：

計算実行後か出口で GPU→CPU にコピー

2.2 PGI Accelerator によるプログラム GPU 化事例
ここでは PGI Accelerator を用いて実際にプログラムを GPU 化した事例を報告する。事例として用いたプログラムは、姫野 BMT、2D-FDTD、分子軌道法、の 3 種類である。姫野 BMT は、開発者
Ａが GPU 化を行い、2D-FDTD と分子軌道法は開発者 B が GPU 化を行った。各開発者のプログラミ
ングスキルは以下の通りである。

開発者Ａ

開発者Ｂ

開発者Ｃ

OpenMP ・ MPI に よ る 並 列 プ ロ グ ラ ム 開 発 経 験 は あ る が 、 PGI
Accelerator は使用したことがない。
計算科学分野のアプリケーション開発経験は豊富。逐次プログラム開発が主体であり、プログラムの並列化はあまり詳しくない。
並列処理のエキスパートで、CUDA・OpenCL でのプログラム開発経験がある。

2.2.1 事例 1：姫野 BMT
ここでは、PGI Accelerator を用いて姫野 BMT を GPU 化した結果を報告する。
プログラムの GPU 化

開発者Ａ

GPU 種類

Tesla C1060

PGI Accelerator

Version 10.0

PGI Accelerator 使用経験のない開発者Ａが、Fortran 版の姫野 BMT を PGI Accelerator (version10.0)で GPU 化したところ、Tesla C1060 で 19.2GFLOPS の性能が得られた。開発者Ｃが開発した
CUDA 版姫野 BMT では、同じ GPU で 50 GFLOPS 程度の性能を達成可能であり、それに比べて 40%程度の性能しか得られないという結果となった。

-6-

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

開発の容易さの観点では、開発に要した時間は 1～2 時間程度であり（PGI Accelerator の directive
学習時間を含めると数日）、また、19.2 GFLOPS を達成したソースに挿入したディレクティブは 8行であった（ディレクティブ挿入後のコードは以下を参照）。これだけ簡単に最高性能の 40%程度の性能が得られており、開発時間あたりの効率は良いと言える。
GFLOPS

Directive 数

14.9

2行

19.2

8行

+30 分

CUDA 版 (by 開発者Ｃ)

50

…

数日

(参考：Xeon X5670 2.93GHz 1 コア)

3.9…

(参考：Xeon X5670 2.93GHz 4 コア)

7.8…

PGI Accelerator 版
Compute Region Directive のみ
PGI Accelerator 版
Loop Mapping Directive 使用

開発工数
1 時間
(学習に数日)[ディレクティブ挿入後の姫野 BMT ソースコード]!$acc region local(wrk2)
do loop=1,nn
gosa= 0.0!$acc do parallel(3)do k=2,kmax-1!$acc do parallel(3)do j=2,jmax-1!$acc do vector(256)do i=2,imax-1
s0=a(I,J,K,1)*p(I+1,J,K)+…(省略)…
ss=(s0*a(I,J,K,4)-p(I,J,K))*bnd(I,J,K)
GOSA=GOSA+SS*SS
wrk2(I,J,K)=p(I,J,K)+OMEGA*SS
enddo
enddo
enddo
!$acc do parallel(16)do k=2,kmax-1!$acc do parallel(16)do j=2,jmax-1!$acc do vector(256)-7-サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

do i=2,imax-1
p(I,j,k)= wrk2(i,j,k)
end doend doend doenddo!$acc end region
開発者Ｃによる CUDA 実装版の性能に及ばない理由の一つは、並列化方法を選択するときの自由度が少ないことである。PGI Accelerator(version 10.0)では、Loop Mapping Directive により、ループ単位で並列化方法（block 並列 or thread 並列）と並列数を指示・設定することができるが、多重ループのときに（姫野 BMT は３重ループ）、複数ループを融合した並列化を指示・選択することができない。例えば、上記のコード例では、同一 block 内のスレッドは１次元上（線上）に配置
されるが、姫野 BMT を GPU 化するときには、同一 block 内のスレッドを 2 次元的に、平面上に配置すると、メモリアクセスパターンがより規則的になり、高い性能が得られることが分かっている。

姫野 BMT におけるブロック形状(左：PGI アクセラレータで選択した形状、右：高い性能が得られる形状)
このように並列化の方法（ブロック形状）を柔軟に指示・選択することができれば、ディレク
ティブ方式でも更に高い性能が得られると期待される。ただ、記述の柔軟性を高めると、得てして指示方法が複雑になってしまう。例えば、GUI を用いてブロック形状を直感的に指示できるようにするなど、並列化方法を柔軟かつ簡単に指示できる方法が求められる。
なお、PGI Accelerator はオプションで CUDA コードを出力することが可能であり、その CUDA
コードをベースとして人手でチューニングを進めると、トータルとしての開発効率を高められるのではないかとの期待もあった。PGI Accelerator が出力する CUDA コードは行数が多く、可読性に問題があった。現状では、人手でチューニングをするのであれば、最初から人手で CUDA コードを既-8-

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書述した方が良い。
2.2.2 事例 2：2D-FDTD
ここでは、PGI Accelerator を用いて 2D-FDTD プログラムを GPU 化した結果を報告する。
プログラムの GPU 化

開発者 B

GPU 種類

Tesla M2050

PGI Accelerator

Version 11.8

GPU の ハ ー ド ウ ェ ア 特 性 に 詳 し く な い 開 発 者 Ｂ が 2D-FDTD プ ロ グ ラ ム を PGI
Accelerator(version 11.8)で GPU 化したところ、Tesla M2050 で 18.5 GFLOPS の性能が得られた。
CPU 1 コアの性能（Xeon X5650 2.66GHz 1 コア）は 7 GFLOPS 程度であり、CPU 1 コア性能を上回ることはできたが、CUDA 実装版の性能、45.5 GFLOPS には及ばず、CUDA 実装版の 40%程度の性能にとどまった。性能が CUDA 実装版より低い理由の一つは、姫野 BMT 同様、並列化手法を柔軟に指示できないためと考えられる。
GFLOPSPGI Accelerator 版

18.5

CUDA 版

45.5(参考：Xeon X5650 2.66GHz 1 コア)

7.0[ディレクティブ挿入後の 2D-FDTD ソースコード]!$acc region&!$acc copy (Ez ( 0:NMESH-1 ))&!$acc copy (Hx ( 0:NMESH-1 ))&!$acc copy (Hy ( 0:NMESH-1 ))&!$acc copyin (cEx, cEy, CHx, cHy)!$acc do host
do it = 0, N -1
do i = 1, NMESHX-1-1!$acc do vector (256)do j = 0, NMESHY-1-1
Hx ( j, i ) = Hx ( j, i )–-cHy * ( Ez ( j + 1, i )–-Ez ( j, i ) )-9-サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

enddo
enddo
do i = 0, NMESHX-1-1!$acc do vector (256)do j = 1, NMESHY-1-1
Hy ( j, i ) = Hy ( j, i ) + CHx * ( Ez ( j, i + 1)–-Ez ( j, i ) )
enddo
enddo
do i = 1, NMESHX-1-1!$acc do vector (256)do j = 1, NMESHY-1-1
Ez ( j, i ) = Ez ( j, i ) + cEx * ( Hy ( j, i )–-Hy ( j, i -1 ) ) &-cEy * ( Hx ( j, i )–-Hx ( j -1, i ) )
enddo
enddo
enddo
!$acc end region

2.2.3 事例 3：分子軌道法（ERI-MO）
ここでは、PGI Accelerator による分子軌道法プログラムの GPU 化の結果を報告する。

- 10 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

分子軌道法プログラムは、2 電子積分計算を含む、G 行列計算と呼ばれる部分が実行時間の 99%以上を占めており、GPU で処理時間を短縮するには、G 行列計算を GPU 化する必要がある。G 行
列計算は、各種積分型に対応した 6 種類の G 行列計算から構成されている。各積分型 G 行列計算
は、それぞれ４重ループ構成の積分駆動型アルゴリズムである。G 行列計算は ERI 計算のコストが- 11 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書大きいため、ある一組の ERI を計算したら (ERI の要素数は、積分型により 1～81 と異なる)、そのERI が寄与する G 行列要素のみ更新することで、ERI 計算回数を削減している。その結果、G 行列
要素の更新箇所が不連続となり、メモリアクセスパターンが不規則となる。更に、最内ループに計算値に依存する条件分岐がある。
事例１の姫野 BMT、事例２の 2D-FDTD は、いわゆるステンシルアクセス系のプログラムで GPU
実行に向いたプログラムであったが、本事例の分子軌道法プログラムは、GPU 向きとは言い難いプ
ログラムである。
プログラムの GPU 化

開発者 B

GPU 種類

Tesla C1060

PGI Accelerator

Version 10.6

開発者 B が PGI Accelerator(version 10.6)で分子軌道法プログラムの GPU 化を試みたが、簡潔に記述できるアプローチは、図の内側２重ループの GPU 化であった。しかし、その方法で GPU 化した場合
でも、ループボディ部におけるポインタ演算等により、アクセラレータ実行用オブジェクトコードが生成されない事が分かった。本プログラムを GPU 上で高速に実行するには、ソースコードを大幅に変更する必要があり、ディレクティブ挿入による GPU 化（高速化）が非常に困難であることが判明した。

本事例は、現状のディレクティブ方式によるプログラム GPU 化の限界・問題を示している。ディレクティブ方式のメリットは、ソースコードの可読性を損ねることなく、簡単にプログラムをGPU 化できることであるが、分子軌道法プログラムのような複雑なプログラムでは、現状のディレ
クティブ方式では、プログラムを GPU 向けに高速化することができない。このような複雑なプロ
グラムでも、ソースコードを修正することなく、ディレクティブ挿入で並列化を指示する方法を確立しないと、ディレクティブ方式の適用可能領域は、姫野 BMT や FDTD などのメモリアクセスパ
ターンが規則的で、条件分岐が少ないプログラムに限定されることになる。
なお、分子軌道法プログラムのソースコードを大幅に変更して CUDA 化した事例に関しては、別章にて報告する。
2.3 まとめ
本章では、ディレクティブ方式による GPU 利用に関して説明した。言語・開発環境にはワーキ
ンググループ開催期間に使用可能であった PGI Accelerator を用い、3 種類の事例、姫野 BMT、
2D-FDTD、分子軌道法を GPU 化することで、ディレクティブ挿入方式の長所・短所を明らかにした。

- 12 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

GPU 向きである姫野 BMT と 2D-FDTD に関しては、どちらもプログラムを短時間で GPU 化することができ、また、CPU 実行と比べて高い性能が達成可能であることを確認した。GPU 向きアプリ
に関しては、ディレクティブ方式は有用であると言える。一方、GPU 専用言語である CUDA で記述した場合の性能と比較すると、どちらの事例でも性能は 4 割程度に止まった。GPU に向いたアプリ
ケーションであっても、少なくとも現時点ではディレクティブ方式で GPU 性能を最大限に生かす
ことは難しいことが明らかになった。
GPU に不向きである分子軌道法に関しては、現行のディレクティブ方式では、分子軌道法のような複雑な問題には容易には対処できないことが明らかになった。

今回の評価では、非常に GPU 向いている事例と、非常に GPU に不向きな事例、ある意味で両極端な事例を扱ったが、実際には、明確には GPU 向き・不向きと判別できないアプリが多いと思われる。今後は、そのような中間的な事例で、ディレクティブ方式の有効性を確認する必要がある。

- 13 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

3 アクセラレータ向け拡張言語による GPU 利用
ここでは、アクセラレータ向け拡張言語でプログラムを記述して GPU を利用する形態の評価結果を報告する。
アクセラレータ向け拡張言語とは、例えば CUDA や OpenCL のことであり、C/Fortran 等の既存言語を拡張することで、アクセラレータ向け処理をオフロードするプログラムを比較的容易に記述できるようになっている。ディレクティブ方式による GPU 利用の OpenMP との類似性と対比すると、CUDA や OpenCL は分散並列環境向けの並列プログラミング言語・環境である MPI との類似性の高いプログラミングモデルとも言える。
この開発環境では、既存の C や Fortran で記述されたソースコードから、開発者がアクセラレ
ータで実行させる部分を選択し、その部分のコード(一般に kernel と呼ばれる)をアクセラレータ向けに記述する。また、ホスト実行部分には、この kernel を呼び出すコードを記述する。コンパイラ
は、ソースコードに記述された内容に従って、ホスト用オブジェクトとアクセラレータ用オブジェ
クトを生成する。CUDA や OpenCL で開発されたプログラムは、MPI と同じように、適切に記述されれば、アクセラレータ性能をフルに引き出すことができるメリットがある。その反面、逐次向けに記述されたソースコードとは記述内容・方式が大きく変わるため、汎用 CPU 向けコードとアクセ
ラレータ向けコードの共通化が難しく、ソースのメンテナンスやポータビリティ面でのデメリットがある。
3.1 CUDA 概要
CUDA は、nVidia が自社 GPU 向けに提供しているプログラミング言語・開発環境である。CUDA は
C/C++をベースとして GPU 向けに言語仕様を拡張しているが、仕様拡張は最小限に抑えられており、C/C++でのプログラム経験のある開発者であれば、それほど違和感無く CUDA を利用することができる。現在は Fortran 版の CUDA の利用可能であり、Fortran 言語をベースとした GPU 利用も進んでいる。CUDA 以前は、グラフィックス処理用の言語である OpenGL を用いることが一般的であったが、OpenGL によるプログラム開発は制約が多く、CUDA の出現により GPU を用いた汎用計算の垣根が大きく下がったことは疑いようのない事実である。
GPU は、演算ユニット単体の性能は低いが、それを大量に搭載することで高い性能を狙うアー
キテクチャであり、GPU の能力を最大限引き出すにはプログラムを並列化する必要がある。しかし、
CUDA 環境では、プログラムの自動並列化機能は提供されない。従って、既存のプログラムを、CUDA
を使って GPU 化する時には、開発者が明示的にプログラムを並列化する必要がある。また、CUDAの並列化モデルは、block 並列と thread 並列から構成される 2 階層の並列モデルとなっている。この 2 階層の並列モデルは nVidia GPU の内部構造(Stream Multi-Processor と CUDA core)を反映したものであり、nVidia GPU 向けにプログラムを開発するモデルとしては適切なものである。
- 14 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

しかし、そもそもプログラムの並列化は困難な作業であるのに、1 階層でなく、2 階層で並列化を考慮する必要があることは、特に、並列処理に不慣れなアプリ開発者が、CUDA でプログラムを開発する上での障壁となっている。また、nVidia GPU の内部構造をある程度理解しないと、性能面で適切なプログラムを開発するのが難しいという問題がある。そして、頻繁に CUDA の version が上がるため、その影響で従来の CUDA コードが動かない、挙動が変わる等の問題がある。更に、CUDA
は nVidia proprietary な言語・開発環境でありポータビリティがないため、プログラムを CUDA 化してしまうと他プラットフォームへのポーティングが難しくなるという問題もあり、ISV アプリのGPU 対応が進まない要因の一つとなっている。
CUDA の問題


開発者がプログラムを明示的に並列化（手動並列）



2 階層の並列化の考慮が必要（block 並列と thread 並列）

マルチ GPU・GPU クラスタ環境では、3 階層の並列化の考慮が必要

頻繁な version up と version 間の非互換性

nVidia proprietary な言語・開発環境（nVidia がなくなったらどうするのか）

3.2 CUDA によるプログラム最適化事例
ここでは CUDA を用いてプログラムを GPU 化した事例を報告する。事例としたプログラムは、
分子軌道法、Linpack の 2 種類である。なお、Linpack は Tsubame 2.0 での Linpack 性能評価事例である。
3.2.1 事例 1:分子軌道法（ERI-MO）
ここでは CUDA プログラム開発経験のある開発者Ｃが分子軌道法を CUDA 化した事例を説明する。分子軌道法プログラムを PGI アクセラレータで GPU 化した事例では、アクセラレータ実行用
オブジェクトを生成できなかったが、GPU の高メモリバンド幅特性を生かす並列化手法を考案・採用し、G 行列計算の全てを CUDA 化することで、汎用 CPU(1 コア)に対して最大 29.2 倍の高速化を達成している。
プログラムの GPU 化

開発者 C

GPU 種類

Tesla C2050

CUDA version

4.0

下記の図は、(ss,ss)型積分の G 行列計算の概要である。

- 15 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

naïve な GPU 化は、1st & 2nd ループの各イタレーションを 1 つのスレッドに割り当てる方法である。この方法では、各スレッドは、一組の ERI 計算を行い、その ERI が寄与する G 行列を更新(部分 G 行列更新)する。この方法の問題は、部分 G 行列更新の部分である。具体的には、G 行列の更新箇所が不連続でメモリアクセスパターンが不規則になること、そして、他スレッドと同時に同一箇所を更新する可能性があり、その場合でも正しい値が得られるようにするための排他制御(atomic 操作)が必要になることである。
この naïve な手法で(ss,ss)型積分 G 行列計算を GPU 化したところ、GPU 化した部分だけの性能比較で、CPU(1 コア)に対して 10 倍弱の性能向上しか得られなかった(入力データは sw03-sto)。
(ss,ss)型は、6 種類の G 行列計算の中で、最もシンプルで GPU 化に向いた積分型である。他積分型
ではこれより性能向上率が大幅に下がると思われるため、naïve な手法で 6 種類の積分型全てを GPU化すると、CPU(1 コア)に対して数倍程度の性能向上しか得られないことになる。GPU で高い性能を達成するには、部分 G 行列更新の処理をなんらかの手法で高速化する必要がある。
そこで、GPU 向け並列化手法として考案したのが、以下に示す kernel 分離方式である。

- 16 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

通常は ERI 計算を行うと、その ERI が寄与する G 行列要素をすぐに更新するが、kernel 分離方式では ERI 計算と部分 G 行列更新を別 kernel に分離する。1 番目の kernel では、各スレッドが一組の ERI 計算を行い、その結果をメモリに格納する。2 番目の kernel では、メモリに格納された
ERI を読み出して、その ERI が寄与する G 行列部分を更新する。ただ、各スレッドが一組の ERI を担当したのでは、G 行列へのメモリアクセスは不規則なままで、排他制御も必要となり性能は向上しない。そこで、2 番目の kernel では、各スレッドに対して更新する G 行列要素を割り当てること
で、G 行列のメモリアクセスを規則化し、排他制御を不要にした。この kernel 分離方式では、ERI計算結果のメモリへの読み書きが必要になるデメリットがあるが、これは GPU の高メモリバンド
幅特性で相殺できるだろうという考えである。
この kernel 分離方式で(ss,ss)型積分 G 行列計算を GPU 化したところ、GPU 化した部分だけの性能比較で、naïve な方式と比較して 4 倍の性能、CPU(1 コア)と比較して 40 倍の性能向上が得られることが確認できた。
(ss,ss)型積分 G 行列計算
１回の実行時間 (ミリ秒)

性能向上(CPU 1 コア)

444

(基準)

GPU naïve 方式

45

x9.9

GPU kernel 分離方式

11

x40.2

(*) 入力データは sw03-sto を使用(下記参照)- 17 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

6 種類の積分型 G 行列計算を全て kernel 分離方式で GPU 化し、プログラム全体の実行時間を、様々な入力データで測定した結果を以下に示す。CPU(1 コア)と比較して、プログラム全体の実行性能が、最大で 29.2 倍向上することが確認できた(入力データとして sw04-sto を使用した場合)。
CPU(1 コア)

入力データ
名前

原子数

sw01-sto3

sw02-sto

9

sw03-sto24

sw04-sto

基底関数

GPU
対 CPU

実行時間(秒)

実行時間(秒)0.0310

0.165

x0.19

0.451

0.248

x1.82

26.3

1.19

x22.1

42

254

8.69

x29.2

sw01-631

30.0590

0.272

x0.22

sw02-631

9

2.72

0.638

x4.26

sw03-631

24

155

13.8

x11.2

STO-3G

6-31G

性能向上

入力データのサイズ(原子数)が小さいケースでは、CPU でも実行時間が高々数秒程度と短いため、GPU による性能向上率は僅かであった(特に原子数が少ない sw01-*では、GPU の方が遅かった)。
これは、CPU・GPU 間のデータ転送時間や、GPU 起動時間がボトルネックになったためと思われ、GPU を活用するには、ある程度以上のサイズの問題を解かせる必要があると言える。基底関数に関しては、STO-3G と比べると、6-31G の場合の GPU による性能向上率はやや低めとなった。これは、STO-3G は最内ループ数がほぼ一定であるのに対して、6-31G は最内ループ数が 1～81 の間で可変であり、スレッド毎の処理量に大きなアンバランスが生じるためである。このような負荷の不均一性は、GPU に限らず、多並列で性能をかせぐアクセラレータ全般の課題である。
3.2.2 事例 2:Linpack
ここでは、東工大 TSUBAME2.0 での Linpack 最適化事例を紹介する。
TSUBAME2.0 は Tesla M2050 を 4200 枚以上備えたアクセラレータ・スパコンで、1.192pFlopsの Linpack 性能を達成している。多数のアクセラレータを効率良く動作させるために Linpack 実装
を大幅に改良している。アクセラレータ・スパコンでは、演算性能の向上に対して、相対的にネッ
トワーク性能が悪くなるため、その対策として、通信コストを隠蔽する技術がきめ細かく適用されている。今後もアクセラレータ化とメニーコア化の進展で、演算性能と通信性能の乖離が進むと考えられ、本事例は、この問題に対処した先行事例の一つである。

- 18 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

4000GPU 超を備えるスーパーコンピュータ TSUBAME 2.0 のLinpack による性能評価
遠藤敏夫 (東京工業大学学術国際情報センター
1.

はじめに
GPU を中心とする数十～数百コアのメニーコアに基づくアクセラレータの，科学技術計算への応用の可能性は早くから期待されていた．しかしアクセラレータを数百以上結合してトータ
ルで高性能が得られるか否かは，東京工業大学が TSUBAME1.0 スパコンの導入を行うまで明らかではなかった．筆者らは 2006 年に，TSUBAME1.0 の 648 計算ノードと 360 枚のClearSpeed アクセラレータを効率的に用いるような，密行列計算ベンチマークである Linpackの新たなアルゴリズム・最適化手法を考案した[1]．それにより CPU を用いる場合の 25%増に相当する 47tFlops を達成し，アクセラレータのスパコン規模でのスケーラビリティを示した．
2010 年にはより最新の，プログラマビリティにおいても演算速度・メモリバンド幅においても優れたアクセラレータである NVIDIA Tesla M2050 GPU アクセラレータを 4200 枚以上備えた，TSUBAME2.0 スパコンが導入された．我々は上記の Linpack 実装を大幅に改良し，
TSUBAME2.0 で効率的な実行を実現した．得られた性能は 1.192pFlops であり，国内で初めて 1pFlops を超え，2010 年 11 月の Top500 ランキングでは世界 4 位を達成した．さらにアクセラレータを効率的に利用することにより電力あたりの性能という指標についても大幅な利点があることを実証した．Linpack 実行時の電力性能比は 958mFlops/Watt であり，2010 年 11
月の Green500 ランキングにおいて世界 2 位を獲得した[2]．
TSUBAME1.0/2.0 に渡って高効率を達成するために対処が必要であった課題は，演算性能と比較したときの相対的な通信バンド幅の悪化である．無論，インターコネクト・ノード内バス・メモリ技術の向上はあるものの，それ以上にアクセラレータの演算性能の伸びが大きく，その性能を活用するには通信コストの隠ぺいが必須であった．この問題は将来のエクサスケール・
スパコンにおいてもさらに悪化する本質的な課題である．
2.

TSUBAME2.0 スーパーコンピュータ
TSUBAME 2.0 は，1400 ノード以上の計算ノード(合計理論演算性能 2.4pFlops)と，合計7.1pBytes のストレージが QDR InifiniBand により接続されている，ペタスケールのスーパーコンピュータである．以下では本稿に関連の深い事項のみ記す．
各計算ノード Hewlett-Packard Proliant SL390s G7 (図 1)は 6 コアの Intel Xeon X5670
2.93GHz プロセッサを二つ，DDR3 メモリ 54GB，および NVIDIA Tesla M2050 GPU を三つ搭載する．各 GPU はストリーミングマルチプロセッサ(SM) を 14 基持ち，各 SM は SIMD 動
作する CUDA core を 32 基持つ．各 GPU の理論演算性能は 515gFlops である．GPU はホス
トメモリとは独立したメモリを 3GB 持ち，GPU 演算に必要なデータはそちらに置く必要がある．ホスト側からの通信経路として，三つの GPU はそれぞれ PCI-Express 2.0 x16 にて接続されている．
- 19 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

図 1: TSUBAME2.0 計算ノードの内部構成

全ノードを接続するインターコネクトは 2 段のスイッチから成るファットツリーであり，フルバイセクション構成である．Dual rail 構成であり，各 rail がファットツリーを構成する．
3.

Linpack ベンチマーク
Linpack は密行列連立一次方程式をガウス消去法を用いて解くベンチマークであり，よく知られた MPI 並列実装として High-performance Linpack (HPL) [3]がある．我々の実装は HPLを大幅に改造したものである．
HPL においては，一辺の長さ N の巨大な密正方行列(TSUBAME2.0 においては N=240 万程
度のものを用いた)を，二次元ブロックサイクリック分割にて，各 MPI プロセスに分散配置する．分散の単位であるブロックのサイズを B とする．計算全体での演算量は O(N3)，通信量はO(N2)であり，演算のほとんどは密行列積(DGEMM)カーネルとなっている．詳細は[2,3]を参照
されたい．
このようなベンチマークの性質から，古典的な(CPU のみの)クラスタシステムで高性能を得るには，以下のような方策をとる．


高性能な DGEMM カーネルを持つ BLAS ライブラリを用いる．



通信量を相対的に軽減するため，行列がシステム全体のメモリに収まる範囲で，N をなるべく大きくする．

4.

GPU を用いる際の課題と解決
HPL から GPU を利用するのに最も単純な手法は，DGEMM カーネルを GPU 上のものに置き換えることである．
図 2: TSUBAME2.0 の 8 ノードを用いた Linpack 性能．

図 2 の”Naïve”は，DGEMM カーネルを単純に GPU 版に置きかえた場合の性能を示す．実験
条件は以下のとおりである：8 ノード 24GPU が用いられている．今回は NVIDIA 社内部の，約 350gFlops の性能を持つカーネルの提供を受けたためそれを用いた．これは CUBLAS ライ- 20 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書
ブラリの DGEMM(バージョン 3.2 で約 290gFlops)より高速である．前提として，N をなるべ
く大きくするため(本実験では N=188,416)に行列はメインメモリにおき，DGEMM のたびに
PCI-Express を介し GPU と通信するものとする．
Naïve は GPU の利用により，オリジナル(CPU only)に比べ 5.3 倍の速度向上を得た．しかしながら 24GPU 分の DGEMM 性能である 8.4tFlops に比べると 6 割程度に留まっている．
問題解析のため，HPL の 1 プロセスの 1 ステップのアルゴリズムの概要を図 3 に示す．各プ
ロセスが持つ部分行列を A’とすると，1 ステップで行うべき仕事は，U および L と呼ばれる行
列の一部を MPI 通信によって獲得し，A’=A’-L*U という DGEMM 演算をすることである．な
お最適化のために，L の通信については前のステップのうちに行う実装が，オリジナル HPL ですでになされている．

図 3: オリジナル HPL の 1 ステップのアルゴリズム

これを単純に GPU 化すると，以下の理由で問題があると分かった．
(1) U, L を MPI 通信で取得する間は計算が進まない．この点は通常のクラスタではあまり問題とされて来なかったが，GPU を用いるスパコンにおいては，演算性能に対し相対的な通信性能が低いため，問題が顕在化する．
(2) 次ステップのための L 通信を非同期に行うため，部分的な DGEMM とメッセージ到着チ
ェックを繰り返す．GPU 版においては，分割された DGEMM カーネルのたびに無駄な PCI
通信が繰り返される．
これらを解決するために，以下のような手法を提案・実装した．
(2)の問題の解決を以下のように行った．GPU メモリサイズがホストメモリより小さいため，
行列 A’を分割して，部分的に DGEMM を行うことは必須である．ただし，Naïve と異なり，Lについては GPU メモリに置きっぱなしにし，毎回 L を PCI 通信することを避ける．また部分
A’および部分 U の PCI 通信と DGEMM 演算をパイプライン化させることにより，PCI 通信のコストを隠ぺいする．さらには，次ステップのための L 通信は，GPU が計算している間に CPUにより行われるため，このコストも隠ぺいされる(問題(1)も部分的に解決)．ここまでを行ったものが，図 2 の Ver1 となる．
さらに U についても問題(1)を解決するため Ver2 を実装した(図 4)．ここでは U の MPI 通信の隠ぺいのために，MPI 通信，PCI 通信，DGEMM 演算を全てパイプライン化する方式を実装した．この Ver2 の 1 ステップのアルゴリズムは図 4 のようになる．消費される時間のほとんどで GPU 上の計算を行えており，時間を有効に使う実装となっている(L の PCI 通信が隠ぺいされず残っているが，その隠ぺいも原理的には容易である)．

- 21 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

図 4: GPU 版 HPL(Ver2)の1 ステップのアルゴリズム

結果として，図 2 のグラフに示された通り 8 ノードで 7.0tFlops の性能となり，事実上の限界(8.4tFlops)の 83%を達成した．なお，8 ノードの理論演算性能 13.5tFlops に比べると 52%程度であるが，この乖離の最大の理由は，DGEMM カーネル性能である 350gFlops がすでに理論値 515gFlops から乖離していることであることに留意されたい．
5.

TSUBAME2.0 全体を用いたペタスケールの性能評価
TSUBAME2.0 正式稼働開始の直前の 2010 年 10 月に，システムの安定性試験の一環として，

TSUBAME2.0 全体で，前節の Linpack GPU 版 Ver2 を用いた評価を行った．実験に用いたシステムソフトウェアは，SUSE Linux Enterprise 11, OpenMPI 1.4.2, GCC 4.3，CUDA 3.1 である．最終的に用いたノード数は 1357 であり，GPU 数は 4071 である．
Linpack 実行時のパラメータは以下の通りとした．行列サイズ N=2,490,368 とした．この行列を 1357 ノードに分散すると，ノードあたり約 34GB となる．MPI プロセス数については，1プロセスが 1GPU を利用することとし，ノードあたり 3 プロセス，全体で 4071 プロセスとした．行列はプロセス間で二次元分割されるが，4071=59*69 となるようにした．ブロックサイ
ズ B=1024 とした．
この実行には約 2.4 時間かかり，結果として得られた速度性能は 1.192pFlops となった．これは国内で初めて 1pFlops を超えた計算となった．なお 1357 ノードの理論演算性能
2.288pFlops に対して 52.1%の相対性能となった．理論値との乖離については前節で議論した通りだが，8 ノードでも 1000 ノード超でも似た相対性能が得られており，スケーラビリティは非常に良好であることが分かる．
同時に実行時のシステム電力も測定しており，消費電力は 1.24mWatt (Green500 の規定に基づく電力)，電力性能比は 958mFlops/Watt という，これも世界最高レベルの値が得られた．
6.

おわりに国内で初めて Linpack 1pFlops 超の達成を，GPU アクセラレータを効率的に利用するソフト

ウェア技術により行った．2008 年に世界で初めて 1pFlops を超えた LANL RoadRunner スパ
コンも，PowerXCell 8i アクセラレータを用いていた．将来のスパコンが異種プロセッサ混在のものであるか否かは議論の待たれるところであるが，プロセッサのメニーコア化と，それに伴う演算性能に比べた通信性能の相対的な悪化はまず間違いない．その問題に対処する，アルゴリズム・プログラミングシステムの研究開発が必要である．
[1] T. Endo and S. Matsuoka, “Massive supercomputing coping with heterogeneity of modern accelerators”,- 22 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書
IEEE IPDPS08, 10pages, 2008.
[2] 遠藤 敏夫，額田 彰，松岡 聡． スーパーコンピュータ TSUBAME 2.0 における Linpack 性能 1 ペタフ
ロップス超の達成． 情報処理学会論文誌コンピューティングシステム， Vol. 4, No.4 (ACS 35), pp.169-179，
2011 年 10 月．
[3] A. Petitet, R. C. Whaley, J. Dongarra, and A. Cleary. HPL–- a portable implementation of thehigh-performance Linpack benchmark for distributed-memory computers.
http://www.netlib.org/benchmark/hpl/.

3.3 OpenCL 概要 (CUDA との比較)
OpenCL は、OpenGL などの仕様策定を行っている Chronos を中心に、多数のベンダーが協力して仕様が策定されたプログラミング言語である。OpenCL の仕様策定には、GPU ベンダー以外に、
汎用 CPU・モバイル・組込みベンダーも参加しており、OpenCL でプログラムを開発すれば、様々なプラットフォームでそのプログラムを実行できる、高いポータビリティが期待されている。
現在の OpenCL 言語仕様は、CUDA、特に CUDA の低レベル API との類似性が非常に高い。CUDAと同様、2 階層の並列化モデルを採用している。このモデルは GPU の内部構造にフィットしたものであり、汎用 CPU や組込み CPU にフィットするかどうかは疑問が残る。また、OpenCL のカーネル・
ルーチン記述方法は CUDA との違いがほとんどないため、CUDA 使用経験のある開発者であれば、OpenCL 利用の障壁は低いことはメリットだが、カーネル・ルーチン呼び出し方法が CUDA ほど直感的ではないこと、カーネル・ルーチン呼び出し前の処理が煩雑であるなど、使い勝手の点で問題がある。CUDA は従来と比べてプログラミングを簡単にする方向性であったが、OpenCL はその流れに逆行しているとも言える。また、CUDA は version up で新しい機能が追加されるが、OpenCL
ではそれら新機能を使うことができない点もデメリットである。そのため、特に nVidia GPU 向けのプログラム開発に関しては、OpenCL を採用するメリットは薄い。OpenCL は開発者が直接プログラムを記述する言語ではなく、中間言語として使用されるのではないかという意見もある。
一方で、AMD は自社 GPU 向けのプログラミング言語を、自社独自言語から OpenCL に変更した
こともあり、AMD GPU の今後の状況次第では、ポータビリティがより重視され、OpenCL が使用される可能性もある。


CUDA と OpenCL の違い(主要点のみ)
用語

最小実行単位- 23 -

CUDA

OpenCLThreadWork-itemサイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

最小実行単位のグループ

Thread-block

Work-group

カーネル・ルーチンの実行方法
CUDA: サブルーチン呼び出しに類似
Dim3 block(…);
Dim3 thread(…);
Gpu_kernel<<block,thread>>(…);
OpenCL: コマンドキューにカーネル・ルーチンを登録
clEnqueueNDRangeKernel( queue, gpu_kernel, … );カーネル実行前に必要な処理
CUDA: 少ない
カーネル引数設定 (block と thread の作成)
OpenCL: 多い
プラットフォーム ID 取得
デバイス ID 取得
コンテキスト作成
コマンドキュー作成
メモリオブジェクト作成
カーネルオブジェクト作成
カーネル引数設定
3.3.1 OpenCL と CUDA の性能比較
ここでは nVidia GPU 上で CUDA と OpenCL の性能を比較した結果を報告する。
性能比較用プログラムには 3D-FDTD を用いた。CUDA で開発した 3D-FDTD プログラムをOpenCL に書き換えて性能を比較した。ベースとした CUDA 版プログラムは、CPU（Core-i7 2.66GHz、
1 コア）に対して 20 倍以上高速な実装である（GPU は GTX-285 を使用）
。CUDA 版と OpenCL 版の性能差は以下の通りであった。
CUDA 版と OpenCL 版の性能差

CUDA version 2.3

CUDA 版が 1.8 倍高速

CUDA version 3.0

CUDA 版が 1.02 倍高速 (ほぼ同性能)

CUDA version 2.3 環境では、CUDA 版と OpenCL 版で 1.8 倍程度の性能差が観測された。原理上は CUDA と OpenCL で性能差が生じる理由はない。CUDA version 2.3 の時点では OpenCL 対応の初- 24 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

期段階であり、コンパイラの最適化機能のレベルが CUDA と OpenCL で差異が性能差の理由と思われる。実際、その後にリリースされた CUDA version 3.0 で、CUDA と OpenCL でほぼ同じ性能が得られることが確認できた。技術的には、CUDA と OpenCL で顕著な性能差が生じることはないが、今後どうなるかはベンダーが OpenCL をどこまで真剣にサポートするかに依存することになる。
参考までに、CUDA 版と OpenCL 版、それぞれのカーネル・ルーチンを以下に示す。OpenCL 版の赤字部分が CUDA 版と記述の異なる部分である。カーネル・ルーチンの記述方法に関しては、CUDAと OpenCL に本質的な違いのないことが、これを見ると良く分かる。
[CUDA で記述した 3D-FDTD のカーネル・ルーチン（一部）]__global__ void gpu_calc_Ex( float *Ex, float *Cex, float *Dex,float *Hx,
float *Hy, float *Hz ){int

x = blockIdx.x * blockDim.x + threadIdx.x;

int

y = blockIdx.y * blockDim.y + threadIdx.y;

int

z;

int xs = 0;
int ys = 1;
int zs = 1;
int xe = NUM_X-2;
int ye = NUM_Y-2;
int ze = NUM_Z-2;
int

flag = 0;

if ( (x >= xs) && (x <= xe) && (y >= ys) && (y <= ye) ) {flag =;}for ( z = z ; z <= z ; z++ ) {if ( flag ) {EX(z,y,x) = CEX(z,y,x) * EX(z,y,x)+DEX(z,y,x)*(HZ(z,y,x)–

HZ(z,y-1,x)–

HY(z,y,x)

HY(z-1,y,x) ;}}}[OpenCL で記述した 3D-FDTD のカーネル・ルーチン（一部）]__kernel void gpu_calc_Ex( __global float *Ex, __global float *Cex, __global float *Dex,__global float *Hx, __global float *Hy, __global float *Hz ){- 25 -+

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

int x = get_global_id(0);int y = get_global_id(1);intz;

int xs = 0;
int ys =;int zs =;int xe = NUM_X- ;int ye = NUM_Y- ;int ze = NUM_Z-2;
int

flag = 0;

if ( (x >= xs) && (x <= xe) && (y >= ys) && (y <= ye) ) {flag =;}for ( z = zs; z <= ze; z++ ) {if ( flag ) {EX(z,y,x) = CEX(z,y,x) * EX(z,y,x)+DEX(z,y,x) * (HZ(z,y,x) –HZ(z,y-1,x) –HY(z,y,x) + HY(z-1,y,x) ;}}}3.3.2 OpenCL のポータビリティ
ここでは OpenCL で記述したプログラムのポータビリティの評価結果を報告する。
ポータビリティ調査の対象としたプログラムは 3D-FDTD であり、nVidia GPU 向けに開発したOpenCL 版の 3D-FDTD プログラムが AMD GPU 上で動作するかどうか、また、そのときの性能を測定した。評価環境は以下の通りである。
nVidia GPU 環境

AMD GPU 環境

GPU

nVidia GTX-285

AMD HD-5870

SDK

CUDA Toolkit 2.3

Stream SDK 2.0

CPUCore-i7 2.66GHz

Core-i7 2.66GHzOS

RHEL 5.3

OpenSUSE 11.0

動作する・しないのレベルであれば、nVidia GPU 向けに開発した OpenCL 版 3D-FDTD プログラ
ムは、ソースコード無修正で動作することが確認できた（開発環境が異なるため、Makefile は若干変更）。ただし、ソースコード無修正で動作したのは、ベンダー提供の OpenCL サンプルコードをベースとせず、純粋に OpenCL 仕様に基づいてプログラムを開発したことが大きく寄与している。
ベンダーの SDK 等に含まれるサンプルコードには、デバッグを容易にするため、また、プログラ- 26 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

ム記述性を高めるためのラッパーライブラリが使われていることが多い。当然、ラッパーライブラ
リはベンダー独自でポータビリティがないため、SDK のサンプルコードをベースにプログラムを開発すると、ポータビリティの面で問題が生じるので注意が必要である。
一方、性能に関しては、全くポータブルでないことが分かった。nVidia GPU 向けに開発したソースコードをそのまま無修正で AMD GPU 上で実行した場合、nVidia GPU に対して AMD GPU では
実行時間が 100 倍以上長くなった（CPU 1 コア実行より 10 倍以上低速）。AMD GPU で適切な性能を得るには、AMD GPU 向けにソースコードを変更する必要があり、具体的には以下の修正を実施することで、CPU 1 コア性能に対して 2 倍弱の性能が得られた。
CPU・GPU 種類

ソース種類

実行時間(秒)Core i7 2.66GHz 1 コア

C

79.14

nVidia GTX-285

OpenCL ベース実装

6.36

OpenCL ベース実装

843.30+ ストリーム化

87.92＋ ベクトル化

47.18

AMD HD-5870

AMD GPU 向けのソースコード変更内容
1. ストリーム化（work-item の担当領域と work-group の形状変更）
2. ベクトル化（データ型を float から float4 に変更）
ソースコード変更内容の 1. ストリーム化は、プログラムの記述方式にも依存するが、基本的にはパラメータ変更で対処可能であり、ソースコードを大幅に修正する必要はない。一方、2. ベク
トル化は、データ型をスカラ型からベクトル型に変更するだけでは駄目で、データ型の変更に合わ
せてカーネル・ルーチンを適切に書き変える必要がある。また、このベクトル化は VLIW 型の AMD
GPU アーキテクチャには適切なチューニング手法だが、スカラ型の nVidia GPU アーキテクチャには不適切なものであり、ベクトル化したカーネル・ルーチンを用いると、nVidia GPU では性能が低下する。 現実的な性能を得るには、GPU 毎にカーネル・ルーチンを実装・用意する必要がある。
OpenCL 仕様に基づいてプログラムを開発すれば、アーキテクチャの異なる GPU 間でも動作レ
ベルのポータビリティを保つことはできる。しかし、性能面のポータビリティは、少なくとも現時点では全く期待できないことが明らかになった。
3.4 まとめ
本章では、アクセラレータ向け拡張言語による GPU 利用に関して説明した。
分子軌道法の事例は、GPU に不向きなプログラムを単体 GPU で利用する事例であり、GPU の長- 27 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

所・短所を理解した上で、その特性を生かす並列化手法を見つけることができれば、一見するとGPU に不向きなプログラムであっても、GPU で高い性能を達成できることを示した。今回、考案・実装した並列化手法は、既存の逐次実装に対してディレクティブを挿入する GPU 利用方法では、残念ながら実現することができない（少なくとも現時点では）。しかし、全てのアプリ開発者に対して並列処理のエキスパートになることを求めるのは現実的ではない。どのようにアプリ開発者の並列化作業を支援するのかは、依然として課題である。
Linpack の事例は、今後普及すると考えられるアクセラレータ・クラスタ構成での GPU 利用事例であり、アクセラレータ・クラスタでは、演算性能に対して相対的に通信性能が悪化するため、通信処理を隠蔽する手法が重要になることが明確となった。ただ、現状では、このような通信処理隠蔽の実現方法はプログラム依存であり、また、その実装やデバッグには大きな労力がかかる。今後は、通信処理隠蔽のアルゴリズム開発、また、それを簡単に実装・デバッグできるプログラミン
グ環境の開発が今後の課題である。
OpenCL は、異種 GPU 間(nVidia GPU と AMD GPU 間)で、ソースレベルのポータビリティがあることが確認できたが、性能面でのポータビリティは無く、使い物になる性能を達成するには、GPU
毎にカーネル・ルーチンを開発する必要があることが分かった。

- 28 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

4. ハードウェア動向
ここではアクセラレータのハードウェア動向に関して報告する。
アクセラレータ分野は、競争が激しいこともあり、継続的にアーキテクチャ更新が行われている。nVidia GPU は本 WG 期間中の 2010 年初に、アーキテクチャが Tesla から Fermi に変更になり、倍精度演算性能の向上、ECC メモリサポート、汎用キャッシュの導入など、HPC 用途向けの機能が強化された。AMD GPU も 2011 年末に大幅なアーキテクチャ変更が行われ、演算ユニットは VLIW型から、nVidia GPU に近いスカラ型に変わるなど、より汎用計算に適した構造に変化した。
一方、汎用 CPU をベースにしたアクセラレータ（いわゆる Many-core）には Intel MIC がある。
MIC は x86 コアをシンプルにして多数並べる構造になると予測されるが、詳細情報は依然として公開されていない。
Intel MIC 情報(Knights Corner)
倍精度性能

1TFlops 以上

コア数

50 以上

SIMD 演算512bit 幅

メモリバンド幅?

プログラミング言語

既存言語で OK (既存アプリがそのまま動作)

また、現在のアクセラレータは、いずれも Host マシンを必要とするタイプであるが、消費電力の観点からは Host マシンレスの stand-alone タイプになることが望ましい。アクセラレータのあるべき構造は stand-alone タイプである。nVidia は stand-alone タイプのアクセラレータ実現を目指したプロジェクトを進めており(Project Denver)、AMD は x86 CPU と GPU の統合を目指した
HSA(Heterogeneous Systems Architecture)を推進している。
今後も、アクセラレータのアーキテクチャ進化は続くと思われる。
4.1 新旧 nVidia GPU アーキテクチャ比較
ここでは、1 世代前の Tesla アーキ GPU（GTX-285）と現世代の Fermi アーキ GPU（GTX-480）の比較結果を報告する。汎用計算の観点から見た Tesla から Fermi での大きな変化は、以下の 3 点である。
倍精度(DP)性能向上： DP/SP 比率が 1/12 から 1/2 に向上
信頼性向上：ECC メモリサポート
使い易さ向上：キャッシュ構成変更（Read-only キャッシュから Read-Write キャッシュに）- 29 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

倍精度性能向上と ECC メモリサポートは、従来から HPC 分野から必須機能として要求されたいたものであるが、Read-Write キャッシュのサポートが GPU プログラムの性能チューニングの簡易化に非常に効果的であることを、姫野 BMT を事例とした評価で確認できた。
Tesla アーキ世代ではキャッシュ用途が限定されていたため、共有メモリ（いわゆるスクラッチパッド）をキャッシュの代替リソースとして利用して、グロバール・メモリへのアクセス回数を削減することが、GPU で高い性能を実現するための必須技術であった。共有メモリはプログラム側か
ら明示的に制御するリソースであり、ソースコードを修正して利用する必要があるが、共有メモリを利用するコードは可読性が低下するため、デバッグやソースコードのメンテナンスが悪化するという問題がある。また、メモリアクセスパターンが規則的なプログラムや、どのデータが再利用されるのか予測可能なプログラムでないと、効果的に共有メモリを利用するのは困難である。Fermi
世代の GPU に導入された汎用キャッシュにより、これら問題の一部は解消される。

共有メモリ使用の有無による姫野 BMT 性能の違い
姫野 BMT を事例に、共有メモリ使用の効果を確認した結果を説明する。以下は、共有メモリを使用せずに CUDA で記述した姫野 BMT と、共有メモリを使用した場合の性能を、Tesla 世代、Fermi
世代、双方の GPU で測定した結果である。

- 30 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

Tesla 世代の GPU である GTX-285 では、共有メモリ未使用時は 30.8 GFLOPS、使用時は 76.9
GFLOPS と、共有メモリ使用の有無で大きな性能差があった。一方、Fermi 世代 GPU である GTX-480
では、共有メモリ未使用時は 72.3 GFLOPS、使用時は 74.0 GFLOPS と、共有メモリ使用の有無は性能に大きく影響しなかった。共有メモリ使用時の方が、若干ではあるが性能は高いので、ピーク性能を追求する場合には、Fermi 世代の GPU でも、共有メモリ使用は依然として有効なチューニング
手法であるが、ソースコードの可読性等を考慮すると、共有メモリ使用はあまり推奨できない。実際、姫野 BMT の場合、コアルーチン行数は、共有メモリ未使用時が 70 行程度、共有メモリ使用時
が 230 行程度であった。同程度の性能を 1/3～1/4 のコード記述量で実現できていることになり、開発効率の観点でも共有メモリ使用は推奨できない。Tesla 世代の GPU では、共有メモリをキャッシュメモリの代わりとして使用することが半ば常識化していたが、Fermi 世代以降では、共有メモ
リはその本来の用途、例えば、近接スレッド間のデータ交換や reduction 処理などに使うのが望ましい。
キャッシュチューニング
Fermi 世代 GPU のキャッシュメモリを有効活用するには、キャッシュチューニングとしてソースコードへの__synchthreads()挿入が必要な場合がある。姫野 BMT は、このキャッシュチューニン
グが有効な事例であり、__synchthreads()を挿入しないと性能は 43.0 GFLOPS に留まったが、__synchthreads()の挿入で 72.3 GFLOPS まで性能が向上した。__syncthreads()は近接スレッド間の同期をとる命令であり、実行オーバーヘッドがあるため一般的には使用しないほうが性能上は望ましいが、姫野 BMT のような全スレッドが規則的に動作するプログラムでは、__syncthreads()により全スレッドの足並みが揃うことで、キャッシュヒット率が向上するという効果がある。本キャッシュチューニングが機能する仕組みに関しては、添付資料 9 を参照願いたい。
メモリアクセス時間- 31 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

基本的なメモリ性能の一つである、メモリアクセス時間を、Tesla 世代 GPU（GTX-285）と Fermi
世代 GPU（GTX-480）で測定した結果を以下に示す。

[単位:ns]

L1 / L2

共有キャッシュメモリ

GTX-285

---

---

351

GTX-480

71

251

402

1.2 / 313

65

Nehalem(2.93GHz)

GTX-285 のメモリアクセス時間は、キャッシュメモリがないため、351ns で一定である。一方、
汎用キャッシュを搭載している GTX-480 は、L1 キャッシュ、共有キャッシュ、グロバール・メモ
リでアクセス時間が異なり、それぞれ、71ns、251ns、402ns であった。L1 キャッシュのアクセス
時間 71ns は、汎用 CPU のキャッシュメモリ・アクセス時間と比べて非常に遅く、汎用 CPU のメモ
リアクセス相当の時間がかかっている。実行スレッドを頻繁に切替えてメモリアクセス時間を隠蔽する GPU なら、これで十分なのかもしれないが、GPU と汎用 CPU では設計コンセプトが違うことがよく分かる結果である。
なお、グローバル・モリへのアクセス時間は、GTX-285 から GTX-480 で 50ns 以上遅くなっているが、原因は不明である。メモリ構造が階層化されたこと（汎用キャッシュが追加されたこと）
が影響しているのかもしれない。

- 32 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

メモリバンド幅
基本的なメモリ性能の一つであるメモリバンド幅を測定した結果を、Tesla 世代 GPU（GTX-285）と Fermi 世代 GPU（GTX-480）で測定した結果を以下に示す。メモリバンド幅はメモリコピー時の実測値である。

GTX-285

GTX-480

Nehalem

GDDR3

GDDR5

DDR3

2.484 GHz3.696 GHz

1.333 GHz

512 bit384 bit

192 bit

ピークバンド幅

159.0 GB/s

177.4 GB/s

32.0 GB/s

CUDA version

3.0

3.0

---

メモリタイプ
データ転送レート
ビット幅

GTX-285 では、標準 API の cudaMemcpy を使用した場合と、自作メモリコピールーチンの
user_memcpy を使用した場合で、ほぼ同じ実効メモリバンド幅となったが、GTX-480 では、標準
API と自作メモリコピールーチンで実効メモリバンド幅が異なり、標準 API のときは性能が低い結果が得られた。これは、cudaMemcpy が、まだ Fermi 世代 GPU 向けに最適化されていなかったためと思われる。
また、ピークメモリバンド幅は GTX-480 の方が上であるにも関わらず、GTX-285 の方が実効メ- 33 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

モリバンド幅は良いという結果となった。メモリバンド幅は多くの HPC アプリケーションにとって重要なポイントであり、Fermi 世代の GPU でメモリバンド幅が強化されなかったのは、HPC 観点
では残念な点の一つである。

- 34 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書

5. おわりに
アクセラレータは、エクサスケール・スパコンに求められる、現在よりも桁違いに良い電力あたり演算性能を実現する一つの手法として有望視されているが、エクサスケール世代におけるアクセラレータのあるべき姿は、まだ明確にはなっていない。本アクセラレータ技術 WG では、アクセ
ラレータの長所・短所、そして課題を浮き彫りにするため、近年普及しつつあるアクセラレータの一形態とも言える GPU を題材に、具体的な事例を用いて調査・議論を重ねた。
第２章では、ディレクティブ方式によるアクセラレータ利用方式に関する調査をまとめた。ディレクティブ方式は GPU 利用の壁を引き下げる効果が大きく、一般的なアプリ開発者が容易に GPUを使うことを可能とし、ループ構成が規則的であり GPU に向いたアプリケーションであれば、汎用 CPU を上回る性能を実際に達成できることを確認した。しかし、性能的にはアクセラレータ向け
拡張言語で開発したものには及ばす、また、ループ構造が複雑なアプリケーションでは、そもそもどのようにディレクティブを挿入してよいのか分からないという課題が浮き彫りになった。ディレクティブの種類を豊富にして様々な並列化手法を記述可能にすることも大切であるが、アプリ開発
者の並列化作業を支援する仕組みが今後は重要になるだろう。
第３章では、アクセラレータ向け拡張言語によるアクセラレータ利用方式に関する調査をまとめた。CUDA や OpenCL 等の拡張言語を用いることで、ディレクティブ方式では達成できない性能を実現可能なこと、本来は GPU に不向きなアプリケーションであっても、実装アルゴリズムを GPU向きに変更することで汎用 CPU を上回る性能を達成できることを事例ベースで確認した。しかし、拡張言語で高い性能を達成するには、GPU の動作原理・特性の理解が不可欠であり、一般的なアプ
リ開発者にこれを求めるのは酷である。
現時点では、アクセラレータ環境において、プログラミングの容易性と高性能を両立する解は残念ながら存在しない。これを解決するアプローチとしては、アクセラレータ HW 強化、ディレク
ティブ方式の記述性向上、並列化支援機能、自動性能チューニング、DSL(Domain specific language)の開発、アクセラレータ向け拡張言語教育など、いろいろな方向性が考えられる。現在はアクセラ
レータを使いこなす障壁は高いが、エクサスケール世代に向けて、今後はこの障壁を下げることがますます重要になる。今後も継続的に、アクセラレータの技術動向・方向性に関する議論を行う必要があるだろう。
最後に、SS 研事務局の WG 運営に感謝の意を表すると供に、各種調査や議論に参加して頂いた
WG メンバの皆様に深く感謝致します。また、本報告書が、アクセラレータ技術に関する議論の一助となれば幸いである。

- 35 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 1]

- 36 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 37 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 38 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 39 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 2]- 40 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 41 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 42 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 3]

- 43 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 44 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 4]- 45 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 46 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 47 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 5]- 48 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 49 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 6]

- 50 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 51 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 7]

- 52 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 8]- 53 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 54 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 55 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 56 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 57 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 58 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 9]- 59 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 60 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 61 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 10]- 62 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 63 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 64 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 11]

- 65 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 66 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 67 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 12]

- 68 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 69 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 70 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 71 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 72 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 73 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 13]- 74 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 75 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 14]

- 76 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 77 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 78 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 15]

- 79 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 80 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 81 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 82 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書[添付資料 16]

- 83 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 84 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 85 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 86 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 87 -

サイエンティフィック・システム研究会 アクセラレータ技術 WG 成果報告書- 88 -

【発行者】サイエンティフィック・システム研究会【編

集】サイエンティフィック・システム研究会
アクセラレータ技術ＷＧ【発行日】2012 年 3 月 26 日
◆

◆ ◆

＜連絡先＞
サイエンティフィック・システム研究会 事務局〒105-7123 東京都港区東新橋 1-5-2 汐留シティセンター
TEL : 03-6252-2582
E-mail : office@ssken.gr.jp※著作権は各原稿の著者または所属機関に帰属します。無断転載を禁じます。