第205回 お試しアカウント付き
並列プログラミング講習会「Wisteria実践」
東京⼤学情報基盤センター
塙 敏博
三⽊ 洋平
2023/2/14

講習会「Wisteria実践」

1

講習会概略• 開催⽇：
2023年6⽉8⽇（木）10：30 - 17：45• 形態：ZoomとSlackを⽤いたオンライン講習会• 講習会プログラム： 講師：塙• 10：30 ‒ 12：00 Wisteria/BDEC-01システム紹介• 13：30 ‒ 15：15 Odysseyノード：A64FXにおけるOpenMP並列化、
MPI+OpenMPハイブリッド並列、性能分析（講義+演習）• 15：30 - 17：15 Aquariusノード: A100 GPUの利⽤、OpenACC、MPI並列化、性能分析（講義+演習）• 17：30 ‒Q&A

2023/2/14

講習会「Wisteria実践」

2

お願い等
• ハンズオンのためのPC，Zoom及びスパコンへ接続するためのネット
ワーク環境は各受講者でご準備ください。
• PCは Windows/Microsoft Update，Apple Security Updateなどで最新のセキュリティアップデートを⾏ってください。
• 必ずウィルス対策ソフトウェアをインストールし，ウィルス検索を実⾏して問題がないことを事前に確認してから受講してください。
• セキュリティ対策未実施の場合はオンライン講習会受講を認めません。

• OSは、Windows、Macどちらでも構いませんが、SSHを⽤いてセンターのスーパーコンピューターへ接続ができることが必要です（後述）。
• 演習の実施に当たり，受講⽣にセンターのスーパーコンピューターを1⽉間利⽤できる無料アカウント（お試しアカウント）を発⾏します。
2023/2/14

講習会「Wisteria実践」

3

Zoom関連• 「⼿をあげる」機能• 質問がある際，全体の状況を確認するため使⽤• ブレークアウトセッション• 画⾯を共有しながらエラー対応する際に使⽤• （なるべく⼝頭でのやりとりやSlackで対応する予定）• https://utelecon.adm.u-tokyo.ac.jp/zoom/how_to_use

2023/2/14

講習会「Wisteria実践」

4

「⼿を挙げる」⽅法
1. Zoomメニュー中の「リアクション」をクリック

2. ポップアップで表⽰された「⼿を挙げる」をクリック

2023/2/14

講習会「Wisteria実践」

5

⼿が挙がっていることの確認⽅法
1. Zoomメニュー中の「参加者」をクリックして，参加者⼀覧を表⽰
2. 表⽰された参加者⼀覧の，⾃分のところを⾒ると⼿が挙がっている

2023/2/14

講習会「Wisteria実践」

6

「⼿を降ろす」⽅法
1. Zoomメニュー中の「リアクション」をクリック

2. ポップアップで表⽰された「⼿を降ろす」をクリック

2023/2/14

講習会「Wisteria実践」

7

ブレイクアウトルーム（1/6）• 演習時に使⽤するかもしれません• 演習中に「ヘルプを求める」ことができます• ホストを招待した後に「画⾯を共有」することで，皆さんの記述したプログラムを⼀緒に⾒ながら問題解決にあたります• Zoomメニュー中の「ブレイクアウトルーム」をクリック

2023/2/14

講習会「Wisteria実践」

8

ブレイクアウトルーム（2/6）• 進⾏中のブレイクアウトルームのリストが表⽰されるので，空いている部屋に「参加」してください• 左の例では5部屋がすべて空室，右の例ではルーム1のみ参加者がいる

2023/2/14

講習会「Wisteria実践」

9

ブレイクアウトルーム（3/6）• 「参加」をクリックすると確認画⾯が出てくるので，「はい」を選択すると⼊室できます

2023/2/14

講習会「Wisteria実践」

10ブレイクアウトルーム（4/6）• 再度メニュー中の「ブレイクアウトルーム」をクリックすると，「ヘルプを求める」が増えているのでクリックしてください

2023/2/14

講習会「Wisteria実践」

11ブレイクアウトルーム（5/6）• ポップアップで出てきた「ホストを招待」をクリック• ホスト（講師）の参加待ちに移⾏します（画⾯はそのまま）• 他の受講者のヘルプ中など，直ちに対応できない場合もあります

2023/2/14

講習会「Wisteria実践」

12

ブレイクアウトルーム（6/6）• 問題解決後は，Zoomメニュー中の「ルームを退出する」• 「ブレイクアウトルームを退出」が表⽰されるのでクリックして，元の講習会会場にお戻りください• 間違えて「ミーティングを退出」すると講習会から退出します

2023/2/14

講習会「Wisteria実践」

13

Slack関連• ブラウザ上で使う場合には：• https://w1590055008bgo338004.slack.com/archives/C05ASCUE293• 注：ログインには，事前にお配りしたリンクからの登録が必要です• 質問対応に使⽤• コードの貼り付け⽅• スレッドの確認⽅法• 以下，ブラウザ版で説明しますがアプリ版でも操作は同じです
2023/2/14

講習会「Wisteria実践」

14質疑応答チャンネルへの移動• 左側のメニューバーのチャンネル⼀覧内に「第170回-」があるので，クリック• 表⽰されていない場合
1. 「チャンネルを追加する」をクリック
2. 「チャンネル⼀覧を確認する」をクリック
3. 「第205回-wisteria実践」があるので，「参加する」
をクリック•

第199回〜それ以前とお間違えなくメッセージの⼊⼒⽅法• 最下部に⼊⼒欄があるので，質問内容を記載して Ctrl+Enter• ⼊⼒後に右下の「メッセージを送信する」をクリックしても同じ（メッセージ⼊⼒前には，「メッセージを送信する」は押せない）
メッセージの⼊⼒欄
コードブロックの⽣成

メッセージを送信する• コードを⼊⼒する際には，「コードブロック」がおすすめ• 枠が⽣成されるので，この中にコピペするのが簡単かつ⾒やすい• ```（JIS配列ならばShift+@を3連打）しても枠が⽣成される
2023/2/14

講習会「Wisteria実践」

16

⾃分が参加したスレッドのみの表⽰⽅法• 左上の「スレッド」をクリックすると，⾃分が参加しているスレッドの⼀覧が表⽰されます• 質問内容には，「スレッドで返信する」形式で回答するので，⾃分が聞いた内容のみが表⽰できます2023/2/14

講習会「Wisteria実践」

17

Wisteria/BDEC-01の紹介
別資料参照

2023/2/14

講習会「Wisteria実践」

18

19
19

2023/2/14

参考：ジョブクラス：インタラクティヴジョブ（トークン消費無し）
https://www.cc.u-tokyo.ac.jp/supercomputer/wisteria/service/job.php• Wisteria-O（Odyssey）:シミュレーションノード群
ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

1
2 ～ 12

30 分
10 分

28
28

ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

interactive-a

1ノード

10 分

448

share-interactive

1GPU

10 分

56

キュー名

interactive-o(interactive-o_n1)(interactive-o_n12)• Wisteria-A（Aquarius）:データ・学習ノード群
キュー名

講習会「Wisteria実践」

参考：インタラクティブ実⾏のやり⽅（本講習会では使えません）
Odyssey• 1ノード実⾏$ pjsub --interact –g グループ名 –L rg=interactive-o,elapse=01:00• 12ノード実⾏$ pjsub --interact –g グループ名 –L rg=interactive-o,node=12,elapse=01:00
Aquarius• 1GPU実⾏$ pjsub --interact –g グループ名 –L rg=share-interactive,elapse=01:00• 1ノード(8GPU)実⾏$ pjsub --interact –g グループ名 –L rg=interactive-a,elapse=01:00※インタラクティブ用のノードがすべて使われている場合、資源が空くまでログインできません。
※このアカウントでは使えません。
2023/2/14

講習会「Wisteria実践」

20

21
21

2023/2/14

参考：ジョブクラス：バッチジョブ（Odyssey）
https://www.cc.u-tokyo.ac.jp/supercomputer/wisteria/service/job.php
ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

debug-o1 ～ 144

30 分

28

short-o1 ～ 728 時間

28

1 ～ 144145 ～ 576
577 ～ 1152
1153 ～ 2304

48 時間
〃
〃
24 時間

28
〃
〃
〃

1 ～ 288

48 時間

28

キュー名

regular-o(small-o)(medium-o)(large-o)(x-large-o)priority-o （優先実行，トークン消費量 1.5倍）

講習会「Wisteria実践」

22
222023/2/14

参考：ジョブクラス：バッチジョブ（Aquarius）
https://www.cc.u-tokyo.ac.jp/supercomputer/wisteria/service/job.php
ノード数
GPU数

制限時間
（Elapsed）

メモリ容量
（GB）

debug-a

1 ノード

30 分

448

short-a1 ～ 2 ノード

2 時間

448

regular-a(small-a)(medium-a)(large-a)1 ～ 2 ノード3 ～ 4 ノード5 ～ 8 ノード

48 時間
〃
24 時間

448
〃
〃

share-debug

1, 2, 4 GPU

30 分

56

share-short

1, 2, 4 GPU

2 時間

56

share
(share-1)(share-2)(share-4)

1 GPU
2 GPU
4 GPU

48 時間
〃
24 時間

56
〃
〃

キュー名

講習会「Wisteria実践」

23
23

2023/2/14

参考：ジョブクラス：プリポストサービス
ログインノードと同じアーキテクチャのノードを使用
https://www.cc.u-tokyo.ac.jp/supercomputer/wisteria/service/job.php
ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

prepost （予約無し）

1

6 時間

340

prepost1_n1～
prepost4_n1

1

1 ～ 6 時間

340

prepost1_n4

1～4

1 ～ 6 時間

340

prepost1_n8

1～8

1 ～ 6 時間

340

キュー名

講習会「Wisteria実践」

24
24

2023/2/14

教育利用のジョブクラス：バッチジョブ（Odyssey）
ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

lecture-o（教育利用全体で共有）

1 ～ 12

15 分

28

tutorial-o（この講習会の時間中のみ使用可能）

1 ～ 12

15 分

28

キュー名

ノード当たり 48コア è 最大576コアまで使用可能！

講習会「Wisteria実践」

25
25

2023/2/14

教育利用のジョブクラス：バッチジョブ（Aquarius）
キュー名

lecture-a（教育利用全体で共有）
tutorial-a（この講習会の時間中のみ使用可能）

GPU数

制限時間
（Elapsed）

メモリ容量
（GB）

1,2,4

15 分

56/GPU

1,2,4

15 分

56/GPU

share-aと同様の使い方講習会「Wisteria実践」

「Wisteria実践」について

2023/2/14

講習会「Wisteria実践」

26

本講習会の⽬的
並列アルゴリズムはひとまず置いといて…

• Odysseyノードの多数のコアの扱い⽅• Odysseyの多数の計算ノードによる効率の良い実⾏⽅法• 富⼠通コンパイラ、ジョブ環境のTIPS• AquariusノードのGPUの扱い⽅等について学ぼう！

2023/2/14

講習会「Wisteria実践」

27

（現代の）スパコンは並列性が命• 京コンピュータ(2011~2019)： 88,128ノード x 8コア/ノード• Oakforest-PACS : 8,208ノード x 68コア/ノード• Sunway TaihuLight: 40,960ノード x 260コア/ノード• Summit : 約4,600ノード x 6 GPU/ノード• 富岳: 158,976ノード x 48コア/ノード• Frontier: 9,248ノード x (64コア + 4 GPU)/ノード• ABCI: 1,088ノードx 4 GPU/ノード+120ノード x 8 GPU/ノード• Wisteria: 7,680ノード x 48コア/ノード + 45ノード x(72コア+8 GPU)/ノード• 1コア当たりではスパコンの⽅が遅いこともある• スパコン向けは，電⼒効率を狙ってクロック周波数を下げている
è その分コア数を増やす
2023/2/14

講習会「Wisteria実践」

28

並列プログラムの実⾏主体• マルチスレッド

共有メモリで動作→並列化の指⽰だけでOK• OpenMP• ユーザが並列化指⽰⾏を記述• マルチプロセス• MPI （Message Passing Interface）• ユーザがデータ分割⽅法を明⽰的に記述• ノードが物理的に異なる• 同⼀ノードでもプロセス毎にメモリは論理的に分離→明⽰的な通信が必要

マルチプロセスとマルチスレッドは共存可能→ハイブリッドMPI/OpenMP実⾏
2023/2/14

講習会「Wisteria実践」

29

Weak ScalingとStrong Scaling
並列処理においてシステム規模を⼤きくする⽅法

MPI通信もポイント
Odyssey: 低レイテンシ
Aquarius: ⾼バンド幅• Weak Scaling: それぞれの問題サイズは変えず並列度をあげる
Ø全体の問題サイズが（並列数に⽐例して）⼤きくなる
Ø通信のオーバヘッドはあまり変わらないか、やや増加する• Strong Scaling: 全体の問題サイズを変えずに並列度をあげる
Ø問題サイズが装置数に反⽐例して⼩さくなるØ通信のオーバヘッドは相対的に⼤きくなる

Weak Scaling

Strong Scalingも重要

それまで解けなかった

同じ問題規模で、短時間に結果を得る（より難しい）

規模の問題が解ける
2023/2/14

講習会「Wisteria実践」

30

並列プログラミングについて学習したい⽅へ
以下の講習会資料を参考にしてください（初級è上級）
https://www.cc.u-tokyo.ac.jp/events/lectures/• 「MPI基礎：並列プログラミング⼊⾨」•

https://www.cc.u-tokyo.ac.jp/events/lectures/203/• 「OpenACCとMPIによるマルチGPUプログラミング⼊⾨」• https://www.cc.u-tokyo.ac.jp/events/lectures/209/• 「OpenMPによるマルチコア・メニィコア並列プログラミング⼊⾨」•

https://www.cc.u-tokyo.ac.jp/events/lectures/192/• 「GPUプログラミング⼊⾨」•

https://www.cc.u-tokyo.ac.jp/events/lectures/206/• 「並列有限要素法で学ぶ並列プログラミング徹底⼊⾨」•

https://www.cc.u-tokyo.ac.jp/events/lectures/194/• 「OpenMPで並列化されたC++プログラムのGPU移植⼿法」•

https://www.cc.u-tokyo.ac.jp/events/lectures/208/• 「MPI+OpenMPで並列化されたFortranプログラムのGPUへの移⾏⼿法」•

https://www.cc.u-tokyo.ac.jp/events/lectures/211/「こんな講習会が欲しい！」という要望があればぜひアンケートへ！
2023/2/14

講習会「Wisteria実践」

31

Wisteria-Odyssey
A64FXの基礎

2023/2/14

講習会「Wisteria実践」

32

A64FXプロセッサ
TofuD コントローラ

PCIe コントローラ

HBM2

1

2

3

4

5

6

7

8

9

10 11

0

1

2

3

4

5

6

7

8

9

10 11

リ
ン
グ
バ
ス

0

1

2

3

4

5

6

7

8

9

10 11

0

1

2

3

4

5

6

7

8

9

10 11

HBM2

0

CMG#1

命令セット

CMG#2
HBM2

HBM2

CMG#0

プロセッサ名Fujitsu A64FX
Arm v8.2+SVE (Scalable Vector Extension)

プロセッサ数 (コア数)周波数
理論演算性能
メモリ容量
メモリ帯域幅1 (48+アシスタントコア2 or 4)
2.2 GHz
3.3792 TFLOPS
32 GB
1,024 GB/s

CMG#3

• 4つのCMG（Core Memory Group），12計算コア/CMG• NUMAアーキテクチャ（Non-Uniform Memory Access）
ü メモリは各CMGに搭載されていて独立，異なるCMGのローカルメモリ上のデータをアクセスすることは可能ü ローカルメモリ上のデータを使って計算するのが効率的• 大規模並列：各CMGに1-MPIプロセス（12-OpenMPスレッド），プロセッサ内4プロ
セスのハイブリッド推奨
2023/2/14

講習会「Wisteria実践」

33

A64FX：CMG（Core Memory Group）
TofuD コントローラ

PCIe コントローラ

0

1

2

3

4

5

6

7

8

9

10 11

0

1

2

3

4

5

6

7

8

9

10 11

CMG#1

2023/2/14

リ
ン
グ
バ
ス

講習会「Wisteria実践」

0

1

2

3

4

5

6

7

8

9

10 11

0

1

2

3

4

5

6

7

8

9

10 11

CMG#3

HBM2

HBM2

CMG#2
HBM2

HBM2

CMG#0

34

NUMAドメイン (1/3)$ numactl -H
available: 8 nodes (0-7)
node 0 cpus:
node 0 size: 723 MB
node 0 free: 704 MB
node 1 cpus: 0
node 1 size: 765 MB
node 1 free: 394 MB
node 2 cpus: 1

アシスタントコア（OS専⽤、ユーザは使えない）

node 2 size: 765 MB
node 2 free: 62 MB

node 4 cpus: 12 13 14 15 16 17 18 19 20 21 22 23
node 4 size: 7345 MB
node 4 free: 7205 MB
node 5 cpus: 24 25 26 27 28 29 30 31 32 33 34 35
node 5 size: 7409 MB
node 5 free: 7320 MB
node 6 cpus: 36 37 38 39 40 41 42 43 44 45 46 47
node 6 size: 7409 MB
node 6 free: 7320 MB
node 7 cpus: 48 49 50 51 52 53 54 55 56 57 58 59
node 7 size: 7402 MB
node 7 free: 7302 MB

node 3 cpus:
node 3 size: 759 MB
node 3 free: 678 MB
2023/2/14

講習会「Wisteria実践」

35

NUMAドメイン（2/3）
CMG#2

12 13 14 15

36 37 38 39

16 17 18 19
20 21 22 23
24 25 26 27
28 29 30 31

リ
ン
グ
バ
ス

32 33 34 35

CMG#1

2023/2/14

40 41 42 43
44 45 46 47
48 49 50 51
52 53 54 55
56 57 58 59

講習会「Wisteria実践」

CMG#3

Memory #7

Memory #5

CMG#0

Memory #6

Memory #4• CMG:#0-#3，Core:#12-59，Memory:#4-#7

36

NUMAドメイン（3/3）
CMG#2

13 14 15
12numactl

37 38 39
36numactl

18 19
16-C1712-23

-m 22
4 23
20 21
25 26 27
24numactl
30 31
28-C2924-35

リ
ン
グ
バ
ス

CMG#1

-m 46
6 47
44 45
49 50 51
48numactl
54 55
52-C5348-59

-m 58
7 59
56 57

-m 34
5 35
32 33
2023/2/14

42 43
40-C4136-47

講習会「Wisteria実践」

CMG#3

Memory #7

Memory #5

CMG#0

Memory #6

Memory #4• CMG:#0-#3，Core:#12-59，Memory:#4-#7

37

Stream Copy性能 (Odyssey, 1.2.36, C)
900
800
700

バンド幅 (GB/sec)• close: 同⼀CMG優先• spread: CMG分散• trad: tradモード• zfill: zfillオプション• demand: ページ
割り当て設定(後述)• clang: clangモード(tradのオプション+α (後述))close trad
close trad, zfill
close trad, deman dclose trad, zfill,demandclose clang
spread trad
spread trad, zfill
spread trad, demand
spread trad, zfill,deman dspread clang

600
500
400
300
200
100
0
0

2023/2/14

5

10

15

講習会「Wisteria実践」

20

25

コア数

30

35

40

45
38

50

Stream Copy性能 (Odyssey, 1.2.36, Fortran)
900
800
700

バンド幅 (GB/sec)• close: 同⼀CMG優先• spread: CMG分散• trad: tradモード• zfill: zfillオプション• demand: ページ
割り当て設定(後述)• C⾔語とほぼ同じ傾向

close std

close zfill

close std,demandclose +fjomplib
spread std

close zfill,demandclose +fjomplib
spread zfill

spread std,demand
spread +fjomplib

spread zfill,demand

600
500
400
300
200
100
0
0

2023/2/14

5

10

講習会「Wisteria実践」

15

20

25

コア数

30

35

40

45
39

50

Stream Copy性能まとめ (Odyssey, 1.2.36)• 性能が最も⾼い組み合わせ (~820 GB)• コンパイル(C, Fortran)：-Kfast,openmp,zfill
実⾏時：export XOS_MMM_L_PAGING_POLICY=demand:demand:demand• コンパイル(C:Clangモード)：-Kfast,openmp,zfill -Nclang -msve-vector-bits=512 -fno-builtin
実⾏時：export XOS_MMM_L_PAGING_POLICY=demand:demand:demand• XOS_MMM_L_PAGING_POLICY=demand:demand:demandが無いと、コンパイルオプションを上の通りにしても ~ 200 GB/s
• zfillがないと ~560GB/s• zfill: 書き込みのみの配列をメモリを参照せずキャッシュに確保• 通常のコードではプリフェッチが効かず逆効果になることも多いので注意• libomp(デフォルト)の場合、spreadの⽅が⾼い性能• fjomplibの場合、spread配置を実現するにはコア番号を明⽰する必要
2023/2/14

講習会「Wisteria実践」

40

コンパイラの種類と実⾏(Odyssey)• ログインノードとOdysseyの計算ノードとで、CPUの命令セットが⼤きく異なる• ログインノード：命令セットアーキテクチャ Intel CascadeLake＋AVX512, x86_64• Odyssey計算ノード： Fujitsu A64FX, 命令セットアーキテクチャ ARM v8.2 + SVE, aarch64

Odysseyノード

ログインノード
バッチ処理
x86_64
ソース
コード

クロス
コンパイラ

aarch64

aarch64

ネイティブ
コンパイラ

a.out

インタラクティブ実⾏• 富⼠通製コンパイラ：module load odyssey• ネイティブコンパイラ： ＜コンパイラの種類名＞• クロスコンパイラ： ＜コンパイラの種類名＞+px• MPI: mpi+コンパイラ名 (例：mpifccpx)
2023/2/14

ソース
コード

講習会「Wisteria実践」

a.out

aarch64

⾔語

ネイティブ
コンパイラ

クロス
コンパイラ

C

fcc

fccpx

C++

FCC

FCCpx

Fortran

frt

frtpx

41

コンパイラ内部動作の切り替え (Odyssey)• OpenMPランタイムライブラリ選択• -Nlibomp : デフォルト (LLVM由来)• -Nfjomplib : 富⼠通OpenMPライブラリ• モード (C⾔語のみ)• -Nnoclang : tradモード (デフォルト)• -Nclang : clangモード(LLVMベース)• libompのとき• clangモードの⽅が性能改善が顕著であるので、試す価値がある• バグも多いかもしれない• libompしか使えない• デフォルトはソフトウェアバリア、ハード
ウェアバリアを試すのが良い（ただしCMG
内のみ有効、バグもあるかも）
export FLIB_BARRIER=HARD• 詳細は• 「プログラミングガイド プログラミング共通編」• p.50-57 2つのOpenMPライブラリ• p.57 OpenMP microbenchmark• fjomplibの⽅が最⼤5倍速いが、OpenMP3.1までしか対応しない = task⽂はlibompを使う

2023/2/14

可能な
組み合わせ

C / C++

Fortrantradモードclangモード

libomp○(default)○ (default)○ (default)

fjomplib○

×○

講習会「Wisteria実践」

42

基本的なコンパイルオプション (Odyssey)• -Kocl• -Kfast,openmp• ⾼速なプログラムを⽣成• -O3 + いろいろ
• 後者はOpenMPを使う場合に指定• 富⼠通の拡張機能を使う場合• -Koptmsg=2• 最適化レポート• いつもつけておくことをお勧めします• -Kzfill[=N]• キャッシュに書き込み領域を直接確保• 明らかにメモリバンド幅ネックな場合に有効• 本来は以下でループ毎に明⽰すべき• 詳細はWebで…
ポータルèドキュメント閲覧

C: #pragma loop [no]zfill [N]Fortran: !ocl loop [no]zfill [N]-Koclオプションが必要
2023/2/14

講習会「Wisteria実践」• C⾔語使⽤⼿引書• C++⾔語使⽤⼿引書• Fortran使⽤⼿引書
43基本的なコンパイルオプション (Odyssey, C, -Nclang)• -msve-vector-bits=512• 詳細はWebで…
ポータルèドキュメント閲覧• SVEベクトル⻑を512bitに固定• A64FXではこれが最適値(デフォルトではscalable)• C⾔語使⽤⼿引書• C++⾔語使⽤⼿引書• -ffj-zfill[=N]• キャッシュに書き込み領域を直接確保• 明らかにメモリバンド幅ネックな場合に有効• 本来は以下でループ毎に明⽰すべき#pragma fj loop [no]zfill [N]• -fno-builtin併⽤が必要な場合もある• -ffj-ocl• 富⼠通の拡張命令を使う場合
2023/2/14

講習会「Wisteria実践」

44

First touch• NUMAにおけるFirst touch• 初めてデータにアクセスする際に、コアが属するドメインローカルのメモリにデータが配置される
初期化等の際にも、実際の計算と同じアクセスパターンで実⾏する必要あり

C(:)=0 ç全てコア0に配置されてしまう!$omp parallel dodo i=1,100
C(i) = C(i)+A(i)enddo!$omp parallel dodo i=1, 100
C(i)=0
enddo!$omp parallel dodo i=1, 100
C(i) = C(i)+A(i)enddo• A64FXでもCMGを超える際の影響は⼤きい• numactlで正しいドメイン(コアとメモリの組）を指定しておく必要がある• 実⾏時に必要: export XOS_MMM_L_PAGING_POLICY=demand:demand:demand
2023/2/14

講習会「Wisteria実践」

45

メモリのアラインメント• SVE命令で512ビットデータを有効に使うためには、データがメモリ境界に揃っていることが必要• 512ビット = 64バイトの倍数のアドレスにアクセスするように• 境界以外では複数回のアクセスになる• 例：倍精度実数(double)を連続8個(=512ビット)読み出す場合
境界からずれているため、実際には複数回アクセスになる

アラインされたアクセス=効率がいい
0

1

2

3

4

5

6

7

0

1

2

3

4

5

6

7

0

1

2

3

4

5

6

64の倍数の境界
2023/2/14

講習会「Wisteria実践」

46

7

メモリアラインメントの指定• C⾔語• __attribute__((aligned(64))) を指定 (64 byte = 512 bit)
• 例： double A[1024] __attribute__((aligned(64)))• 動的確保の際は posix_memalign() を使⽤• 上と同様の例：#include <stdlib.h>double *A;
posix_memalign(&A, 64, 8192);• Fortran• 実⾏時に FLIB_ALLOC_ALIGN 環境変数を指定
export FLIB_ALLOC_ALIGN=64• Allocate⽂に適⽤される
2023/2/14

講習会「Wisteria実践」

47

ポインタ指定の制限緩和 (restrict, C⾔語)
• ポインタを経由すると最適化の制限がかかる• アドレスがsourceとdestinationで重複した場合，ベクトル化するとおかしくなる(aliasing)• restrict修飾⼦でaliasingはないことを宣⾔
void mycopy (double * restrict a, double * restrict b, int n){for(int i=0; i<n; i++)b[i] = a[i];}• またはコンパイル時にtradモードでは -Krestp を指定• clangモードではrestrictを指定する以外⽅法がない
2023/2/14

講習会「Wisteria実践」

48

環境変数
• ページ割り当て⽅式：NUMAをうまく使うには、以下の設定が必要
export XOS_MMM_L_PAGING_POLICY=demand:demand:demand• ラージページかつデマンドページング（3つのメモリセクションともに）• 他にもメモリ関連、OpenMP関連などで多様なオプションがある• 参考：• ジョブ運⽤ソフトウェア エンドユーザ向けガイド HPC拡張機能編

2023/2/14

講習会「Wisteria実践」

49

まとめ：逐次実⾏レベル
1. CMGをまたぐ処理は性能が低下する場合が多い。MPIの併⽤を検討
しましょう。
2. ラージページ＋デマンドページングのオプションをまずは⼊れましょう。
export XOS_MMM_L_PAGING_POLICY=demand:demand:demand
3. First Touchは必須。これも2.の設定がないと機能しません。
4. メモリのアラインメントをなるべく64 byteに揃える。

2023/2/14

講習会「Wisteria実践」

50

最適化の弊害(?)• あるパターンを見つけるとライブラリ関数への置き換えを勝手にやってしまう•for(i = 0; i < N; i++)B[i] = A[i];#pragma omp parallel for private(j,k)
for(i=0; i<N; i++)for(j=0; j<N; j++)for(k=0; k<N; k++)C[i][j]+=A[i][k]*B[k][j]2023/6/8

memcpyClangモードの場合に起こる(zfillが効かなくなる)è -fno-builtinで抑制

matmul• BLASが呼ばれるが少しでも形が崩れると適用されない• スカラ変数に代入したりとか• 講習会・講義で演習をしていて何も工夫しない方が早い、という現象が起きる、、• -Kopenmpつけずにコンパイルしても、BLASがOpenMP
並列化されているため、OMP_NUM_THREADSに連動
FX1000 運用技術検討会

51コンパイルオプション• -Kfast,openmp,restp=all -x0 -Koptmsg=2• -Krestp=all: 全てにrestrictを仮定する è これがないとそもそも最適化が効かない• -x0: インライン展開を無効にする è こうしないとmatmulパターンを見つけてくれない• Optimization messages
jwd6002s-i "mat-mat-openmp.c", line 35: このループをSIMD化しました。
jwd8204o-i "mat-mat-openmp.c", line 35: ループにソフトウェアパイプライニングを適用しました。
jwd8205o-i "mat-mat-openmp.c", line 35: ループの繰返し数が48回以上の時、ソフトウェアパイプライニングを適用したループが実行時に選択されます。
jwd8663o-i "mat-mat-openmp.c", line 35: ソフトウェアパイプライニングの効果がないループと判断したため、ソフトウェアパイプライニングを抑止しました。
jwd8202o-i "mat-mat-openmp.c", line 35: このループを展開数2回でループアンローリングしました。
jwd8330o-i
jwd8101o-i
jwd8209o-i
jwd6131s-i
jwd8671o-i
jwd6004s-i
jwd8204o-i
jwd8205o-i
jwd8663o-i
jwd6208s-i
jwd8208o-i"mat-mat-openmp.c", line 36: ループ制御変数'i',...,'j'の多重ループを1重化しました。
"mat-mat-openmp.c", line 57: 利用者定義の関数'MyMatMat'をインライン展開しました。
"mat-mat-openmp.c", line 67: 多項式の演算順序を変更しました。
"mat-mat-openmp.c", line 74: ループ出口が2箇所以上あるため、このループはSIMD化できません。
"mat-mat-openmp.c", line 74: ループの形状が最適化の対象外であるため、ソフトウェアパイプライニングを適用できません。
"mat-mat-openmp.c", line 95: リダクション演算を含むループ制御変数'k'のループをSIMD化しました。
"mat-mat-openmp.c", line 95: ループにソフトウェアパイプライニングを適用しました。
"mat-mat-openmp.c", line 95: ループの繰返し数が256回以上の時、ソフトウェアパイプライニングを適用したループが実行時に選択されます。
"mat-mat-openmp.c", line 95: ソフトウェアパイプライニングの効果がないループと判断したため、ソフトウェアパイプライニングを抑止しました。
"mat-mat-openmp.c", line 96: 変数'C'を定義・参照する順序がわからないため、定義・参照する順序が逐次実行と変わる可能性があり、このループはSIMD化できません。
"mat-mat-openmp.c", line 96: ループ内の総和または乗積演算の計算方法を変更しました。

jwd8220o-i
jwd8331o-i

2023/6/8"mat-mat-openmp.c", line 93: 副作用の可能性のある最適化を行いました。
"mat-mat-openmp.c", line 93: このループをライブラリ呼出し(matmul)に変換しました。

FX1000 運用技術検討会

52

OpenMPによるSIMD化

2023/2/14

講習会「Wisteria実践」

53

OpenMP+SIMD

OpenMPの基本的な内容については、以下の講習会「OpenMPによるマルチコア・メニィコア
並列プログラミング⼊⾨」参照
https://www.cc.u-tokyo.ac.jp/events/lectures/154/• SIMD (Single Instruction - Multiple Data)• 1命令で複数データを処理• SVE命令：1命令で倍精度浮動⼩数点数8個を処理• A64FXではSVE命令を最⼤限活⽤しなければ⾼い性能が得られない• SIMD節はOpenMP 4.0からサポート• コンパイラが⾃動ベクトル化もしてくれるが，指⽰をした⽅が確実• コンパイラが誤った答えを出すこともあるので要確認（経験済み）• コードの書き換えを必要とする場合も（しばしば）ある• メモリのアラインメントが重要（前述）2023/2/14

講習会「Wisteria実践」

54

OMP For/Do+SIMD• C⾔語：#pragma omp parallel for simdfor (i=0; i < num; i++) {sum[i] = sum[i] + a[i];}• Fortran:!$omp parallel do simd
do i=1, numsum(i) = sum(i) + a(i)!$omp end parallel
2023/2/14

講習会「Wisteria実践」

倍精度： simdlen = 8
単精度： simdlen = 16
num=>各ループが上の倍数になるように割り当て55

OMP SIMD単独• C⾔語：• Fortran:#pragma omp parallel forfor (i=0; i < num; i++) {#pragma omp simdfor (j=0; j < num; j++) {sum[i] = sum[i] +a[i][j];}}2023/2/14!$omp parallel dodo i=1, num!$omp simd
do j=1, numsum(i) = sum(i) + a(j,i)
enddo!$omp end simd
enddo!$omp end parallel

講習会「Wisteria実践」

56

Declare SIMD：富⼠通コンパイラは苦⼿• 関数をまるごとSIMD化したい場合• その関数はOMP SIMDの中で呼ばれる場合にSIMD化される• インライン展開もあまり、、#pragma omp declare simd notinbranchfloat min (float a, float b){ return a < b ? a : b; }void minner (float *a, float *b, float *c) {#pragma omp parallel for simdfor (i=0; i<N; i++)c[i] = min(a[i], b[i]);}2023/2/14

講習会「Wisteria実践」

57

最適化の例: 「バリアフリー」#pragma omp parallel forfor (i=0; i<N; i++)c[i] = a[i] + b[i];/* 暗黙のバリア同期 *//* ⼀度single regionに戻る*/#pragma omp parallel forfor (i=0; i<N; i++)#pragma omp parallel{#pragma omp for nowaitfor (i=0; i<N; i++)c[i] = a[i] + b[i];#pragma omp for
forの終了時にバfor (i=0; i<N; i++) リア同期しない
f[i] = d[i] + e[i];f[i] = d[i] + e[i];} // parallel終わり

2023/2/14

講習会「Wisteria実践」

58

まとめ：OpenMPレベル• C⾔語：OpenMPをちゃんと使うなら、tradコンパイラより
clangコンパイラの⽅がよい• 今後の開発はclangベースで⾏われる

2023/2/14

講習会「Wisteria実践」

59

Wisteria-Odyssey
MPI+OpenMPハイブリッド

2023/2/14

講習会「Wisteria実践」

60

MPIに関する情報• 「MPI基礎：並列プログラミング⼊⾨」講習会資料• https://www.cc.u-tokyo.ac.jp/events/lectures/189/• 「MPI上級編」講習会資料• https://www.cc.u-tokyo.ac.jp/events/lectures/190/• MPI仕様書 (3.1)• https://www.mpi-forum.org/docs/mpi-3.1/mpi31-report.pdf• MPI仕様書 (4.0)• https://www.mpi-forum.org/docs/mpi-4.0/mpi40-report.pdf

2023/2/14

講習会「Wisteria実践」

61

Wisteria-OdysseyのTofu-D• Tofu-D 形状: 全20ラック• (X, Y, Z, a, b, c) = (10, 8, 8, 2, 3, 2)è 3次元、2次元、1次元に展開して指定可能
典型的な例（全系） 参考：
ジョブ運⽤ソフトウェア• 3次元: 20 x 24 x 16
エンドユーザ向けガイド• X*a, Y*b, Z*c

シェルフ内に24枚搭載
シェルフ当たりは(1, 1, 4, 2, 3, 2)
c軸

MPI使⽤⼿引書• 2次元: 60 x 128

X軸(2シェルフ x 5ラック= 10) トーラス

• X*a*b, Y*Z*c• 1次元: 7680

Z軸
(4x2シェルフ=8)
トーラス

2023/2/14

Y軸 (2シェルフ x 4ラック = 8) メッシュ
講習会「Wisteria実践」

1シェルフ
62

JOBスクリプト（Odyssey, フラットMPI）#!/bin/bash#PJM -L rscgrp=lecture-o#PJM -L node=12#PJM --mpi proc=576#PJM -L elapse=00:01:00#PJM -g gt00
module load odyssey
mpiexec ./hello

リソースグループ名：lecture-o
利用ノード数、
MPIプロセス数
実行時間制限：１分
利用グループ名：gt00

MPIジョブを48*12 = 576 プロセスで実行する。
2023/2/14

講習会「Wisteria実践」

63

JOBスクリプト（Odyssey, ハイブリッドMPI: 1プロセス/ノード）#!/bin/bash#PJM -L rscgrp=lecture-o#PJM -L node=12#PJM --mpi proc=12#PJM --omp thread=48#PJM -L elapse=00:01:00#PJM -g gt00

利用ノード数
利用コア数（MPIプロセス数）
１MPIプロセス当たり
48スレッド生成
実行時間制限：１分

module load odyssey
mpiexec ./hello_omp
MPIジョブを1＊12 = 12 プロセスで実行する。
2023/2/14

リソースグループ名：lecture-o

講習会「Wisteria実践」

利用グループ名：gt00
64

JOBスクリプト（Odyssey, ハイブリッドMPI: 4プロセス/ノード）#!/bin/bash#PJM -L rscgrp=lecture-o#PJM -L node=12#PJM --mpi proc=48#PJM --omp thread=12#PJM -L elapse=00:01:00#PJM -g gt00

利用ノード数
利用コア数（MPIプロセス数）
１MPIプロセス当たり
12スレッド生成
実行時間制限：１分

module load odyssey
mpiexec ./hello_omp
MPIジョブを4＊12 = 48 プロセスで実行する。
2023/2/14

リソースグループ名：lecture-o

講習会「Wisteria実践」

利用グループ名：gt00
65

MPI性能 (Odyssey, pt2pt, 1.2.37)• バンド幅: osu_bw• レイテンシ：osu_latency• 最⼤(隣接): 6.3 GB/s
• 8KBの落ち込み:• 最⼩(隣接): 1.06us-mca btl_tofu_eager_limit 8192で改善も
16K-32Kは悪化
7000

100

バンド幅 (MB/s)

レイテンシ (usec)

1000

10

1
1

16

256

4096

65536

1048576

メッセージ長 (byte)

6000

default

5000

8K

4000
3000
2000
1000
0
1

2023/2/14

7000

講習会「Wisteria実践」
default

8

64

512

4096

32768

メッセージ長 (byte)

26214466

2097152

ノード形状の指定とrank-map• 形状の指定 (-L node)• MPIオプション (--mpi)• node=N1• node=N1xN2• node=N1xN2xN3• proc=プロセス数• rank-map-bychip,rank-maphostfile=ホストファイル• 割り当て⽅の指定• ホストにN個ずつ割り当て(proc/node)• node=N1xN2xN3:torus• rank-map-bynode,rank-maphostfile=ホストファイル• TOFU単位 (12ノード: 2x3x2単位)で切り出す• node=N1xN2xN3:mesh• 隣接していればいい• node=N1xN2xN3:nocont• ホストに1個ずつ割り当て• node=2x3の例• (0,0) (0,1) (0,2) (1,0)(1,1) (1,2)
実際には各⾏に1要素• 離れていても良い• 性能は低下してもいいから
とにかく実⾏したい場合
2023/2/14

講習会「Wisteria実践」

67

性能プロファイラ (Odyssey)• 富⼠通コンパイラには、性能プロファイラ機能がある• 富⼠通コンパイラでコンパイル後、実⾏コマンドで指定し利⽤する• 以下の3種類があります• 基本プロファイラ• プログラム全体で、最も時間のかかっている関数を同定するなど、傾向の把握• 詳細プロファイラ• 最も時間のかかっている関数内の特定部分において、メモリアクセス効率、キャッシュヒット率、スレッド実⾏効率、MPI通信頻度解析等の詳細分析• CPU 性能解析レポート• 特定区間のCPU性能解析情報の可視化• 詳細は、• 「利⽤⼿引書」• J2UL-2483-02Z0(10) 2022年3⽉, Technical Computing Suite V4.0L20、Development Studio プロファイラ使⽤⼿引書
2023/2/14

講習会「Wisteria実践」

68

性能最適化情報の提⽰（1/2）• C, Fortranを問わず、コンパイラが⾏った最適化情報を知ることは性能最適化で重要である• 各ファイルのソースコードごとに、どのような最適化が⾏われたか出
⼒する、コンパイラオプションがある• C, Fortranで共通の翻訳情報• -Nlst=p ：標準の最適化情報を出⼒(デフォルト)• -Nlst=t ：詳細な最適化情報を出⼒

2023/2/14

講習会「Wisteria実践」

69

性能最適化情報の提⽰（2/2）• Fortran のみ、p、t以外も指定可能• -Nlst=a ：名前の属性情報を出⼒• -Nlst=d ：派⽣型の構成情報を出⼒• -Nlst=I ：インクルードされたファイルのプログラムリスト
およびインクルードファイル名⼀覧を出⼒を出⼒• -Nlst=m ：⾃動並列化の状況をOpenMP指⽰⽂によって表現した原始プログラムを出⼒• -Nlst=x ：名前および⽂番号の相互参照情報を出⼒• 詳細は、オンラインマニュアルを参照のこと• C⾔語使⽤⼿引書• C++⾔語使⽤⼿引書• Fortran使⽤⼿引書

2023/2/14

講習会「Wisteria実践」

70

Wisteria-Aquarius
GPU利⽤を中⼼に

2023/2/14

講習会「Wisteria実践」

71

Aquariusの構成• Intel Xeon Platinum 8360Y (36c 2.4GHz) x 2ソケット, 512GBメモリ
メモリバンド幅：• ノード当たり8基のNVIDIA A100 GPU
409.6 GB/sec
DDR4
DDR4
DDR4
DDR4
DDR4
DDR4
DDR4
DDR4

x16

IBHDR

IB HDR:
200 Gbps x 4 link
(ノード当たり)x16

PCIe Gen4
x16

A100

UPI

CPU1
Intel Xeon Platinum
8360Y
(ICX)

CPU2
Intel Xeon Platinum
8360Y
(ICX)

x16

x16

PCIe
sw

x16

x16

IBHDR

A100

x16

PCIe
sw

PCIe
sw

x16

x16

x16

A100

A100

A100

PCIe
sw

x16

x16

A100

NVSwitch
NVSwitch
NVSwitch
NVSwitch
NVSwitch
NVSwitch
2023/2/14

DDR4
DDR4
DDR4
DDR4
DDR4
DDR4
DDR4
DDR4

講習会「Wisteria実践」

IBHDR

x16

A100

PCI Express:
Gen4 x16 = 16 GHz x 16 (128b130b)= 31.5 GB/s (⽚⽅向)
x16

IBHDR

PCIe Gen4
x16

A100

NVLink3:
GPU当たり 50 GHz x4x2 x6= 300 GB/s (片方向)72

CPUの構成• 最新のIceLakeだが、トークン
消費がGPU基準なのでCPUのみの利⽤はお勧めできない(ノード当たり、Odysseyの24倍の消費)• ソケット内も2 NUMAドメインに分割 (18コアずつ)• shareでは、1GPU当たり9コア

2023/2/14

nvidia-smi topo -mpの出⼒(lecture-aの場合)
GPU0 mlx5_0 mlx5_1 mlx5_2

GPU0

mlx5_3 CPU Affinity NUMA Affinity

X

PXB

SYS

SYS

SYS

mlx5_0 PXB

X

SYS

SYS

SYS

mlx5_1 SYS

SYS

X

SYS

SYS

mlx5_2 SYS

SYS

SYS

X

SYS

mlx5_3 SYS

SYS

SYS

SYS

X

講習会「Wisteria実践」

0-8

0-3

73

NVIDIA A100 Tensor Core GPU (1/2)
NVIDIA A100 Tensor コア GPU アーキテクチャの詳細• 108 SM (Streaming Multiprocessor)

図4.

2023/2/14
講習会「Wisteria実践」
74
出典:
NVIDIA
128
個の SM を搭載した GA100
フル GPU (A100 Tensor
コア
GPU A100は Tensor core GPUアーキテクチャ

NVIDIA A100 Tensor Core GPU (2/2)

NVIDIA A100 Tensor コア GPU アーキテクチャの詳細• 倍精度にも対応したTensor Coreを搭載• 19.5 TF @ FP64, FP32
• 156 TF @ TF32 (実質19bit)• 312 TF @ FP16, BF16
• 624 TF @ INT8
• 1248 TF @ INT4

• メモリ HBM2 40GB 搭載• 1.555 TB/s

2023/2/14

講習会「Wisteria実践」

図 5.

GA100 ストリーミング マルチプロセッサ (SM)75
出典: NVIDIA A100 Tensor core GPUアーキテクチャ

Volta世代のおさらい• Independent thread schedulingが導⼊され，（デフォルトでは）warpを構成する32スレッド内の暗黙の同期がなくなった• コンパイル時に arch=compute_60,code=sm_70 を指定すれば，Pascal世代以前と同様に暗黙の同期が復活（勝⼿にPascalモードと呼ぶことにしている）• A100ではarch=compute_60,code=sm_80で暗黙同期が復活• 重⼒ツリーコードの場合，Pascalモードの⽅が最⼤2割程度⾼速だった• 暗黙同期の有効化は最適化⼿法の1つ• A100の場合，Pascalモードは最⼤1割程度⾼速（差が縮まったのは，以降で紹介する新機能による⾼速化が効いている）• 整数演算ユニットがCUDAコアから独⽴したため，独⽴な演算の同時実⾏による演算時間の隠ぺいによって⾼速化された
2023/2/14

講習会「Wisteria実践」

76

Independent thread scheduling

2023/2/14

講習会「Wisteria実践」

77

Independent thread scheduling

2023/2/14

講習会「Wisteria実践」

78

A100/CUDA 11から導⼊された新機能• （正式名称NVIDIA A100 Tensor Core GPUにも⼊っている）
テンソルコアは⾏列の積和算専⽤機であり万⼈向けではないため，以下では関連機能も含めてすべてスキップ• Sparsity, TF32（実態はFP19）がテンソルコア関連機能• テンソルコアなしでは，理論ピーク性能⽐はV100 SXM2⽐で1.24倍• Asynchronous copy/barrier• グローバルメモリからシェアードメモリに直接（レジスタを介さずに）データを置ける• L2 cache residency control• 配列要素の⼀部をL2キャッシュに置きっぱなしにしておける• Warp-wide reduction• Warp内でのリダクション演算⽤の⾼速な命令が提供された
2023/2/14

講習会「Wisteria実践」

79

Asynchronous copy (memcpy_async)• グローバルメモリからシェアードメモリに直接（レジスタを介さずに）データを置ける• 処理はバックグラウンドで⾛る（したがって，同期も必要）

Svedin et al. 2021
(arXiv:2106.04979)• MPI_Send/MPI_Isend と対応させるとイメー
ジしやすい• A100では，hardware accelerated• つまり，⾮同期で動かせない状況でも速くなる• Svedin et al. (2021):• Pipeline API は Barrier API よりも⾼速• メモリ律速な問題では性能向上• 演算律速な問題では性能低下
2023/2/14

講習会「Wisteria実践」

80実装例（複数通り実装⽅法がある）#include <cooperative_groups.h>#include <cooperative_groups/memcpy_async.h>
Using namespace cooperative_groups;// デバイス関数内での実装：
thread_block_tile<TSUB> tile = tiled_partition<TSUB>(this_thread_block());// TSUB は，この関数内で memcpy_async に関与するスレッド数（この書き方の場
合は32以下）
cooperative_groups::memcpy_async(tile, &shared_mem, &global_mem,
sizeof(uint) * num);// 何か裏で回したい処理があれば，ここに書く
cooperative_groups::wait(tile);// ここでコピーの完了待ちをする
2023/2/14

講習会「Wisteria実践」

81

NVIDIA Developer Blogでの実装例• https://developer.nvidia.com/
blog/cuda-11-features-revealed///Without async-copy__shared__ extern int smem[];// algorithm loop iteration
while ( ... ) {__syncthreads();// load element into shared memfor (int i = ... ) {smem[i] = gmem[i];}__syncthreads();/* compute on smem[] *///With async-copy
using namespace nvcuda::experimental;__shared__ extern int smem[];pipeline pipe;// algorithm loop iteration
while ( ... ) {__syncthreads();// load element into shared memfor (int i = ... ) {// initiate async memory copy
memcpy_async(smem[i], gmem[i], pipe);}// wait for async-copy to complete
pipe.commit_and_wait();__syncthreads();}/* compute on smem[] */2023/2/14

講習会「Wisteria実践」}

82

L2 cache residency control• Best Practice Guide中の記述："A portion of the L2 cache can be set aside for persistent accesses
to a data region in global memory"• L2キャッシュの容量は40 MB（V100では6 MBだった）• 1/16（=2.5 MB）刻みで調整できる（white paperより）• CUDAストリームごとに1つの配列中の連続領域を指定可能• 16個設定できても良さそうだが，残念ながらそうなってはいない

2023/2/14

講習会「Wisteria実践」

83

NVIDIA Developer Blogでの実装例• https://developer.nvidia.com/blog/cuda-11-featuresrevealed/
cudaStreamAttrValue attr ;attr.accessPolicyWindow.base_ptr = /* beginning of range in global memory */ ;attr.accessPolicyWindow.num_bytes = /* number of bytes in range */ ;// hitRatio causes the hardware to select the memory window to designate as per
sistent in the area set-aside in L2
attr.accessPolicyWindow.hitRatio = /* Hint for cache hit ratio: 1.0で良い */attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);
2023/2/14

講習会「Wisteria実践」

84

Warp-wide reduction• Throughputは16（FP32演算が64なので，4演算相当）• Warp shuffleが32（2演算相当）なので，5段（32スレッド）必要な旧実装に⽐べて（演算コストを無視しても）圧倒的に速い• __reduce_*_sync(unsigned mask, T value)• T: unsigned/intに対してはadd, min, max• ⼩細⼯すればfloatに対してmin/maxを返させることは可能• T: unsignedに対してはand, or, xor も可能• Supported by devices of compute capability 8.x or higher• Pascalモード（arch=compute_60,code=sm_80）の場合には，コンパイルが通らなかった（compute_80の指定が必要）• Pascalモード使⽤による⾼速化か，Ampereモード＋warp-wide reductionによる⾼速化か，の競争（コード次第なので，どちらが速いかはお試しください）
2023/2/14

講習会「Wisteria実践」

85

シェアードメモリ使⽤時の注意点• 静的に確保できるのはブロック当たり48 KBまで（この制約はずっと昔からあった）• V100ではSMあたり96 KBまでシェアードメモリに使えたが，（常識的には）SMあたりに複数ブロック⽴てるため，実質的に制約なしだった• A100では，SMあたり164 KBまで使える（1ブロックでは163 KBまで）• Dynamic allocationを使えば48 KB越えでも問題ない
extern __shared__ real dynamic_shared_memory[];__global__ void calc_acc(int num){
float4* pj = (float4*)dynamic_shared_memory;uint* queue = (uint *)&pj[num];.. Skip ..
}2023/2/14// how to launch
講習会「Wisteria実践」
calc_acc<<<blck, thrd, size_of_shared_memory,
CUDA_stream>>>(num);86GPU情報の表⽰• nvidia-smi• nvidia-smi -q+-----------------------------------------------------------------------------+

Timestamp: Wed Sep 14 20:09:53 2022

| NVIDIA-SMI 470.57.02

Driver Version: 470.57.02

CUDA Version: 11.4

Attached GPUs

: 1

Driver Version: 470.57.02

CUDA Version: 11.4||-------------------------------+----------------------+----------------------+
| GPU

Name

| Fan

Temp

Persistence-M| Bus-Id
Perf

Pwr:Usage/Cap||

Disp.A | Volatile Uncorr. ECC |Memory-Usage | GPU-Util||

Compute M. |MIG M. ||===============================+======================+======================|
|

0

| N/A

NVIDIA A100-SXM...
24C

P0

On

| 00000000:C7:00.0 Off |51W / 400W |0MiB / 40536MiB |

|

|

|

0 |
0%

Default |GPU 00000000:C7:00.0
Product Name……
Applications Clocks

Disabled |+-------------------------------+----------------------+----------------------+|

Graphics: 1095 MHz

Memory: 1215 MHz

Default Applications Clocks+-----------------------------------------------------------------------------+| Processes:: NVIDIA A100-SXM4-40GB

Graphics: 1095 MHz

Memory: 1215 MHz

Max Clocks

GPU Memory |Graphics: 1410 MHz

Usage

SM

: 1410 MHz|=============================================================================|

Memory: 1215 MHz

|

Video

: 1290 MHz

|
|

GPU

GI

CI

ID

ID

PID

Type

Process name

No running processes found||+-----------------------------------------------------------------------------+Max Customer Boost Clocks
Graphics

2023/2/14

講習会「Wisteria実践」: 1410 MHz

87コンパイラの種類と実⾏(Aquarius)• ログインノードとAquarius計算ノードとでは、CPUの命令セットが（ほぼ）同じ• ログインノード：命令セットアーキテクチャ Intel CascadeLake＋AVX512, x86_64• Aquarius計算ノード： 命令セットアーキテクチャ Intel IceLake＋AVX512, x86_64• 様々なコンパイラが利⽤可能: GPU向けには gcc+CUDAか NVIDIAを推奨$ module load aquarius cuda ompi-cuda(2023/2現在、gcc=8.3.1, CUDA=11.4, OMPI-CUDA=4.1.1-11.4)$ module load nvidia cuda ompi-cuda(2023/2現在、 NVIDIA=22.7, CUDA=11.4, OMPI-CUDA=4.1.4-11.4)
⾔語

GNUコンパイラ

Intelコンパイラ

NVIDIA コンパイラ(旧PGI)

C

gcc

icc

nvc (pgcc)

C++

g++

icpc

nvc++(pgc++)

Fortrangfortran

ifort

OpenACC

CUDAコンパイラ
nvcc

nvfortran (pgfortran)○

2023/2/14

講習会「Wisteria実践」

88

NVIDIAコンパイラ(OpenACC)

8.
9.
10.
11.

• コンパイラオプションとして-ta=tesla,cc80 を付ける

12.

ソースコード

13.
14.
15.
16.
17.

サブルーチン名

18.
19.
20.

コンパイラメッセージ(fortran)

21.

subroutine acc_kernels()
double precision :: A(N,N), B(N,N)
double precision :: alpha = 1.0
integer :: i, j
A(:,:) = 1.0
B(:,:) = 0.0!$acc kernels
do j = 1, N
do i = 1, N
B(i,j) = alpha * A(i,j)
end doend do!$acc end kernels
end subroutine acc_kernels

pgfortran -O3 -acc -Minfo=accel -ta=tesla,cc80 -Mpreprocess acc_compute.f90 -o
acc_computeacc_kernels:
配列aはcopyin, bはcopyoutとして扱われます
14, Generating implicit copyin(a(:,:))
Generating implicit copyout(b(:,:))
15, Loop is parallelizable
16, Loop is parallelizable
Accelerator kernel generated
Generating Tesla code
15, !$acc loop gang, vector(4) ! blockidx%y threadidx%y
16, !$acc loop gang, vector(32) ! blockidx%x threadidx%x….

15, 16行目の２重ループは(32x4)のスレッドでブロック分割して扱います。
講習会「Wisteria実践」
892023/2/14

JOBスクリプトサンプルの説明（Aquarius, MPIなし）#!/bin/bash#PJM -L rscgrp=lecture-a#PJM -L gpu=4#PJM -L elapse=00:01:00#PJM -g gt00
module load aquarius cuda./a.out

リソースグループ名：lecture-a
利用GPU数
実行時間制限：１分
利用グループ名：gt00

2023/2/14

講習会「Wisteria実践」

90

Wisteria-Aquarius
GPU+MPI環境

2023/2/14

講習会「Wisteria実践」

91

OpenMPI• Aquariusでは、MPIとGPUを使うのは OpenMPIがお勧め• CUDA Aware: GPUメモリを直接送受信可能(論理的に)• GPU Direct RDMA: さらにRDMA機能で⾼速に転送• デフォルトでこれらの機能がONになっている• export UCX_MAX_RNDV_RAILS=1 （ノード内通信では増やすべき）• export UCX_RNDV_THRESH=32K• 特定メッセージサイズで性能が落ち込むときに調整• 不具合が起こる時：以下をお試しください• export UCX_MEMTYPE_CACHE=n• export UCX_IB_GPU_DIRECT_RDMA=n2023/2/14

講習会「Wisteria実践」

92

JOBスクリプトサンプルの説明（Aquarius, MPI）
リソースグループ名：lecture-a#!/bin/bash#PJM -L rscgrp=lecture-a#PJM -L gpu=4#PJM --mpi proc=4#PJM -L elapse=00:01:00#PJM -g gt00

利用GPU数、
MPIプロセス数
実行時間制限：１分
利用グループ名：gt00

module load aquarius cuda ompi-cuda
mpiexec -machinefile $PJM_O_NODEINF -n $PJM_MPI_PROC ¥-npernode 4 ./wrapper.sh ./a.out
MPIジョブをノードあたり4プロセスで実行する。
2023/2/14

講習会「Wisteria実践」

93

Aquariusノード内で複数ランク使う場合• プログラム側で複数GPUを考慮していない場合には、以下のようなシェルスクリプトを⽤意• wrapper.sh (chmod +x wrapper.sh を忘れずに)#!/bin/shexport LOCAL_RANK=$OMPI_COMM_WORLD_LOCAL_RANK
export CUDA_VISIBLE_DEVICES=$LOCAL_RANK$*

2023/2/14

講習会「Wisteria実践」

94

Aquariusノード内で複数ランク使う場合(2)• どのネットワークデバイスを使うか考慮• UCX_NET_DEVICES 環境変数
• InfiniBandのポート番号を使い分ける• 例えば1レールずつ分けるとwrapper.shはこうなる#!/bin/shcase $OMPI_COMM_WORLD_LOCAL_RANK in
0|1) export UCX_NET_DEVICES=mlx5_0:1 ;;2|3) export UCX_NET_DEVICES=mlx5_1:1 ;;4|5) export UCX_NET_DEVICES=mlx5_2:1 ;;6|7) export UCX_NET_DEVICES=mlx5_3:1 ;;esac$*
2023/2/14

講習会「Wisteria実践」

95

Aquariusノード内で複数ランク使う場合(3)• まとめると以下の通り#!/bin/shcase $OMPI_COMM_WORLD_LOCAL_RANK in
0|1) export UCX_NET_DEVICES=mlx5_0:1 ;;2|3) export UCX_NET_DEVICES=mlx5_1:1 ;;4|5) export UCX_NET_DEVICES=mlx5_2:1 ;;6|7) export UCX_NET_DEVICES=mlx5_3:1 ;;esac
export LOCAL_RANK=$OMPI_COMM_WORLD_LOCAL_RANK
export CUDA_VISIBLE_DEVICES=$LOCAL_RANK
numactl -l $*2023/2/14

講習会「Wisteria実践」

ç必要に応じて96

ノード間通信 (CUDA: 12.0, OpenMPI 4.1.5) NEW!!
• バンド幅: osu_bw• レイテンシ：osu_latency• HtoH: 1.57 us (CPUと同じ)• 最⼤: ~25 GB/s
• DtoH, HtoDはほぼ同じ: 2.28 us• 8KB~16KBの *toDが落ち込み• DtoD: 2.98 us
export UCX_MAX_RNDV_RAILS=1

1000

30000

HtoD

100

Normal

25000

HtoH

バンド幅 (MB/sec)

レイテンシ (usec)

CPU

HtoH

20000

DtoH

HtoD
DtoH

15000

DtoD

DtoD

10000

10

5000
0

1
1

32

1024

32768

2023/2/14(byte)
メッセージ⻑

1048576

1

講習会「Wisteria実践」

32

1024
32768
メッセージ⻑(byte)

104857697

ノード間通信 (CUDA: 12.0, OpenMPI 4.1.5,パラメータ調整後)• バンド幅: osu_bw• レイテンシ：osu_latency• 最⼤: ~25 GB/s
• HtoH: 1.57 us (CPUと同じ)• *toDの8KB~16KBの落ち込みは改善、• DtoH, HtoDはほぼ同じ: 2.28 us*toHが悪化
export UCX_MAX_RNDV_RAILS=1• DtoD: 2.98 us
export UCX_RNDV_THRESH=4K

1000

30000
HtoH

100

バンド幅 (MB/sec)

レイテンシ (usec)

CPU
HtoD
DtoH
DtoD
10

1
1

32

1024

32768

メッセージ⻑ (byte)
2023/2/14

1048576

25000

Normal

20000

HtoD

HtoH
DtoH

15000

DtoD

10000
5000
0
1

講習会「Wisteria実践」

32

1024

32768

メッセージ⻑(byte)

104857698

ノード間通信 (CUDA: 12.0, OpenMPI 4.1.5)• バンド幅: osu_bw

export UCX_MAX_RNDV_RAILS=1,2,3,4
export UCX_MAX_EAGER_RAILS=1,2,3,4
export UCX_RNDV_THRESH=16K• 最⼤: 25, 49, 74, 98 GB/s(CUDA なし)
CUDAの場合は1railと変化なし
100000

1rail

2rail

3rail

4rail

8

64

512
4096
メッセージ⻑(byte)

バンド幅 (MB/sec)

80000
60000
40000
20000
0
1
2023/2/14

講習会「Wisteria実践」

32768

262144209715299

ノード内通信 (CUDA: 12.0, OpenMPI 4.1.5)• バンド幅: osu_bw• バンド幅: osu_bw• デフォルトでInfiniBandを使う• メモリ経由より速い• マルチレールが効く(2 railまで)• HtoD• DtoD: ~280 GB/sec, NVLink3 (300GB/s)の性
能• DtoH, HtoDでは ~10, ~15GB/s 、HtoHでは~36GB/sexport UCX_MAX_RNDV_RAILS=1• 変な性能低下が⾒られる
30000

300000

CPU
HtoH
HtoD
DtoH
DtoD

200000150000

1rail

25000

2rail

バンド幅 (MB/sec)

バンド幅 (MB/sec)

250000

20000
15000
10000

100000
50000

5000

0

0
1

128
16384
2097152 268435456
メッセージ⻑(byte)
2023/2/14

1

講習会「Wisteria実践」

16

256

4096

メッセージ⻑(byte)

65536

1048576100

ノード内通信 (CUDA: 12.0, OpenMPI 4.1.5)• バンド幅: osu_bw• デフォルトでInfiniBandを使う•

メモリ経由より速い• DtoD: ~280 GB/sec, NVLink3 (300GB/s)の性能• DtoH, HtoDでは ~10, ~15GB/s 、HtoHでは~36GB/s• 変な性能低下が⾒られる

export UCX_MAX_RNDV_RAILS=2

300000CPU

バンド幅 (MB/sec)

250000

HtoH

200000

HtoD
DtoH

150000

DtoD

100000
50000
0
1

16
2023/2/14

256

4096
65536
メッセージ⻑(byte)

講習会「Wisteria実践」

1048576

16777216 268435456
101ノード内通信 (CUDA: 12.0, パラメータ調整)• Dto*パターンのとき短メッセージで性能が出ないè以下で改善• export UCX_RNDV_THRESH=1• 他のパターンでは副作⽤あり

export UCX_MAX_RNDV_RAILS=2
35000

250000CPU
200000

バンド幅 (MB/sec)

バンド幅 (MB/sec)

HtoH
HtoD

150000

DtoH

30000

DtoH
DtoD

100000
50000

DtoD

25000

DtoH (RNDV=1)

20000

DtoD (RNDV=1)

15000
10000
5000
0

0
1

32

1024

32768

1048576

16

256

4096

65536

メッセージ⻑(byte)

メッセージ⻑(byte)
2023/2/14

1

講習会「Wisteria実践」

102

Wisteria/BDEC-01のストレージ• FEFS (Fujitsu Exabyte File
System)• 並列ファイルシステム(4 rack)• 25 PB, 500 GB/s
• 4 MDS• OSS: DDN SFA7990XE x16• /work• ⾼速ファイルシステム(2 rack)• 1 PB, 1 TB/s
• 2 MDS• OSS: DDN SFA400NVXE x16• NVMe SSD• /data/scratch: スクラッチ• /data/perm: 恒久
2023/2/14

講習会「Wisteria実践」

103

ストレージ性能• メタデータ性能 (mdtest)• 共有FSに対し，複数のクライアントから，個別のディレクトリに対するファイル操作• ⾼速FSも並列FSと同等の性能• O: 4シェルフ 100*4=400 Gbps有効?
• A: 8ノード 200*4*8=6.4 Tbps
Odyssey

Aquarius

4x4x8ノード
1024プロセス

8ノード
128プロセスFilecreation

63,59653,856File stat120,498

281,756

Fileremoval26,72092,4722023/2/14• ファイル書き込み性能 (IOR)• 共有FSに対し独⽴なファイルの書
き込み性能 (MiB)• O: 極⼒GIOノードを割り当て• A: 全ノード
Odyssey
⾼速FS

20x12x16
ノード確保
960
プロセス

並列FS

16x24x16
ノード確保
3072
プロセス

講習会「Wisteria実践」

Aquarius

455,372

45ノード
720
プロセス

322,115

281,480

45ノード
720
プロセス

236,643

104

Singularityの利⽤

2023/2/14

講習会「Wisteria実践」

105

Singularity: コンテナ仮想化
Singularity Hub

shub://
Docker Hub

docker://

コンテナイメージ
ファイル

Production向け (Read only)

コンテナ環境内で実⾏ (shell)$ singularity build
container.img <source>
container.img

Singularity recipe

環境構築 (writable image)

コンテナ環境のコマンドで実⾏(exec)$ sudo singularity build--writable container.img<source>
コンテナイメージ
ファイル$ singularity shell
container.img> python mnist.py$ singularity exec
container.img python
mnist.py

環境構築 (sandbox)

定義された通りに実⾏(run)$ sudo singularity build--sandbox container.img<source>$ singularity run
container.img
or$ ./container.img

2023/2/14

講習会「Wisteria実践」

106

Singularity向けジョブスクリプト(1ノード)#!/bin/bash#PJM -L rg=tutorial-a#PJM -L gpu=1#PJM -g gt00
#PJM -L jobenv=singularity#PJM -L elapse=00:15:00#PJM -jSingularity使⽤時に必要！！

cd $PJM_O_WORKDIR
module load gcc cuda singularity
コンテナの外にアクセスす
SIF=/work/gt00/share/tf-image.file
るパス毎に必要
singularity exec $SIF which python
singularity exec $SIF python --version
singularity exec --nv --bind `pwd` --bind /work/share/MNIST $SIF python keras-tf-mnist.py ¥>& keras-tf-mnist.log.$PJM_JOBID
2023/2/14

講習会「Wisteria実践」

107Wisteria/BDEC-01の
便利な使い⽅：JupyterHub, Guacamoleの紹介

2023/2/14

講習会「Wisteria実践」

108

JupyterHub• JupyterNotebookなどの環境をポータルとして提供• 実はターミナル機能も使える• https://wisteria08.cc.utokyo.ac.jp:8000/jupyterhub/• Windows: InternetExplorerでは不具合が起こることがある• MacOS: Safariではキー⼊⼒を受け付けないことがある、Firefox or Chrome
推奨2023/2/14

講習会「Wisteria実践」

109

JupyterNotebookからのジョブ投⼊• 準備:• pysubコマンドでジョブ投⼊可能~/.notebook/template_job.sh を作
成[4] %cd /work/gt00/tABCDE/job_dir#!/bin/sh[5] import os#PJM -L rscgrp=tutorial-a[6] print(os.getcwd())#PJM -L gpu=1[7] f = open(’./tmp.txt’,’a’)[8] （何か処理）[9] f.close()[10] !pysub -f Notebook.ipynb 5,6,7,8,9

2023/2/14

講習会「Wisteria実践」

110

Pythonの利⽤• 「Python環境の利⽤」スパコンニュース，2021年11⽉• https://www.cc.utokyo.ac.jp/public/VOL23/No6/12_202111Python.pdf• minicondaの利⽤⽅法• JupyterHubのカスタマイズ• Singularityの利⽤

2023/2/14

講習会「Wisteria実践」

111

Apache Guacamole
ブラウザ上でリモートデスク
トップ実現• VNC (Virtual Network
Computing)のクライアント• X-Windowクライアントを⼿元に⾶ばすよりずっと⾼速• https://wisteria09.cc.utokyo.ac.jp:8443/guacamole/2023/2/14

講習会「Wisteria実践」

112

Apache Guacamoleの準備• ログインノードでvncserver起動(vnc⽤独⾃パスワードの設定が求められるかも)$ vncserver• 起動したポート番号確認: 5900+番号=> 5903
$ vncserver -list
X DISPLAY #PROCESS ID:3891292• 接続情報登録: set_vncserver -h 起動ホスト名 -p 起動ポート番号$ set_vncserver -h wisteria03 -p 5903
Enter LDAP Password: ポータルのパスワード
modifying entry "cn=z30105-vnc,…”• DISPLAY環境変数の設定$ export DISPLAY=:3
2023/2/14

講習会「Wisteria実践」

113

お疲れ様でした！
アンケートにご協⼒お願いします• https://forms.office.c
om/r/LW6RCqt7Bp

2023/2/14

講習会「Wisteria実践」

114