OpenAccを用いた地震動シミュレーションの並列化 - 早稲田大学リポジトリ

2012 年度修士論文
OpenAcc を用いた地震動シミュレーションの
並列化
Parallelization of Ground Motion Simulator by OpenAcc
早稲田大学大学院基幹理工学研究科
情報理工学専攻
須田昌和
学籍番号 5111B057-6
提出 2013 年 2 月 1 日
指導笠原博徳教授
目次
概要
5
Abstract
6
第 1 章はじめに
7
第 2 章 GPU アーキテクチャ、及びプログラミング環境概要
2.1 CUDA 概要 . . . . . . . . . . . . . . . . . . . . . . .
2.2 GPU アーキテクチャ . . . . . . . . . . . . . . . . . .
2.2.1 GPU のプロセッサ構成 . . . . . . . . . . . .
2.2.2 デバイスメモリとオンチップメモリ . . . . . .
2.3 CUDA プログラミングモデル . . . . . . . . . . . . .
2.3.1 GPU の利用法 . . . . . . . . . . . . . . . . .
2.3.2 CUDA における並列処理 . . . . . . . . . . . .
2.3.3 メモリアクセスと CUDA API . . . . . . . . .
2.3.4 CUDA プログラミングの例 . . . . . . . . . .
2.4 OpenAcc の概要 . . . . . . . . . . . . . . . . . . . .
2.4.1 PGI コンパイラの概要 . . . . . . . . . . . . .
2.4.2 OpenAcc の機能 . . . . . . . . . . . . . . . .
2.4.3 OpenAcc の主な指示文 . . . . . . . . . . . . .
2.4.4 OpenAcc の主なプログラミング例 . . . . . .
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
8
8
9
9
9
11
11
11
11
12
14
14
14
14
15
第 3 章評価アプリケーション
17
第 4 章性能評価
4.1 評価環境 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
4.2 評価手法 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
20
20
20
第 5 章 Discussion
22
第 6 章まとめと今後の課題
23
まとめと今後の課題
6.1 まとめ . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
6.2 今後の課題 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
23
23
23
謝辞
24
1
参考文献
24
2
図目次
2.1
2.2
2.3
2.4
2.5
2.6
GPU のハードウェア構成 . . . . . . . . . . . . . . . . .
CUDA 並列処理の概念と,GPU アーキテクチャとの対応 .
Doall ループの例 . . . . . . . . . . . . . . . . . . . . . .
図 2.3 の CUDA 書き換え後 . . . . . . . . . . . . . . . .
Doall ループの例 . . . . . . . . . . . . . . . . . . . . . .
OpenAcc の例 . . . . . . . . . . . . . . . . . . . . . . . .
.
.
.
.
.
.
9
11
12
13
16
16
3.1
計算に用いる不連続格子 . . . . . . . . . . . . . . . . . . . . . . . . .
18
4.1
4.2
評価結果 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
評価結果 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
20
21
3
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
.
表目次
2.1
GPU 上で使用可能なメモリの種類 . . . . . . . . . . . . . . . . . . .
10
4.1
Tesla C2050 の仕様 . . . . . . . . . . . . . . . . . . . . . . . . . . . .
20
4
概要
これまで、阪神大震災を契機に地震に関する調査研究が行われてきた。その調査の
中で地震のリスク評価の基礎となり得る精度で地震ハザードを予測できるよう高精度
かつ汎用性のある地震波伝播（強震動）シミュレーション手法の開発に関する研究が
行われている。地震波伝播シミュレーションにおいて詳細な 3 次元地下構造を十分な
精度で離散化し、短周期の地震波まで計算するためには細かな格子が必要であるため、
実用的な計算においては格子数が数億から数十億に及ぶ規模のモデルを扱うことにな
り、計算には大きな計算機リソースが必要であり、数日以上の計算となることもある。
また、近年, グラフィック専用のハードウェアである GPU を汎用科学計算に用いる
GPGPU が注目を集めており、これは GPU 内の多数のコアを用いることで大規模
のデータに対して並列化し、高速化を実現している。しかし、GPU プログラミング
には独自の専用言語の知識が必要で、プログラミングが困難であるという問題があっ
た。そこで、本稿では地震波伝搬シミュレーションを GPU を用い並列化することで
シミュレーション時間の高速化を目指した。
5
Abstract
Since Great Hanshin Earthquake,investigation of earthquake has been started. In
that investigation, to estimate earthquake in accuracy to base of evaluating risk,
the reserch about Ground Motion Simulator with high accuracy and flexibility has
been studied. We need fine lattice in order to disperse three dimensions structure
in enough accuracy and caluculate short cycle earthquake. In practical caluculation
we must consider from hundreds of millions to billions lattices. We need huge caluculating memory and it takes a few days to caluculate. Recently, GPUs, which are
the hardware of graphic processing, have been applied to general problems. They
are called GPGPU and have attracted much attention, using many cores of GPU
we can parallelize huge datas and get speed up. But we must study special language for GPU programming, GPU programming is very diﬃcult. This paper says
improvment in the speed with GPU parallelization for Groud Motion Simulator.
6
第1章
はじめに
近年, 従来グラフィック専用のハードウェアであった GPU(Graphics Processing Unit)
を, 高い演算能力を活かすことで, 汎用計算に利用する GPGPU(General Purpose
Computing on GPUs) という考え方が注目を集めている.[1]GPGPU は, 大量の計算
量や, 大量のデータを扱う分野で有効であり, 流体計算,N 体問題, 金融工学などの分
野での有効性が確認されている. 近年のスパコンの TOP500 を見ると, 中でも膨大な
数の GPU を搭載した東工大の TSUBAME や, 中国の天河は高い演算性能を見せて
いる。このように GPGPU は極めて演算性能が高いハードウェアであり, 同時に一般
に広まりつつあるアクセラレータである. 加えて、NVIDIA や PGI をはじめとする
複数の企業が合同で発表した新しいプログラミング規格である OpenACC という規格
ができた。[2] これは従来のプログラムにコンパイラ向けの指示文を挟むことによって
自動的に GPU 用に並列化されたコードを生成する手法である。これにより、GPU プ
ログラミングにおける独自言語に対する知識がなくても GPU を用いたプログラミン
グが可能になった。GPU 内の多数のプロセッサを用いることで計算時間の高速化が
可能になる一方で、GPU の場合,CPU に PCI-Express を介して外付けをする形をと
るため, データ転送速度が遅く, 性能向上を行う上でのボトルネックとなる可能性があ
る。本稿では地震動シミュレーションを GPU を用いて並列化し、かつデータ転送を
極力抑制し評価を行った。本論文の構成は以下のようになる．第 2 章で GPU アーキ
テクチャについて、第 3 章では評価に用いた地震動シミュレーションについて、第 4
章は本手法の性能評価について考察する．第 6 章で本稿のまとめを述べる.
7
第2章
GPU アーキテクチャ、及びプログ
ラミング環境概要
本章では,GPU アーキテクチャと CUDA の概要、及び OpenAcc の概要について述
べる。
2.1
CUDA 概要
CUDA とは,NVIDIA 社が提唱する,GPU 上で汎用計算を行わせるためのハードウェ
アアーキテクチャ, およびソフトウェア環境のことである.[3] シェーダ言語と呼ばれる
GPU 専用の言語を用いず,C 言語に似た言語仕様になっている. また, 数値計算でよく
使われる初等関数や FFT(高速フーリエ変換), 行列計算の基本的なライブラリも実装
されている.
8
2.2
GPU アーキテクチャ
ここでは、GPU アーキテクチャの説明を行う.
2.2.1
GPU のプロセッサ構成
GPU のハードウェア構成を図 2.1 に示す.
図 2.1: GPU のハードウェア構成
図 2.1 のように,GPU は「マルチプロセッサ」(一般的には, プロセッサアレイと呼
ばれるので, 以降プロセッサアレイと呼ぶ) が, 複数個並んだ構成をとっている. 一つの
プロセッサアレイ内には,「ストリームプロセッサ」が複数搭載されている. プロセッ
サアレイとストリームプロセッサの数は GPU のバージョンによって異なる.
2.2.2
デバイスメモリとオンチップメモリ
GPU のメモリは２種類存在する. １つは,GPU ボード上に実装された「デバイスメ
モリ (グローバルメモリ, コンスタントメモリ, テクスチャメモリローカルメモリ)」, も
う一つは GPU 内部のプロセッサアレイ毎に実装されている「オンチップメモリ (シェ
アードメモリ, レジスタ, テクスチャキャッシュコンスタントキャッシュ) である. この
うち, テクスチャメモリとコンスタントメモリについては, それぞれテクスチャキャッ
シュとコンスタントキャッシュにデータがキャッシュされるが, ローカル・メモリとグ
ローバル・メモリはキャッシュングされない.CPU が読み書きできるのは, グローバル
メモリだけなので,GPU の計算結果はグローバル・メモリに格納し,GPU の計算が終
了した後で CPU が計算結果をグローバルメモリから回収する. メモリに関する API
9
表 2.1: GPU 上で使用可能なメモリの種類
メモリの種類
デバイスメモリ
オンチップメモリ
グローバルメモリ
コンスタントメモリ
テクスチャメモリ
ローカルメモリ
シェアードメモリ
レジスタ
テクスチャキャッシュ
コンスタントキャッシュ
GPU 側からのアクセス
R/W
R
R
R/W
R/W
R/W
R
R
ホスト側からのアクセス
R/W
W
W
N/A
N/A
N/A
N/A
N/A
は後述する. 各メモリの一覧を表 2.1 に示している. 表中の N/A は, アクセス不可を意
味する.
このためグローバル・メモリは頻繁に使用する必要があるが,GPU との間のレイテ
ンシは 400∼600 サイクルと大きいことと, キャッシングされないことを考えると, 性
能が低下する要因となりうる. また, グローバル・メモリと CPU の主メモリとのデー
タ転送速度は 2GB/sec と (理論値は 4GB) 小さいことも, 性能低下の要因となりうる.
10
2.3
CUDA プログラミングモデル
本節では,CUDA の API や, 具体的なプログラミングモデルの説明を行う.
2.3.1
GPU の利用法
CUDA を用いて GPU を制御するとき,GPU を装備した CPU を管理コア (以降ホス
ト) として利用する. ホストで main 関数を実行し, そこから GPU で実行する関数 (以
降カーネル) を呼び出す形をとる.
2.3.2
CUDA における並列処理
CUDA では,「スレッド」と呼ばれる並列処理概念が用いられる. カーネルを GPU
で実行する際には, ハードウェアが処理をスレッドと呼ばれる単位に分割し, 各スレッ
ドを GPU 内のプロセッサに割り当てる. カーネルは「グリッド」とも呼ばれる.
GPU がグリッドを実行する際に, 各プロセッサアレイに並列処理をさせるために, グ
リッドを複数の「ブロック」と呼ばれる単位に分割する. ブロックは一つのプロセッ
サアレイに割り当てられる. ただし, 一つのプロセッサアレイにひとつのブロックとい
うことではなく, プロセッサアレイ内のレジスタやシェアード・メモリに余裕がある
限り, 複数のブロックがプロセッサアレイに割り当てられる.
ブロックはさらに細粒度の「スレッド」に分割される. スレッドが一つのストリーム
プロセッサに割り当てられ, 並列に計算を行う. また, ブロック内のスレッド全体で同
期をとることはできるが, グリッド内でのスレッド同期はとれない. 以上のイメージを
図 2.2 に示す.
図 2.2: CUDA 並列処理の概念と,GPU アーキテクチャとの対応
2.3.3
メモリアクセスと CUDA API
GPU のメモリ空間の種類については,2.2.2 で述べたとおりである. 先述の通り, ホ
ストがアクセスできる GPU のメモリ空間は, グローバルメモリのみである.
11
for (i = 0; i < N; i++) {
for (j = 0; j < N; j++) {
C[i][j] = A[i][j] + B[i][j];
}
}
図 2.3: Doall ループの例
ホストと GPU は異なるメモリ空間をもつので,GPU で計算を行うために, また計算
結果をホストが保持するためには両者間のデータ転送をおこなう必要がある. ホスト
から GPU へのデータ転送が行われる前に, グローバルメモリ上のデータ領域の確保を
行うことになる. グローバルメモリ空間におけるデータ領域は,cudaMalloc 関数をホ
スト側で実行することで行われる. 確保されたデータ領域はカーネルの開始/終了にか
かわらず,cudaFree 関数によって解放されるまで保持される.
データ転送は, ホスト側で cudaMemcpy 関数を実行することで行われる. 転送方向
は, 関数の引数で明示する.
2.3.4
CUDA プログラミングの例
ここでは,C のコードを CUDA に実行させるときの書き換え手法の具体的な例を示す.
図 2.3 のような 2 つの float 型の配列の要素同士の合計を求める C のソースは, 図
2.4 のように書き換える.
この書き換えでは,1 つのスレッドには, 図 2.3 での内側の 1 イタレーション分の処理
が割り当てられる.
12
%%カーネルの記述
__global__ vectorAdd(float (*d_A)[N], float (*d_B)[N], float (*d_C)[N])
{
int i = blockIdx.x;
int j = threadIdx.x;
d_C[i][j] = d_A[i][j] + d_B[i][j];
}
%%グローバルメモリ上でのメモリ領域確保
cudaMalloc((void**) &d_A, sizeof(float) * N);
cudaMalloc((void**) &d_B, sizeof(float) * N);
cudaMalloc((void**) &d_C, sizeof(float) * N);
%%ホストからグローバルメモリへのロード
cudaMemcpy(d_A, A, sizeof(float) * N, cudaMemcpyDeviceToHost);
cudaMemcpy(d_B, B, sizeof(float) * N, cudaMemcpyDeviceToHost);
%%カーネルの実行
vectorAdd <<< N, N >>>(d_A, d_B, d_C);
%%グローバルメモリからホストへ, 計算結果のストア
cudaMemcpy(C, d_C, sizeof(float) * N, cudaMemcpyHostToDevice);
%%グローバルメモリ上に確保した領域の解放
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
図 2.4: 図 2.3 の CUDA 書き換え後
13
2.4
OpenAcc の概要
本節では OpenAcc の説明を行う。OpenAcc とは NVIDIA 社や PGI 社を始めとす
る複数の企業が合同で発表した新しい GPU プログラミング規格である。これにより、
GPU プログラミングにおける独自言語によらず、指示文を挿入するだけで GPU プロ
グラミングが可能になる。
2.4.1
PGI コンパイラの概要
本節では PGI コンパイラの説明を行う。PGI コンパイラは OpenAcc に対応したコ
ンパイラである。PGI コンパイラは The Portland Group Inc 社が研究開発している
自動並列化コンパイラであり、 Intel のプロセッサと AMD のプロセッサにも最適化
対応しており, 32 ビット/64 ビットの双方のマルチプロセッサにハイパフォーマンス
な並列化を行う. また, 2009 年に業界で初めての GPGPU に対応したアクセラレータ
コンパイラを提供した. スパコン TOP10 に入った東工大の TSUBAME2.0 にも導入さ
れている。
2.4.2
OpenAcc の機能
本節では OpenAcc が提供する機能について述べる。GPU で処理させたい部分につ
いて、既存のプログラム内に指示文を追記すると、コンパイラは以下の機能を組み込
んだホスト（CPU）側コードとアクセラレータ（GPU）側コードを生成する。
• GPU をイニシャライズ
• CPU ⇔ GPU 間の「データ（値）」と「プログラム」の転送の管理
• CPU ⇔ GPU 間の「処理」を管理
• その処理を GPU のアーキテクチャ（ハードウェア構造）に合わせて計算できる
よう並列分割
• その他キャッシュ等の性能最適化を図る
2.4.3
OpenAcc の主な指示文
本節では OpenAcc の主な指示文について述べる。
OpenAcc の指示文は Data 構文（データ移動指示),Accelerate Compute 構文（oﬄoad
領域指示,）Loop 構文（並列化のためのマッピング）と３つからなっており、これら
を組み合わせることで詳細な指示を与えることができる。以下では Fortran の場合の
構文を示しているが, C の場合は!$acc を#pragma と置き換えればよい。
• !$acc data [clause]
Data 構文は CPU と GPU の間の値の受け渡しに関する指示文である。
14
– copyin
ホストからデバイスへデータ転送
– copyout デバイスからホストへデータ転送
– create
デバイスのみで使用するローカル変数の割り当て指示
– present
すでにデバイス側に変数が存在していることを指示
– deviceptr
変数がデバイス側のポインタであることを宣言
• !$acc kernels
kernel 構文は指定領域に複数の独立ループ構造がある場合、それぞれを独立の
カーネルを生成し、順に実行する。
• !$acc loop
– collapse
下に続く n 段数の nested（入れ子）ループに指示文を適用する.
– seq
GPU 内で loop をシーケンシャルに実行する。
– private
loop の各 iteration に対して変数のコピーを生成し、プライベート変数を
生成
– reduction (operator: list)
iteration 間でプライベートコピーを使ってリダクション処理を行う。
2.4.4
OpenAcc の主なプログラミング例
本節では OpenAcc を用いた主なプログラミング例について述べる。ここでは,fortran
のコードを OpenAcc を用いて実行させるときの具体的な例を示す. 図 2.5 のような 2
つの float 型の配列の要素同士の合計を求める fortran のソースは, 図 2.6 のように書
き換える.
このように CUDA と比較したときに、OpenAcc では GPU の初期化や、元のプロ
グラムを変更する必要がなく容易にプログラミングできることがわかる。しかし、一
方で抽象度が高くスレッドとメモリの割り当てができなかったりと CUDA に比べて
高速化しづらいという面もある。
15
do i = 0, n
do j = 0 , j
C(i,j) = A(i,j) + B(i,j);
end do
end do
図 2.5: Doall ループの例
!$acc data copyout(C)
!$acc data copyin(A,B)
!$acc kernels
do i = 0, n
do j = 0 , j
C(i,j) = A(i,j) + B(i,j);
end do
end do
!$acc end kernels
図 2.6: OpenAcc の例
16
第3章
評価アプリケーション
本章では地震動シミュレータ (GMS: Ground MotionSimulator) の概要について述
べる．
GMS は防災科学技術研究所によってパッケージ化された、3 次元有限差分法 (FDM)
により地震波伝搬シミュレーションを行うためのツール群でおり、主に Fortran90 で
書かれた差分計算ソルバはソースコードも公開されている。
Aoi and Fujiwara（1999）や青井・他（2004）は、大きさの異なる格子を組み合わせ
ることにより効率的かつ高精度に計算を行うことの出来る不連続格子による差分法の
定式化を提案した。差分法による地震波伝播シミュレーションを行う場合、格子サイ
ズは計算すべき最短波長により決定されるため、モデルのごく一部のみが低速度の媒
質である場合でも計算領域全体を小さな格子に分割せざるを得ず、大規模なモデル計
算の大きな障害となっていた。軟弱（＝地震波速度が低い）な表層が存在する浅い部
分（領域 I）の格子点間隔は細かく、深い部分（領域 II）の格子点間隔は領域 I の 3
倍の粗い格子点間隔を有する格子モデルを用いている。計算に用いる不連続格子の図
を図 3.1 に表す。2 つの領域はオーバーラップしており、波動場の連続性が保たれる
よう内挿される。典型的な盆地構造モデルの計算において、均質な格子による場合と
比較し数倍から十数倍程度効率がよいことが分かっている。
17
図 3.1: 計算に用いる不連続格子
また、GMS ソルバの処理時間は main となるループで約９９％の割合を占める。
Do step=1,100
地震上部の移動速度の計算
地震下部の移動速度の計算
地震上部の応力の計算
地震下部の応力の計算
計算結果を HDF ５フォーマットで出力
End do
このように移動速度や応力の計算、計算結果を出力といったステップを１００回行
うモデルとなっている。また、移動速度、応力を計算する部分は以下のような構造に
なっている。
Do z =1,depth
Do y =1,length
Do x=1,width
計算部分
End do
End do
End do
上記のループは並列化可能なループとなっており、指示文を挿入することで GPU で
18
の並列化が可能である。
19
第4章
性能評価
本章では, 評価環境、評価手法及び評価結果について述べる。
4.1
評価環境
今回の評価では GPU は NVIDIA の Tesla C2050 を用いた。Tesla C2050 の仕様は
表 4.1 である。
4.2
評価手法
評価は CPU での実行、単純に GPU に実行させた場合、CPU,GPU 間の無駄なデー
タ転送を極力抑制した場合での比較と再内側ループと再外側ループでそれどれ並列化
した場合で比較した評価結果を図 4.1 と 4.2 に示す。
図 4.1: 評価結果
図 4.1 に示したとおり、CPU で実行した場合 250 秒、GPU で単純に処理をした場
合 326 秒、GPU でデータ転送を抑制した場合は 149 秒となった。これは GPU で計算
表 4.1: Tesla C2050 の仕様
コア数
シェーダクロック
メモリ容量
448
1180MHz
3GBytes
20
図 4.2: 評価結果
部分を処理することで高速化が期待されたがデータ転送によるボトルネックが大きく
転送を抑制しない場合は CPU 実行より遅くなった。また図 4.2 に示したとおり、再内
側ループで並列化した場合ループの内側の並列性しか生かせておらず、データ転送な
ども考慮するとかなり遅くなってしまった。逆に再外側ループで並列化した場合はき
ちんとループの並列性を生かせていた。
21
第5章
Discussion
本章では既存の関連研究について議論する。
B. CloutierB. Cloutier らはナビエストークス方程式などの求解にかかる時間を CUDAFortran,CUDAC,OpenAcc で比較した研究を発表している。[5] 結果としては CUDAFortran,CUDAC,OpenAcc の順で実行時間が早かった。OpenAcc は抽象度が高く高
速化のためのチューニングがしづらいのが原因であるようだ。CUDAFortran の方が
CUDAC より早かったのは CUDAFortran の方が使用できる指示文の種類が多く高速
化しやすい。
フランスの CAPS entreprise 社が開発した GPGPU アプリケーション開発ツールであ
る HMPP は OpenAcc 同様指示文ベースのプログラミングを提供する。[6] また、こ
ちらはマルチ GPU にも対応している。HMPP が生成した GPU コードを、CUDA ／
OpenCL を用いてチューニングすることで、最適化を行うことができる。
指示文ベースの GPU プログラミングは CUDA に比べてプログラミングがしやすく、
どの程度高速化が期待できるのかを調査するのに役に立つ。今後も指示文ベースの
GPU プログラミングの研究が広がっていくだろう。
22
第6章
6.1
まとめと今後の課題
まとめ
これまで GPU プログラミングに使われていた CUDA や OPENCL のような専用言
語を使わずに OPENACC による指示文の挿入のみで GPU 実行が可能になり CPU で
250 秒かかるプログラムが 149 秒で実行できた。OPENACC は CUDA と比べて関数
呼び出しに対応していない、マルチ GPU に対応していない、等のデメリットがある
一方で、プログラミングが CUDA に比べて容易、オプションの有無のみで CPU 実行
と GPU 実行を切り替えることができる等のメリットがあり、生産性が高い。
6.2
今後の課題
また、本研究で用いたモデルは 420 × 420 × 100 の格子に分割したが、より細かい
格子に分割するモデルでは GPU の高い計算力をより生かすことができ、さらなる処
理時間の向上が見込まれる。また、GPU での計算処理が高速であるため、相対的に
ファイル出力に要する時間の割合が高くなってしまう。そこで、GPU に計算処理を
投入した後、前のタイムステップのファイル出力を行うことで、GPU での計算処理
とファイル出力をオーバーラップさせ、ファイル出力に要する時間を隠蔽することが
出来ると考えられる。
23
謝辞
本研究を行う機会, 及び世界でもトップクラスの研究環境を与えてくださり，多大
な御助言，適切な御指導を頂きました笠原博徳教授と木村啓二准教授に深く感謝致し
ます．
本研究を進めるに当たり，基礎的なことから丁寧な御指導をいただき，アドバイス
をくださった島岡護氏に心から感謝致します。
また，研究に限らず様々なことを教えて下さったり，日頃から色々と面倒を見て下
さった林明宏氏を始めとする多くの諸先輩方，そして、何より共に笑い共に涙した同
期の皆さんに感謝致します．
最後に，私に大学で勉強する機会を与えて下さった両親に深く感謝致します．
簡単ではありますがこの場をお借りして皆様に深く御礼申し上げます．有難う御座
いました．
24
参考文献
[1] http://www.gpgpu.org/.
[2] http://www.openacc-standard.org/.
[3] http://www.nvidia.com/.
[4] http://www.pgroup.com/.
[5] B. Cloutier, B.K. Muite, P. Rigge, 2012 Symposium on Application Accelerators in High Performance Computing , Performance of FORTRAN and C GPU
Extensions for a Benchmark Suite of Fourier Pseudospectral Algorithms
[6] http://www.jcc-gimmick.com/gpugen.html
[7] 青井真 , 早川俊彦, 藤原広行地震動シミュレータ:GMS 物理探査第 57 巻第 6 号
(2004) 651-666p [8] 青井真:GPGPU の地震ハザード予測シミュレーションへの適応
性評価学際大規模情報基盤共同利用・共同研究拠点平成 23 年度共同研究最終報告
書 2012 年 5 月
25

Download Report