1

九州大学情報基盤研究開発センター
2019.06.04 13:30-17:30
2019.06.13 13:30-17:30

2• 第一部（本資料を用いる）– ITOの説明– 実際にログインしてみる、PGIコンパイラをインストールする– GPUと並列計算についての説明• 第二部（別資料を用いる）– OpenACCの説明と実習– ～17:30：質疑・自由演習• 配付したITOのアカウントは1週間有効です– 受講後にも色々と試してみてください3• ITOにログインできることを確認する（別紙）• PGIコンパイラ Community Edition の導入（別紙）4• 以下の内容を適当なファイルに保存する（仮にtest.shとする）#!/bin/bash#PJM -L "rscunit=ito-b"#PJM -L "rscgrp=ito-b-lecture"#PJM -L "vnode=1"#PJM -L "vnode-core=9"• コピペする場合は記号が全角に化#PJM -L "elapse=00:05:00"
けるなどしないか気を付けること#PJM -o out.txt• 書かれている内容については後で#PJM -j
説明する#PJM -S

module load ~/opt/pgi/modulefiles/pgi/19.4
pgaccelinfo• pjsub test.shを実行する• しばし待つ（バックエンドでジョブが実行される）• out.txtファイルが作られたら中身を確認する– （作られてすぐは中身が書き込まれていないことあり、もう少し待つ）

5• この建物の地下1Fに設置されているスーパーコンピュータ

バックエンド サブシステムA

バックエンド サブシステムB

フロントエンド サブシステム• スーパーコンピュータ＝「パソコンのようなもの」が大量に並べてつなげてある、ようなイメージ– どの程度「パソコンのようなもの」であるかはシステム次第– ITOは「普通のパソコンに近い」部分が多い• 普通のパソコンで使っているものより「ちょっとスゴイ」ものを多数並べて全体として「とてもスゴイ」性能を得ている

共有ストレージ

ノード間高速ネットワーク

6より詳細な情報はWebを参照
https://www.cc.kyushu-u.ac.jp/scp/system/ITO/01_intro.html• 主に計算に使う部分：バックエンドとフロントエンド– バックエンド サブシステムA• バックエンド1台の性能はPC• 1CPUあたり18コア×3GHzと比べて「ちょっとスゴイ」• （2CPU + 192GBメモリ）×2000• 多数並べると全体として「とてもスゴイ」• 合計 6.91 PFLOPS– バックエンド サブシステムB• 1CPUあたり18コア×2.3GHz
今回の利用対象• （2CPU + 384GBメモリ + 4GPU）×128• 合計 3.05 PFLOPS– フロントエンド• 1CPUあたり18コア×2.3GHz、（2CPU+384GBメモリ+1GPU）×160• 1CPUあたり22コア×2.2GHz、（16CPU+12TBメモリ+1GPU）×4➢ 1PFLOPS＝ 1秒間に1000兆回の実数計算ができる• 2018年のハイエンドAndroid端末は500GFLOPS程度• 1000=1K, 1000K=1M, 1000M=1G, 1000G=T, 1000T=1P

7• 画像処理用のハードウェア– CPUやマザーボードに組み込まれたチップ、または拡張スロットに搭載するビデオカードとして普及– 本来の役割：高速・高解像度描画、3D描画処理（透視変換、陰影・照明）、画面出力
3次元画像描画の手順① (2, 2)② (8, 3)③ (5, 7)
オブジェクト単位、頂点単位、ピクセル単位で同時（並列）処理が可能– 同時にできることが多いため、それにあわせたハードウェアへと進化– 数値シミュレーション（科学技術計算）を高速に行えるハードウェアと両立することがわかり、GPUの重要な市場の一つとなってきた– 現在では特にビッグデータ、機械学習、AI処理などで性能を発揮

8• GPUはCPUと比べて単純な処理を（順序を問わず）たくさん行うことが求められるため、それに適したハードウェアへと進化してきた– HW構成バランスのイメージ
マルチコアCPU

メニーコアプロセッサ

GPU

計算ユニット
メモリ
キャッシュなど
制御部など– GPUの特徴：たくさんの計算ユニット、高速なメモリ、OSレス• OSを動かし様々な仕事をせねばならないCPUとは大きく異なる• 単体で利用できない、CPUによるサポートが必要• 現代の計算加速装置（アクセラレータ）の代表格• CPUよりスゴイ、というよりも、得意とする仕事が違うことが重要

9※SPARC64XIfxは名古屋大に設置されているFX100に搭載されているモデルの値

Tesla P100

Xeon Gold 6140
（サブシステムB）

Xeon Gold 6154
（サブシステムA）

SPARC64 XIfx（「京」後継機）

アーキテクチャ
名PascalSkylake-SP

SPARC64 XIfx

コア数3584 FP32/1792 FP64
(56 SMs)

18（HTにより36スレッド実行可能だが、HTは無効化している）

32（ + 2アシスタン
トコア）

動作周波数

1.328 GHz – 1.480 GHz

2.3 GHz - 3.7 GHz

3.0 GHz – 3.7 GHz

2.2 GHz

搭載メモリ

HBM2
16 GBDDR4192 GB/socket

DDR496 GB/socket

HMC
32GB

HPL
（ピボット選択付きLU分解）

4 TFLOPS程度※NVIDIA提供バイナリ
実測値

0.9 TFLOPS程度※MKL実測値

1.1 TFLOPS程度※MKL実測値

0.97 TFLOPS※名古屋大提供値

STREAM Triad
A[i]=B[i]+C[i]550 GB/s

95 GB/s

95 GB/s

210 GB/s↑非常に高い性能を持つ、使わなければもったいない！

10• 56のStreaming Multiprocessor (SM)• SMあたり32のFP64コアと64のFP32コア、キャッシュなど• コアのスケジューリングはWARP単位（SIMD長32のようなもの）
SM

11• TOP500 1~10位中、5システムがNVIDIAのGPUを活用– 上位50中の13、上位100中の24、全500中の128
Total Cores

ACC Cores

Rmax[TFlop/s]

Rpeak[TFlop/s]NameSiteCountry

1 Summit

ORNL

US2397824

2196480

143500200794.88Volta GV100

2 Sierra

LLNL

US

1572480

1382400

94640

125712 Volta GV100

3 Sunway TaihuLight Wuxi

China

10649600

4 Tianhe-2A

Guangzhou

China4981760455475261444.5

100678.66Matrix-2000

5 Piz Daint

CSCS

Switzerland

38787231942421230

27154.3Tesla P100

6 Trinity

LANL/SNL

US

9790727 ABCI
8 SuperMUC-NG

AIST
LeibnizJapan
Germany

391680305856

9 Titan

ORNL

US

56064010 Sequoia

LLNL

US

1572864

93014.594

ACC

125435.9None20158.7

41461.15None

34816019880
19476.6

Tesla V100
32576.635 SXM226873.856None

26163217590

27112.55Tesla K20x

17173.224

20132.659None※2018年11月のランキング情報

12• 何故GPUを活用する必要があるのか？
– ITOにGPUが搭載されているから……というのも事実ではあるが– GPUは非常に高い性能をもつハードウェアであり、うまく活用できれば大変強力な研究の道具となる：成果を得るための加速装置– GPUは国内外の様々な計算環境に導入されているため、利用スキルを得ておくことはきっとプラスになる• どうやって使う？
1. 並列化されたソフトウェア（アプリケーション）を使う
2. （並列化されていないプログラムから）並列化されたライブラリやフレームワークを使う
3. 自分で並列計算を実装する• 自分が扱いたい任意の問題（アプリ）をGPU化できる• 扱いたい問題が既にGPU化されているならそれを使っても良いが、GPUのことを理解しているかどうかでさらに高い性能が得られるかどうかが決まる、こともある– 例：特定の形状・大きさの行列しか扱わない

13• GPU（などのアクセラレータ）向けのプログラムを簡単に書くことのできる並列化プログラミング言語– NVIDIAのGPUに限定されない• C/C++やFortranで書かれたプログラムに簡単な「指示文」を追加するだけでGPU対応が可能– コンパイラ向けのコメントを記述する– CPU向けの並列化に広く用いられているOpenMPを踏襲、理解も利用も容易– 最適化を行う上ではある程度GPUの知識があった方が良い• 「どんなプログラムでもGPU化できる」「どんなプログラムでも高速化できる」わけではないが、多くのプログラムに対して有用– 自作コードのGPU化を行うことができる上で最も簡単でコストパフォー
マンスの高い方法• コストパフォーマンスが高い＝労力に対して十分な性能が得られる

14• 規格などの情報– https://www.openacc.org/• 主な対応コンパイラ– 商用：PGI（無償版あり）、Cray、国家超级计算无锡中心– 研究：Omni、OpenARC、OpenUH、ROSEACC– OSS：GCC• プロファイラやデバッガなど– allinea MAP/DDT、PGI pgprof、NVIDIA nvprof/nvvp• 日本語で書かれた資料– ソフテック社（PGIコンパイラの代理店）の資料• OpenACC ディレクティブによるプログラミング by PGI Compilers
https://www.softek.co.jp/SPG/Pgi/OpenACC/

15• 並列計算とは何か？並列プログラミングとは何か？
– 並列計算：何らかの計算処理を同時並行的に行うこと• 並列：同じ処理を同時に行う、ある処理を幾つかのサブ処理に分けて同時並行的に行う• 平行：何らかの処理を同時並行的に行う– 並列プログラミング（並列化プログラミング）：並列化・並列計算を行
うためのプログラミング• なぜ並列計算を行うのか？
– 短時間で終わらせたい計算があるから、計算機（CPUやパソコンなど計算を行うHW）の持つ能力をフル活用したいから– 現代の計算機は並列計算により高い性能を達成するシステム• 逐次計算性能を上げることができなくなっている– PC用CPUもスマートフォン向けCPUもマルチコアCPUが主流• 大規模なスーパーコンピュータの総コア数は100万以上➢ HWの性能を十分に引き出すには並列計算が必須

16• 元となる逐次計算– 単純な繰り返しループ計算
for(i=0; i<N; i++){A[i] = B[i] + C[i];D[i] = E[i] + F[i];}• ループ内の処理を分割し、同時に計算
PE1

for(i=0; i<N; i++){A[i] = B[i] + C[i];}PE2for(i=0; i<N; i++){D[i] = E[i] + F[i];}PE: Processor Element、CPUやGPUなど計算をする「何か」• ループそのものを分割し、同時に計算
PE1

for(i=0; i<N/2; i++){A[i] = B[i] + C[i];D[i] = E[i] + F[i];}PE2for(i=N/2; i<N; i++){A[i] = B[i] + C[i];D[i] = E[i] + F[i];}• 本講習会で扱うOpenACCによる並列化はこちらのイメージ– OpenMPもこちら

17• 単純なループによる並列化とは異なり、やや大きい粒度の並列計算などでよく活用• OpenMPにおいて近年サポートが活発• GPU（OpenACC）にはあまり適していないため今回は扱わない

18• 並列化できないプログラムも少なくない– 計算順序に制約がある場合は困難• 工夫により可能となることもある

for(i=1; i<N; i++){A[i] = A[i-1] + B[i];C[i] = A[i] + D[i];}• （プログラム高速化全般に言えることだが）速くした部分しか速くならない
逐次処理A

一部並列 処理A

処理B
B

例：ループイタレーション間にも、同じループイタレーション内にも、依存関係がある（並列化が困難）

time
頑張って処理Bを並列高速化したが、全体の実行時間的には……

– 探索問題などでは分割数（並列度）よりずっと速くなることもある★解• 探索対象★
• 逐次探索★
• 並列探索– プログラムのどこをGPUに担当させるかをちゃんと考える必要がある

19• 多数のコア全てを満たすに十分な仕事を与える– 数千のコア全てをフル稼働させる並列度が必要– コア数の数倍以上の仕事があることが望ましい• メモリアクセスを待つ間に別の仕事ができる• WARPを意識する– 内部的には32演算器単位で動作しており、分岐の際などに注意が必要• Voltaから変わってしまった• 連続メモリアクセスについて考える– コアレスなメモリアクセスが高速• 詳細は終盤（OpenACCプログラムの最適化）にて解説する– 基本的には、並列度が高くて分岐の少ない単純なプログラムが良い20• 並列計算環境が普及した今日では様々な「並列化のための道具」が使われている– バラバラだと提供側・利用側ともに不便なため共通化されることが多い• ある環境向けに書いたものが別の環境でも利用できる（互換性）• 性能まで互換性があるとは限らない点には注意が必要（性能可搬性）– ある環境では有効であった最適化が別の環境では性能低下要因になることも– 環境が変わってもその環境ごとに常に最大の性能が得られるようにするための研究も行われている→自動チューニング

• 自動化はできないのか？
– 全く不可能なわけではない、ある程度はコンパイラ等が行ってくれる– 「コンパイラ様のご機嫌を伺う」コードを書く必要がある、コンパイラによって差が大きい、簡単なコードでないとうまくいかないことが多い• 単純な行列積などコードの形状と最適化のパターンが決まっているものは人が書くよりコンパイラやライブラリの方が得意– OpenMPやOpenACCは簡単・少量の記述で効果的な並列高速化を可能とする（良いところ取り、規格化により汎用性・可搬性も担保）

21• MPI：プロセス間の通信規格、特に複数ノード利用時に必須– MPI_Send, MPI_Recv, MPI_Gatherなどの通信関数を使う• pthread：スレッド並列処理のための関数群– pthread_createなどのスレッド操作関数を使う– 現在ではアプリケーションコードで直接利用することはあまりない• OpenMP：スレッド並列処理、主にループ並列化向けの指示文規格– #pragma omp parallel for 、!$omp parallel do– 最近の規格ではタスク処理やGPU対応などを拡充• GPUに対応した実装はまだ少ないが、将来的にはOpenACCを吸収？

• CUDA：NVIDIA社のGPU向け、GPUを普及させた最大要因の一つ、
NVIDIA GPUの能力をフル活用したければ必須• OpenCL：「汎用版CUDA」のようなもの、FPGAなどでも利用可能• コンパイラによる自動並列化：SIMD並列化など一部の処理で有用

22• 単純な配列コピーの例
GPU上で行われる処理__global__ void gpukernel(int N, float* C, float* A){（GPUカーネル）
int tid = blockIdx.x*blockDim.x + threadIdx.x;
int nt = gridDimd.x * blockDim.x;
for(int i=tid; i<N; i+=nt){• 同時にたくさん実行される
C[i] = A[i];• 自分のIDを元に計算するべき範囲を特定する}}
int main(int argc, char **argv){
int N = 100000;
float *A, *C;
CPU上で行われる処理
float *d_A, *d_C;
A = (float*)malloc(sizeof(float)*N);C = (float*)malloc(sizeof(float)*N);• 専用の関数などを用いて様々な処理を行う
cudaMalloc((void**)&d_A, sizeof(float)*N);
cudaMalloc((void**)&d_C, sizeof(float)*N);
cudaMemcpy(d_A, A, sizeof(float)*N, cudaMemcpyHostToDevice);
cudaMemcpy(d_C, C, sizeof(float)*N, cudaMemcpyHostToDevice);gpukernel<<<4,4>>>(N, d_C, d_A);
cudaMemcpy(C, d_C, sizeof(float)*N, cudaMemcpyDeviceToHost);
cudaFree(d_A);
return 0;
慣れてしまえばそれほど難しいものではないが、記述量がそ}れなりにあり「カジュアルなGPU利用」にはあまり向かない23• GPUの性能を活用できる様々なライブラリが公開されている– https://developer.nvidia.com/gpu-accelerated-libraries– 例• Deep Learningライブラリ– cuDNN, TensorRT, DeepStream SDK• 数値計算・数学ライブラリ– cuBLAS, cuSPARSE, cuRAND, cuFFT, ...

• 通信、C++クラス、etc.
– NCCL, Thrust, OpenCV, MAGMA, ...

– NVIDIA社自ら手がけるものは「CUDA-X」と呼ばれるようになった– 各自が行いたい処理とマッチしていれば容易にGPUの性能を活用することができる