KOBE HPC サマースクール（初級） 2022 講義資料
2022/09/02
兵庫県⽴⼤学情報科学研究科
安⽥ 修悟

アクセラレータとは︖• CPUの処理を⼀部代替して全体の処理の効率を向上させる装置（ディバイス）。
– GPGPU (NVIDIA), Xeon Phi (Intel), APU (AMD)など
ホスト

コ
ア
０

コ
ア
１

コ
ア
2

ホスト側メモリ

アクセラレータコ
ア３

制御

Many coresMulti Threads
デバイスメモリ

PCIバス

アクセラレータの実⾏イメージ
ハイブリット構成 （CPU＋アクセラレータ）

ホスト

重たい計算部分をオフロード

GPU・Phi等

CPU
Mainメモリ

アクセラレータ

PCIバス

デバイスメモリ※データのオフロード通信がボトルネックとなる• CPUのメモリ空間とディバイスのメモリ空間が異なる• 一般的なメモリ帯域幅に比べてPCIバスでの低速な転送OpenACC• Accelerators⽤のAPIの⼀つ。
• CPUからディバイスに処理をオフロードする際に使⽤する。
• Fortran/C/C++⾔語でディレクティブを指定（OpenMPと同様）• CUDA等の専⽤プログラムに⽐べて、導⼊が容易。

OpenACCプログラミング• 基本フォーマット
#pragma␣acc␣ディレクティブ␣節１␣節2␣・・・{
構造化ブロック（途中で抜け出したり、終了しない）}
例#paragma ␣acc ␣kernels␣copy(a)for(i=0;i<imax;i++){a[i]=a[i]*a[i];・・・・}

OpenACCプログラミング• 基本フォーマット (Fortranの場合)!$acc␣ディレクティブ␣節１␣節2␣・・・
構造化ブロック!$acc end ディレクティブ!$acc ␣kernels␣copy(a)do i=0, imax
a[i]=a[i]*a[i];・・・・
end do!$acc end kernels

基本的なディレクティブ構⽂① Accelerator compute構⽂– オフロードするループ対象部分を指定する。
– kernels構⽂やparallel構⽂がある② DATA構⽂– ホストとアクセラレータのデータ転送を明⽰的に指⽰する。
– ⾼速化の鍵を握る③ Loop構⽂–①のオフロード領域のループのベクトル⻑や並列分割の詳細を明⽰的に指⽰する。

NVIDIA HPC SDK コンパイラー• Module環境の設定
module load nvhpc– module␣listでintelが設定されている場合
module␣unload␣intel や module␣purge
でIntel環境を先にunloadする。

• OpenACCのコンパイル
nvc –acc prog.c –o acc.out– “-acc”はOpenACCディレクティブを有効にするためのオプション。

NVIDIA HPC SDK コンパイラー• OpenMPのコンパイル
nvc –mp prog.c –o mp.out“-mp”はOpenMPディレクティブを有効にするためのオプション。
• 最適化オプション– “-Mvect=simd” : SSEやAVXのSIMD命令を⽣成．
– “-fast”:⼀般的な最適化オプション– “-O3”, “-O4”: より⾼度な最適化オプション（”-fast”より後に⼊れる
例: nvc -fast -O3 program.c ）NVIDIA HPC SDK コンパイラー• 環境変数（実⾏前に設定，run_acc.sh参照）– NVCOMPILER_ACC_TIME• 実⾏後に簡易プロファイル情報を標準出⼒に出⼒。
• 0（ゼロ）を指定すると機能を抑制（デフォルト）、ゼロ以外の整数値で機能する。

– NVCOMPILER_ACC_NOTIFY• ディバイスの実⾏イベントを表⽰する• 整数値１︓Kernel launch のイベントを出⼒、整数値２︓デー
タ転送のイベントの出⼒、など• 0（ゼロ）を指定すると機能を抑制（デフォルト）GPUキューの実⾏• インタラクティブ実⾏

qsub –I –q G –l select=1:ncpus=N2:ngpus=N4– N2: CPUコア数（1〜40），N4:GPUディバイス数（1〜8）– Ngpus=N4を指定しない場合，GPUディバイスは利⽤しない．
– GPUを搭載したノードにてプログラムを実⾏。

GPUキューの実⾏• ホストプログラムの並列化に応じてジョブスクリプトを作成．
(以下の例は，ホストが並列化無しの場合）#!/bin/bash#PBS␣-q␣G#PBS␣-I␣select=1:ncpus=4:ngpus=1 ・利用するGPUカード数を指定#PBS␣-N␣job_acc#PBS␣-j␣oesource␣/etc/profile.d/modules.shmodule␣purge
・現在のモジュール環境を全て破棄
module␣load␣nvhpc
・NVIDIA HPCコンパイラの環境設定
cd␣${PBS_O_WORKDIR}./a.out

OpenACCの演習 1• OpenMPプログラム(pi_omp.c)をNVIDIA HPCコンパイラーでコンパイルし，スレッド並列を実⾏せよ．
• OpenACCプログラム(pi_acc.c)をコンパイルし，GPU計算を実⾏せよ．スレッド並列とGPU計算の計算時間を⽐較せよ．
* 注）pi_acc.cのnmaxは40000程度以下に設定してください．

• 環境変数（NVCOMPILER_ACC_TIME or
NVCOMPILER_ACC_NOTIFY）を設定してディバイスの実⾏プロファイルを確認しなさい．

OpenACCの演習 2• pi2_omp.c と pi2_acc.c をそれぞれ実⾏し，先の演習の結果と⽐較しなさい．
– pi2_acc.cのディバイス実⾏プロファイルを分析し，全体の計算時間が短縮されない原因について考えなさい．

pi_acc.cの実⾏プロファイル
Accelerator Kernel Timing data
main NVIDIA devicenum=0 time(us): 41
16: compute region reached 1 time （compute構文が１回，kernel構文・reduction演算）
18: kernel launched 1 time
grid: [65535] block: [128]elapsed time(us): total=12,175 max=12,175 min=12,175 avg=12,175
18: reduction kernel launched 1 time
grid: [1] block: [256]elapsed time(us): total=70 max=70 min=70 avg=70
16: data region reached 2 times （data構文が２回，copyin&copyout）
16: data copyin transfers: 1
device time(us): total=8 max=8 min=8 avg=8
23: data copyout transfers: 1
device time(us): total=33 max=33 min=33 avg=33

pi2_acc.cの実⾏プロファイル
main NVIDIA devicenum=0 time(us): 2,026,328
２回のdata構⽂が17: compute region reached 1 time
ボトルネック!!
19: kernel launched 1 timegrid: [65535] block: [128]elapsed time(us): total=19,479 max=19,479 min=19,479 avg=19,479
17: data region reached 2 times
24: data copyout transfers: 763
device time(us): total=987,059 max=1,393 min=1,235 avg=1,293
29: compute region reached 1 time
31: kernel launched 1 time
grid: [65535] block: [128]elapsed time(us): total=15,533 max=15,533 min=15,533 avg=15,533
31: reduction kernel launched 1 time
grid: [1] block: [256]elapsed time(us): total=114 max=114 min=114 avg=114
29: data region reached 2 times
29: data copyin transfers: 764
device time(us): total=1,039,257 max=1,372 min=19 avg=1,360
34: data copyout transfers: 1
device time(us): total=12 max=12 min=12 avg=12

DATA構⽂• #pragma␣acc␣data␣節• 節には以下のようなものがある．（Kernel構⽂の節にも使える．）– copyin(a,b): 配列a, bをホストからアクセラレータにコピーする．
– copyout(c): 配列cをアクセラレータからホストにコピーする．
– copy(A)︓copyinとcpyoutの双⽅を⾏う．
– create(B): アクセラレータでローカルに使⽤する配列Bの領域を作成．
– present(C): 対象となる処理に⼊ったときにすでにデバイスメモリにある配列Cのデータを使⽤する．

DATA構⽂• pi2_acc.cのプログラムに対して，1. GPU内だけで必要な変数を予め作成する．
#pragma acc data create(f,x,y){.......
}
2. Kernel構⽂でデータ属性を指⽰する．
– kernels present(f) copyin(dh)– kernels present(f) copy(sum)修正pi2_acc.cの実⾏プロファイル
main NVIDIA devicenum=0 time(us): 41
17: data region reached 2 times
data構⽂のボトルネックが解消
20: compute region reached 1 time
スカラ変数のデータコピーのみ
22: kernel launched 1 time
grid: [65535] block: [128]elapsed time(us): total=162,324 max=162,324 min=162,324 avg=162,324
20: data region reached 2 times
20: data copyin transfers: 1
device time(us): total=15 max=15 min=15 avg=15
32: compute region reached 1 time
34: kernel launched 1 time
grid: [65535] block: [128]elapsed time(us): total=15,465 max=15,465 min=15,465 avg=15,465
34: reduction kernel launched 1 time
grid: [1] block: [256]elapsed time(us): total=107 max=107 min=107 avg=107
32: data region reached 2 times
32: data copyin transfers: 1
device time(us): total=7 max=7 min=7 avg=7
37: data copyout transfers: 1
device time(us): total=19 max=19 min=19 avg=19

OpenACCの演習3
1. laplace_omp.c をもとに，ラプラス⽅程式のOpenACCのプログラムを作成せよ．
ü OpenMPのディレクティブを全てOpenACCのディレクティブに変更．
#pragma omp parallel for => #pragma acc kernels
ü データ構⽂を使って，データコピーの回数が少なくなるように⼯夫する．

2. ベンチマークを実⾏してOpenMPとOpenACCで計算時間を⽐較せよ．

OpenACCの⾃習問題• ラプラス⽅程式のガウスザイデル法，Even Odd法のプログラム(laplace_gs.c, laplace_gs_eo.c)についてOpenACCのプログラムを作成せよ．

ü ガウスザイデル法の反復式では，両辺のデータに依存関係がある．⼀⽅，Even Odd法では，この依存関係を解消するアルゴリズムとなっている．
ü GPUでも，OpenMPやベクトル演算と同様に依存関係のない式において⾼速化が達成される．

OpenACC Runtime関数• #include <openacc.h>- ディバイスタイプを取得
acc_device_t acc_get_device_type(void)- 利⽤するディバイスの総数を取得
int acc_get_num_devices(acc_device_t t)- 利⽤するディバイスの番号をセット
void acc_set_device_num(int n, acc_device_t t)- などなどMPI / OpenACC• Module環境
module load mpt• コンパイルnvc –acc –Mmpi=sgimpi prog.c –o prog.out• 実⾏ (4プロセス = 4GPU）#PBS –l select=1:ncpus=4:mpiprocs=4:ngpus=4…
mpiexec_mpt ./prog.out