第220回 お試しアカウント付き
並列プログラミング講習会「Wisteria実践」
東京大学情報基盤センター
塙 敏博
三木 洋平
2024/1/19

講習会「Wisteria実践」

1

講習会概略• 開催日：
2024年1月19日（金） 13：00 - 17：15• 形態：ZoomとSlackを用いたオンライン講習会• 講習会プログラム： 講師：塙• 13：00 – 13：45 Wisteria/BDEC-01システム紹介• 14：00 – 15：15 Odysseyノード：A64FXにおけるOpenMP並列化(ごく簡単に)、
MPI+OpenMPハイブリッド並列、性能分析（講義+演習）• 15：30 - 17：00 Aquariusノード: A100 GPUの利用、OpenACC、MPI並列化、性能分析（講義+演習）• 17：00 –Q&A

2024/1/19

講習会「Wisteria実践」

2

是非こちらも参考に: 講習会資料＋ビデオ• Odyssey• 第218回 「A64FXプログラミング入門」• 第222回 「OpenMPによるA64FX並列プログラミング入門」• Aquarius• 第208回 「OpenMPで並列化されたC++プログラムのGPU移植手法」• 第209回 「OpenACCとMPIによるマルチGPUプログラミング入門」• 第211回 「MPI+OpenMPで並列化されたFortranプログラムのGPUへの移行手法」• Odyssey+Aquarius連携• 第191回 「異種システム間連成アプリケーション開発を学ぶ:WaitIO/MP講習会」• 第217回 「異種システム間連成アプリケーション開発を学ぶ:WaitIO/MP講習会 - シミュレーションと
機械学習融合編 - 」• MPI全般• 第213回 「MPI基礎：並列プログラミング入門」• 第214回 「MPI上級編」
2024/1/19

講習会「Wisteria実践」

3

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

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

講習会「Wisteria実践」

4

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

2024/1/19

講習会「Wisteria実践」

5

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

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

2024/1/19

講習会「Wisteria実践」

6

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

2024/1/19

講習会「Wisteria実践」

7

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

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

2024/1/19

講習会「Wisteria実践」

8

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

2024/1/19

講習会「Wisteria実践」

9

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

2024/1/19

講習会「Wisteria実践」

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

2024/1/19

講習会「Wisteria実践」

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

2024/1/19

講習会「Wisteria実践」

12

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

2024/1/19

講習会「Wisteria実践」

13

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

2024/1/19

講習会「Wisteria実践」

14

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

講習会「Wisteria実践」

15

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

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

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

講習会「Wisteria実践」

17

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

講習会「Wisteria実践」

18

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

2024/1/19

講習会「Wisteria実践」

19

20

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

ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

interactive-o(interactive-o_n1)(interactive-o_n12)

1
2 ～ 12

30 分
10 分

28
28

ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

interactive-a

1ノード

10 分

448

share-interactive

1GPU

10 分

56

• Wisteria-A（Aquarius）:データ・学習ノード群
キュー名

2024/1/19

講習会「Wisteria実践」

20

参考：インタラクティブ実行のやり方（本講習会では使えません）
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※インタラクティブ用のノードがすべて使われている場合、資源が空くまでログインできません。
※このアカウントでは使えません。
2024/1/19

講習会「Wisteria実践」

21

22

参考：ジョブクラス：バッチジョブ（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倍）

2024/1/19

講習会「Wisteria実践」

22

23

参考：ジョブクラス：バッチジョブ（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
〃
〃

キュー名

2024/1/19

講習会「Wisteria実践」

23

24

参考：ジョブクラス：プリポストサービス
ログインノードと同じアーキテクチャのノードを使用
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

キュー名

2024/1/19

講習会「Wisteria実践」

24

25

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

キュー名

ノード数

制限時間
（Elapsed）

メモリ容量
（GB）

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

1 ～ 12

15 分

28

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

1 ～ 1215 分

28

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

2024/1/19

講習会「Wisteria実践」

25

26

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

GPU数

制限時間
（Elapsed）

メモリ容量
（GB）

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

1,2,4

15 分

56/GPUtutorial-a（この講習会の時間中のみ使用可能）

1,2,4

15 分

56/GPU

share-aと同様の使い方

2024/1/19

講習会「Wisteria実践」

26

「Wisteria実践」について

2024/1/19

講習会「Wisteria実践」

27

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

• Odysseyの多数の計算ノードによる効率の良い実行方法• A64FX単体のチューニング、OpenMPについては別の講習会にて• AquariusノードのGPUの扱い方等について学ぼう！

2024/1/19

講習会「Wisteria実践」

28

（現代の）スパコンは並列性が命• 京コンピュータ(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コア当たりではスパコンの方が遅いこともある• スパコン向けは，電力効率を狙ってクロック周波数を下げている
è その分コア数が増えているがフルに使えないと性能が出ない
2024/1/19

講習会「Wisteria実践」

29

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

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

マルチプロセスとマルチスレッドは共存可能→ハイブリッドMPI/OpenMP実行
2024/1/19

講習会「Wisteria実践」

30

Weak ScalingとStrong Scaling MPI通信もポイント
並列処理においてシステム規模を大きくする方法

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

Weak Scaling

Strong Scalingも重要

それまで解けなかった

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

規模の問題が解ける
2024/1/19

講習会「Wisteria実践」

31

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

https://www.cc.u-tokyo.ac.jp/events/lectures/213/• 「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/216/• 「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/「こんな講習会が欲しい！」という要望があればぜひアンケートへ！
2024/1/19

講習会「Wisteria実践」

32

Wisteria-Odyssey
A64FXの基礎
以前のWisteria実践で使ったものも残しておきます。
A64FX講習会の資料も参照してください。
第218回 「A64FXプログラミング入門」
第222回（予定) 「OpenMPによるA64FX並列プログラミング入門」

2024/1/19

講習会「Wisteria実践」

33

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

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

0

1

2

3

4

5

6

7

8

9

10 11

HBM2

HBM2

命令セット

CMG#2
HBM2

HBM2

CMG#0

CMG#1

プロセッサ名PCIe コントローラ

プロセッサ数 (コア数)周波数
理論演算性能
メモリ容量
メモリ帯域幅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プ
ロセスのハイブリッド推奨
2024/1/19

講習会「Wisteria実践」

34

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

2024/1/19

リ
ン
グ
バ
ス

講習会「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

35

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
2024/1/19

講習会「Wisteria実践」

36

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

2024/1/19

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

37

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
2024/1/19

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

38

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

2024/1/19

5

10

15

講習会「Wisteria実践」

20

25

コア数

30

35

40

45
39

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

2024/1/19

5

10

講習会「Wisteria実践」

15

20

25

コア数

30

35

40

45
40

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配置を実現するにはコア番号を明示する必要
2024/1/19

講習会「Wisteria実践」

41

コンパイラの種類と実行(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)
2024/1/19

講習会「Wisteria実践」

a.out

aarch64

言語

ネイティブ
コンパイラ

クロス
コンパイラ

C

fcc

fccpx

C++

FCC

FCCpx

Fortranfrt

frtpx42

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

2024/1/19

fjomplib

講習会「Wisteria実践」

C / C++

Fortrantradモードclangモード○ (default)○ (default)○ (default)

○

×

○

43

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

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

講習会「Wisteria実践」• C言語使用手引書• C++言語使用手引書• Fortran使用手引書
44

基本的なコンパイルオプション (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• 富士通の拡張命令を使う場合
2024/1/19

講習会「Wisteria実践」

45

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
2024/1/19

講習会「Wisteria実践」

46

メモリのアラインメント• 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の倍数の境界
2024/1/19

講習会「Wisteria実践」

47

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文に適用される
2024/1/19

講習会「Wisteria実践」

48

ポインタ指定の制限緩和 (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を指定する以外方法がない
2024/1/19

講習会「Wisteria実践」

49

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

2024/1/19

講習会「Wisteria実践」

50

まとめ：逐次実行レベル
1. CMGをまたぐ処理は性能が低下する場合が多い。MPIの併用を検討しましょう。
2. ラージページ＋デマンドページングのオプションをまずは入れましょう。
export XOS_MMM_L_PAGING_POLICY=demand:demand:demand• Pythonではエラーになるケースもあるようで、万能ではありません。

3. First Touchは必須。これも2.の設定がないと機能しません。
4. メモリのアラインメントをなるべく64 byteに揃える。

2024/1/19

講習会「Wisteria実践」

51

最適化の弊害(?)• あるパターンを見つけるとライブラリ関数への置き換えを勝手にやってしまう•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]2024/1/19

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

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

52

コンパイルオプション• -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

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

講習会「Wisteria実践」

53

OpenMPによるSIMD化
第222回（予定) 「OpenMPによるA64FX並列プログラミング入門」が実施されるのでそちらもご検討ください

2024/1/19

講習会「Wisteria実践」

54

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からサポート• コンパイラが自動ベクトル化もしてくれるが，指示をした方が確実• コンパイラが誤った答えを出すこともあるので要確認（経験済み）• コードの書き換えを必要とする場合も（しばしば）ある• メモリのアラインメントが重要（前述）2024/1/19

講習会「Wisteria実践」

55

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
2024/1/19

講習会「Wisteria実践」

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

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];}}2024/1/19!$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実践」

57

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]);}2024/1/19

講習会「Wisteria実践」

58

最適化の例: 「バリアフリー」#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 forfor (i=0; i<N; i++) forの終了時にバリア同期しない
f[i] = d[i] + e[i];f[i] = d[i] + e[i];} // parallel終わり

2024/1/19

講習会「Wisteria実践」

59

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

2024/1/19

講習会「Wisteria実践」

60

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

2024/1/19

講習会「Wisteria実践」

61

MPIに関する情報• 「MPI基礎：並列プログラミング入門」講習会資料• https://www.cc.u-tokyo.ac.jp/events/lectures/213/• 「MPI上級編」講習会資料• https://www.cc.u-tokyo.ac.jp/events/lectures/214/• 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

2024/1/19

講習会「Wisteria実践」

62

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)
トーラス

2024/1/19

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

1シェルフ
63

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 プロセスで実行する。
2024/1/19

講習会「Wisteria実践」

64

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 プロセスで実行する。
2024/1/19

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

講習会「Wisteria実践」

利用グループ名：gt00
65

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 プロセスで実行する。
2024/1/19

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

講習会「Wisteria実践」

利用グループ名：gt00
66

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

1000

100

バンド幅(MB/s)

レイテンシ (usec)

7000

10

6000

default

5000

8K

4000
3000
2000
1000

1
1

16

256

4096

65536

メッセージ長 (byte)
2024/1/19

104857601
講習会「Wisteria実践」

16

256

4096

65536

104857667

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

講習会「Wisteria実践」

68

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

講習会「Wisteria実践」

69

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

2024/1/19

講習会「Wisteria実践」

70

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

2024/1/19

講習会「Wisteria実践」

71

Wisteria-Aquarius
GPU利用を中心に

2024/1/19

講習会「Wisteria実践」

72

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
2024/1/19

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

2024/1/19

nvidia-smi topo -mpの出力(lecture-aの場合)
GPU0 mlx5_0 mlx5_1 mlx5_2 mlx5_3 CPU Affinity NUMA Affinity
GPU0

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

74

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

図4.

2024/1/19
講習会「Wisteria実践」
75
出典:
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

2024/1/19

講習会「Wisteria実践」

図 5.

GA100 ストリーミング マルチプロセッサ (SM)76
出典: 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コアから独立したため，独立な演算の同時実行による演算時間の隠ぺいによって高速化された
2024/1/19

講習会「Wisteria実践」

77

Independent thread scheduling

2024/1/19

講習会「Wisteria実践」

78

Independent thread scheduling

2024/1/19

講習会「Wisteria実践」

79

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内でのリダクション演算用の高速な命令が提供された
2024/1/19

講習会「Wisteria実践」

80

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 よりも高速• メモリ律速な問題では性能向上• 演算律速な問題では性能低下
2024/1/19

講習会「Wisteria実践」

81

実装例（複数通り実装方法がある）#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);// ここでコピーの完了待ちをする
2024/1/19

講習会「Wisteria実践」

82

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[] */2024/1/19

講習会「Wisteria実践」}

83

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

2024/1/19

講習会「Wisteria実践」

84

NVIDIA Developer Blogでの実装例• https://developer.nvidia.com/blog/cuda-11-features-revealed/
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);
2024/1/19

講習会「Wisteria実践」

85

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による高速化
か，の競争（コード次第なので，どちらが速いかはお試しください）2024/1/19

講習会「Wisteria実践」

86

シェアードメモリ使用時の注意点• 静的に確保できるのはブロック当たり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 ..
}2024/1/19// how to launch
講習会「Wisteria実践」
calc_acc<<<blck, thrd, size_of_shared_memory,
CUDA_stream>>>(num);87GPU情報の表示• nvidia-smi -q• nvidia-smi

Timestamp: Thu Jan 18 17:26:25 2024

Driver Version: 525.60.13

CUDA Version: 12.0

Attached GPUs
GPU 00000000:27:00.0: 8+-----------------------------------------------------------------------------+| NVIDIA-SMI 525.60.13

Driver Version: 525.60.13

CUDA Version: 12.0||-------------------------------+----------------------+----------------------+
| 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. |Product Name
Applications Clocks|===============================+======================+======================|
|

0

| N/A

NVIDIA A100-SXM...
24C

P0

On

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

|

|

|

Graphics
Memory0 |0%Default |+-------------------------------+----------------------+----------------------+| Processes:
|
|

GPU

GI
ID

|
CI

PID

ID

Type

Process name

No running processes found

Graphics: 1095 MHz

Memory: 1215 MHz...
Max Clocks

GPU Memory |Graphics: 1410 MHz

Usage

SM

: 1410 MHz

Memory: 1215 MHz||=============================================================================||

Video
Max Customer Boost Clocks|+-----------------------------------------------------------------------------+

Graphics

2024/1/19: 1095 MHz: 1215 MHz

Default Applications Clocks

Disabled |+-----------------------------------------------------------------------------+: NVIDIA A100-SXM4-40GB

……

講習会「Wisteria実践」: 1290 MHz: 1410 MHz

88

コンパイラの種類と実行(Aquarius)• ログインノードとAquarius計算ノードとでは、CPUの命令セットが（ほぼ）同じ• ログインノード：命令セットアーキテクチャ Intel CascadeLake＋AVX512, x86_64• Aquarius計算ノード： 命令セットアーキテクチャ Intel IceLake＋AVX512, x86_64• 様々なコンパイラが利用可能: GPU向けには gcc+CUDAか NVIDIAを推奨$ module load aquarius cuda/12.2 ompi-cuda(2024/1現在、gcc=8.3.1, CUDA=12.2, OMPI-CUDA=4.1.5-12.2)$ module load nvidia/22.7 cuda/11.4 ompi-cuda(2024/1現在、 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)○

2024/1/19

講習会「Wisteria実践」

89

NVIDIAコンパイラ(OpenACC)

8.
9.
10.
11.

• コンパイラオプションとして-gpu=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

nvfortran -O3 -acc -Minfo=accel -gpu=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実践」
902024/1/19

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

2024/1/19

講習会「Wisteria実践」

91

Wisteria-Aquarius
GPU+MPI環境

2024/1/19

講習会「Wisteria実践」

92

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=n2024/1/19

講習会「Wisteria実践」

93

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/12.2 ompi-cuda/4.1.5-12.2
mpiexec -machinefile $PJM_O_NODEINF -n $PJM_MPI_PROC ¥-npernode 4 ./wrapper.sh ./a.out
MPIジョブをノードあたり4プロセスで実行する。
2024/1/19

講習会「Wisteria実践」

94

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

2024/1/19

講習会「Wisteria実践」

95

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$*
2024/1/19

講習会「Wisteria実践」

96

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 $*2024/1/19

講習会「Wisteria実践」

ç必要に応じて97

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

30000

HtoH

100

CPU

25000

HtoH

バンド幅 (MB/sec)

レイテンシ (usec)

CPU
HtoD

20000

DtoH

HtoD
DtoH

15000

DtoD
10

DtoD

10000
5000

1
1

2024/1/19

32

1024

32768

メッセージ長 (byte)

104857601

講習会「Wisteria実践」

32

1024

32768

メッセージ長(byte)

104857698

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

1000

30000

HtoH
HtoD

100

HtoH

20000

DtoH
DtoD

10

CPU

25000
バンド幅 (MB/sec)

レイテンシ (usec)

CPU

HtoD
DtoH

15000

DtoD

10000
5000
0

1
1
2024/1/19

32

1024

32768

メッセージ長 (byte)

1048576

1
講習会「Wisteria実践」

32

1024

32768

メッセージ長(byte)

104857699

ノード間通信 (CUDA: 12.2, 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と変化なし

10000090000
バンド幅 (MB/sec)

80000

1rail

2rail

8

64

3rail

4rail

70000
60000
50000
40000
30000
20000
10000
0
1

512

4096

32768

2621442097152

メッセージ長(byte)
2024/1/19

講習会「Wisteria実践」

100

ノード内通信 (CUDA: 12.2, 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/s• 変な性能低下が見られる
export UCX_MAX_RNDV_RAILS=1
300000
バンド幅 (MB/sec)

250000200000150000

25000

HtoH
HtoD

20000

DtoH

15000

DtoD

50000

5000

0

0
64

4096

262144 16777216

1

16

256

4096

65536

1048576

メッセージ長(byte)

メッセージ長(byte)
2024/1/19

2rail

10000

1000001

1rail

バンド幅 (MB/sec)

CPU

30000

講習会「Wisteria実践」

101

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

export UCX_MAX_RNDV_RAILS=2• メモリ経由より速い• DtoD: ~280 GB/sec, NVLink3 (300GB/s)の性能• HtoD, DtoHでは ~25GB/s 、HtoHでは~52GB/s(概ねCPUと同じ)• 変な性能低下が見られる

300000CPU

バンド幅 (MB/sec)

250000

HtoH

200000

HtoD

150000

DtoH

100000

DtoD

50000
0
1

8

64

512

4096

32768

2621442097152 16777216134217728

メッセージ長(byte)
2024/1/19

講習会「Wisteria実践」

102

ノード内通信 (CUDA: 12.2, OpenMPI 4.1.5 パラメータ調整)• Dto*パターンのとき短メッセージで性能が出ないè以下で改善(まだ変)300000• export UCX_RNDV_THRESH=1• 他のパターンでは副作用あり

export UCX_MAX_RNDV_RAILS=2
35000

CPU

250000200000

バンド幅 (MB/sec)

バンド幅 (MB/sec)

HtoH
HtoD
DtoH

150000

DtoD

100000

30000

DtoH

25000

DtoD
DtoH (RNDV=1)

20000

DtoD (RNDV=1)

15000
10000

50000

5000

0
1
2024/1/19

64

4096

262144

メッセージ長(byte)

16777216
講習会「Wisteria実践」

0
1

16

256

4096

メッセージ長(byte)

65536

103

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: 恒久
2024/1/19

講習会「Wisteria実践」

104

ストレージ性能• メタデータ性能 (mdtest)• 共有FSに対し，複数のクライアントから，
個別のディレクトリに対するファイル操作• ファイル書き込み性能 (IOR)• 共有FSに対し独立なファイルの書き込み性能 (MiB)• O: 極力GIOノードを割り当て• A: 全ノード• 高速FSも並列FSと同等の性能• O: 4シェルフ 100*4=400 Gbps有効?
• A: 8ノード 200*4*8=6.4 Tbps
Odyssey

4x4x8ノード
1024プロセス

Aquarius

Filecreation

63,59653,856File stat120,498

281,756

Fileremoval26,72092,4722024/1/19

Odyssey8ノード
128プロセス

高速FS

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

並列FS

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

講習会「Wisteria実践」

Aquarius

455,372

45ノード
720
プロセス

322,115

281,480

45ノード
720
プロセス

236,643

105

Singularityの利用

2024/1/19

講習会「Wisteria実践」

106

Singularity: コンテナ仮想化
Singularity Hub

shub://
Docker Hub

docker://Singularity recipe

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

Production向け (Read only)$ singularity build
container.img <source>
container.img

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

コンテナ環境内で実行 (shell)$ singularity exec
container.img python
mnist.py

環境構築 (sandbox)

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

講習会「Wisteria実践」

107

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-tfmnist.py ¥>& keras-tf-mnist.log.$PJM_JOBID
2024/1/19

講習会「Wisteria実践」

108

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

2024/1/19

講習会「Wisteria実践」

109

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

2024/1/19

講習会「Wisteria実践」

110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,92024/1/19

講習会「Wisteria実践」

111

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

2024/1/19

講習会「Wisteria実践」

112

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

2024/1/19

講習会「Wisteria実践」

113

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

2024/1/19

講習会「Wisteria実践」

114

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

2024/1/19

講習会「Wisteria実践」

115