XcalableMP Workshop 2019
へようこそ!

PCクラスタコンソーシアム XcalableMP規格部会長
朴 泰祐（筑波大学 計算科学研究センター）

第７回XcalableMPワークショップ
1

2019/11/05

XcalableMPについて2

XcalableMP
分散メモリ環境を対象とした指示文ベースの並列言語
n 次世代並列プログラミング言語検討委員会 → 当部会において仕様を検討、提案。
n 2つの並列プログラミングモデルをサポート
nグローバルビューモデルによる定型的な並列化
n ローカルビューモデルによる自由度の高い並列化
nMPIとの inter-operabilityn 一般的なPCクラスタの他、京コンピュータ・ポスト京コンピュータでも稼働（「ポスト京」は予定）
n

第７回XcalableMPワークショップ
3

2019/11/05

PCクラスタコンソーシアムXcalableMP規格部会の活動
年3回の部会開催 (6/5, 10/5, 2/25, いずれも東京)

nnnn

新仕様V2.0に向けた新しい議論
n

動的なタスクモデルの実現検討

nMPIとの連携、C等の言語とのインタラクションのための関数群

現仕様V1.4規格のリリース
http://xcalablemp.org/download/spec/xmp-spec-1.4.pdf
アクセラレータに向けたOpenACCとの融合⇨XcalableACCの規格化の検討

第6回XMPワークショップ開催（11/1@筑波大）

nnn

共催：理化学研究所、筑波大学計算科学研究センター、高性能Fortran推進協
議会
Invited Talk 5件 (日独仏共同研究SPPEXAにおけるMYXプロジェクト支援による)

XMPの仕様・研究内容を本にまとめ出版予定nn

4

第７回XcalableMPワークショップ
Springerより
2019/11/05

XcalableMPの状況
nPCクラスタコンソーシアムで規格を議論
nn

2018/10 Version 1.4仕様を公開。
nCombined directivenTasklet機能 (draft、V2.0に向けて）

2015年より、次期仕様「XMP2.0」の検討を開始。
PGAS + Multitasking for Multicore
n Code transformation for Optimizationn XcalableACC for Acceleratornn

理研・筑波大で、リファレンス実装n

Omni XMP コンパイラ

第７回XcalableMPワークショップ
5

2019/11/05

www.xcalablemp.org
omni-compiler.org

Omni XcalableMP
.....
.....

n

理研R-CCSと筑波大で開発中のXMP処理系
nnnnn

XMP/C
XMP/Fortranオープンソース
トランスレータ + ランタイム(MPIベース)
OpenACC、XcalableACC対応

XMPソース

フロントエンド

.....
.....

中間表現
(XcodeML)

トランスレータ

.....
.....

中間表現
(XcodeML)

Omni XcalableMP
バックエンド(デコンパイラ)
XMPランタイム
.....
.....

C/Fortran
ソース

MPIライブラリ

ネイティブ
コンパイラ

第７回XcalableMPワークショップ
6

2019/11/05

実行形式

Data Distibution Using TemplatenTemplatenn

virtual array representing data(index) spacearray distribution, work-sharing must be done using templateExample)
0100

double array[100];#pragma xmp nodes p(4)declare node set#pragma xmp template t(0:99)

declare template0100

template t(0:99)#pragma xmp distribute t(BLOCK) on p
025p(1)50p(2)#pragma align array[i] with t(i)0array[]p(1)2019/11/05

75p(3)

100p(4)

distribute array : owner of t(i) has a[i]25

第７回XcalableMPワークショップ
7

distribute template50p(2)

75p(3)

100p(4)

Data Synchronization of Array(shadow)
nShadow Regionin XMP, memory access is always local
n duplicated overlapped data distributed onto other nodes
n data synchronization: reflect directive
n012345678910 1112 13 14 15a[]#pragma xmp shadow a[1:1]

declare shadow

NODE1
NODE2
NODE3
NODE4
#pragma xmp reflect a

第７回XcalableMPワークショップ
82019/11/05

synchronize shadow

Data Synchronization of Array(gather)pgather array data (collect entire elements)p#pragma xmp gather(var=list)#pragma xmp gather(var=array)

array[]
process0
process1
process2
process3第７回XcalableMPワークショップ
9

2019/11/05

all elements of the array get correct dataPut operation in taskletlput̲ready clause: indicates that the specified data may be written by theassociated PUT operationll

This clause has the dependence-type out for the specified data on a node since its values areoverwritten by the remote node.

put clause: indicates that the PUT operation may be performed in theassociated structured block.
l

At the beginning of the block, the task waits to receive the post notification with the tag by theput̲ready clause to indicates that the data is exposed in the target node for the PUT operations.

• When output dependencies for the data are satisf
ied before executing the block, the clause expose
s the data for the PUT operation from the specifie
d set of nodes by sending the post notifications tothese nodes, starting the PUT operations eventua
lly in remote nodes. Then, it waits until remote op
erations are done. When the task receives the completion notification of the PUT operation, the block is immediately scheduled.
• When the post notification is received, the task isscheduled to execute the calculation and PUT ope
ration in the block. When the execution of the blo
ck is finished, the data written by the PUT operati
on is flushed and the completion notification is se
nt to the node matched by the tag.

#pragma xmp nodes P(2)int A:[*], B, C, D, tag;

Node 1in(A)Node 2in(A)#pragma xmp tasklet in(A) out(B) on P(1)B = A;/* taskA */

start
post

#pragma xmp tasklet in(A) out(C) on P(1)C = A;/* taskB */completeout(A)#pragma xmp tasklet out(A) put(tag) on P(2)A:[1] = 1; /* taskC */waitin(A)#pragma xmp tasklet in(A) out(D) ¥put_ready(A, P(1), tag) on P(1)D = A;/* taskD */out(A)taskA

puttaskB

taskC

taskD

10

Block-Cholesky Benchmark: Results on OFPllComparison with “Parallel Loop” (PL) and Task-based“Put” and “Send/Recv”(OMPSS)
8000

Performance [GFLOPS]lMPI+OMP, PL, Send/Recv

7000

MPI+OMP, PL, Put

6000

MPI+OMP, Task, Send/Recv

5000

MPI+OMP, Task, Put
MPI+OMPSS, Task, Send/Recv

4000

3000
2000
1000
012481632# of nodes⇨ taskletの詳細については村井さんのトークで
11

XcalableMPのアクセラレータ対応(GPU & FPGA plan)12

XcalableACC (XACC)の基本的な考え方
XcalableACC = XcalableMP + OpenACC
XMPプログラム中にOpenACCのdirective記述を許す
n XMPが提供する分散メモリ並列処理 (PGAS)に従い、各ノードに分散された配列イメージを対象に、OpenACCのデータ移動・演算オフローディングを行う。
n OpenACC + MPI で記述する必要があった並列演算加速（GPUコンピューティング）をすっきりした形で記述でき、大規模並列化にも適している。
n XMP directiveと OpenACC directive は基本的に直交関係。
n第７回XcalableMPワークショップ
13

2019/11/05

Processing model of XACC

Array/Work#pragma acc device d = nvidia(0:3)

Distribution among nodes#0#1#pragma xmp reflect_init (a) deviceComm.
between CPUs

CPU

Distribution among ACCs.

ACCDirect Comm.
between ACCs

第７回XcalableMPワークショップ
14

2019/11/05#pragma xmp loop (i) on t(i)for (int i = 0; i < 100; i++){#pragma acc kernels loop on_device(d)for (int j = 0; j < 100; j++){a[i][j] = ...
}}#pragma xmp reflect_do (a)Example of XcalableACC program
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE];#pragma xmp nodes p(x, y)#pragma xmp template t(0:YSIZE−1, 0:XSIZE−1)#pragma xmp distribute t(block, block) onto p#pragma xmp align [j][i] with t(i,j) :: u, uu#pragma xmp shadow uu[1:1][1:1]…
#pragma acc data copy(u) copyin(uu){for(k=0; k<MAX_ITER; k++){#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)uu[x][y] = u[x][y];2-D Laplace Eq.

#pragma xmp reflect (uu) acc#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)u[x][y] = (uu[x-1][y]+uu[x+1][y]+uu[x][y-1]+uu[x][y+1])/4.0;} // end k} // end data1

Example of XcalableACC program
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE];#pragma xmp nodes p(x, y)#pragma xmp template t(0:YSIZE−1, 0:XSIZE−1)#pragma xmp distribute t(block, block) onto p#pragma xmp align [j][i] with t(i,j) :: u, uu#pragma xmp shadow uu[1:1][1:1]…
#pragma acc data copy(u) copyin(uu){for(k=0; k<MAX_ITER; k++){#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)uu[x][y] = u[x][y];#pragma xmp reflect (uu)

2-D Laplace Eq.
array distribution and “sleeve”
declarationexchange sleeves on array “uu”

#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)u[x][y] = (uu[x-1][y]+uu[x+1][y]+uu[x][y-1]+uu[x][y+1])/4.0;} // end k} // end data1

Example of XcalableACC program
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE];#pragma xmp nodes p(x, y)#pragma xmp template t(0:YSIZE−1, 0:XSIZE−1)#pragma xmp distribute t(block, block) onto p#pragma xmp align [j][i] with t(i,j) :: u, uu#pragma xmp shadow uu[1:1][1:1]…
#pragma acc data copy(u) copyin(uu){for(k=0; k<MAX_ITER; k++){#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)uu[x][y] = u[x][y];2-D Laplace Eq.
array distribution and “sleeve”
declaration
copy partial (distributed) array to
device memory
distributed array by XMP isprocessed according to OpenACCdirective#pragma xmp reflect (uu) acc

exchange sleeves on array “uu”

#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)u[x][y] = (uu[x-1][y]+uu[x+1][y]+uu[x][y-1]+uu[x][y+1])/4.0;} // end k} // end data“acc” clause indicates to target
the array on device memory

1

Example of XcalableACC program
double u[XSIZE][YSIZE], uu[XSIZE][YSIZE];#pragma xmp nodes p(x, y)#pragma xmp template t(0:YSIZE−1, 0:XSIZE−1)#pragma xmp distribute t(block, block) onto p#pragma xmp align [j][i] with t(i,j) :: u, uu#pragma xmp shadow uu[1:1][1:1]…
#pragma acc data copy(u) copyin(uu){for(k=0; k<MAX_ITER; k++){#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)uu[x][y] = u[x][y];2-D Laplace Eq.

copy partial (distributed) array todevice memory
distributed array by XMP isprocessed according to OpenACC
directive#pragma xmp reflect (uu) acc#pragma xmp loop (y,x) on t(y,x)#pragma acc parallel loop collapse(2)for(x=1; x<XSIZE-1; x++)for(y=1; y<YSIZE-1; y++)u[x][y] = (uu[x-1][y]+uu[x+1][y]+uu[x][y-1]+uu[x][y+1])/4.0;} // end k} // end data1

XACCにおける通信の選択
現時点ではGPU（NVIDIA GPU）を演算加速器として扱う。
n 一般的な並列通信機構によるn 通信のためにMPIコードが生成される。
n XMPとほぼ同じ配列分散を行い、デバイスメモリ上のデータに対する通信の最適化はGDR (GPU Direct RMA)等に対応したMPIによって行う。
n 特殊な並列通信機構向け実装
n CRESTで開発したTCA (Tightly Coupled Accelerator)
コンセプトに基づくGPU間直接通信支援機構である
PEACH2を対象とした実装
n 通信最適化はPEACH2に依存n

第７回XcalableMPワークショップ
19

2019/11/05

XACC on FPGAFPGA for HPC 研究が活性化しつつある
n FPGAは直球勝負でGPUと競ってもFLOPS的な勝ち目はないn

GPUとFPGAを補完的に使えるのでは？
n AiS (Accelerator in Swtich): FPGAを accelerator とinterconnection の両方に使う
FPGAの自律的動作によりGPU memoryのDMAとノード間通信までを行うことで従来の方法より低レイテンシ化
nn

ORNL (by J. Vetterʻs group)でOpenACCからOpenCL for
FPGAのトランスレータ開発が進んでいる
GPUとFPGAの両方のカーネルをOpenACCで記述し、適材適所的にデバイス割り当てをして連携させる
n XACCまで持っていけばFPGAの高速通信チャネル（最大100Gbps x 3ぐらい）を通信レイヤに使うことが可能に
n 共同研究実施に向け調整中n

第７回XcalableMPワークショップ
20

2019/11/05

Omni-OpenACCの拡張（筑波大＋理研R-CCS）
Omni-OpenACC: XcalableACC実装において開発された
n multi-device support のための新機能を実装
n“#pragma accomn ondevice(...)”
Device option: V100, Stratix10, ... (GPU and FPGA currently)
n 2種類のバックエンドコンパイラを使用
nn
nn

PGI OpenACC compiler for GPU
ORNL’s OpenARC research compiler for FPGA (under collaboration withORNL)

文科省「次世代領域研究開発」プログラムにおける「次世代演算通信融合型スーパーコンピュータの開発」プロジェクト（2017-2021, 代表：朴泰祐）

第７回XcalableMPワークショップ
21

2019/11/05

Multi-device OpenACC compilation

openacc.c
int main() {~~~~~~~~~~#pragma accomn ondevice(v100)~~~~~~~~~~~~~~~~~~~~

OpenACC

gpu.c

PGI Compiler(pgc++ -c)

objects

gpu.oGPU
executable
pgc++

OpenCLhost code#pragma acc ~~~=fpga~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
return 0;}

fpga.cfpga.cpp

OpenARC
OpenACC translator to split
each part for GPU and FPGA
(based on Omni OpenACCcompiler for XcalableACC)

a.out

OpenCLkernelfpga.og++ -c

fpga.claoc

aoc -cfpga.aoco

Intel OpenCL SDK for FPGA

第７回XcalableMPワークショップ
22

2019/11/05

fpga.aocx

FPGAmoduleMM & MV on GPU and SpMV on FPGA
#pragma acc data copyout(c[:N*N]) copyin(a[:N*N], b[:N*N])#pragma acc kernels#pragma acc data copyin(VAL[0:VAL_SIZE], COL_IND[0:VAL_SIZE],
ROW_PTR[0:N+1], B[0:N], N, K, VAL_SIZE) copyout(X_result[0:N]){#pragma acc parallel num_gangs(1) num_workers(1) vector_length(1) {#pragma acc loop gang(N/32) vector(32) independent〜for (i = 0; i < N; ++i)// 初期化計算#pragma acc loop{for(int i = 0; i < K; ++i){#pragma acc loop gang(N/32) vector(32) independent

temp_pap = 0.0f;

for (j = 0; j < N; ++j)

int j = 0, l = ROW_PTR_local[0];{#pragma acc loop reduction(+:temp_pap, temp_sum)

float sum = 0.0;
#pragma acc loop reduction(+:sum)for(int m = 0; m < VAL_SIZE-ROW_PTR_local[0]; ++m){for (k = 0; k < N; ++k)

temp_sum += p[COL_IND_local[l]] * VAL_local[l];sum += a[i * N + k] * b[k * N + j];++l;c[i * N + j] = sum;}}}if(l == ROW_PTR_local[j + 1]) {y[j] = temp_sum;

MM on GPU

temp_pap += p[j] * temp_sum;
temp_sum = 0.0f;++j;l = ROW_PTR_local[j]; } }...

SpMV on FPGA

Making matrix by MM on GPU, MV on GPU then transferred to FPGA for SpMV on FPGA
(Sp matrix data is kept on FPGA, only vector is transferred)
第７回XcalableMPワークショップ
23

2019/11/05(by R. Tsunashima, U. Tsukuba)

CUDA/OpenCL vs Multi-Device OpenACC

host CUDA
host OpenCL

kernel GPU

kernel FPGA

othersExecution time#chars in code
(by R. Tsunashima, U. Tsukuba)

Omni-OpenACCへのFPGA対応組み込みを行い、XcalableACCに展開⇨ FPGA間通信にFPGAが持つ光インターコネクトを直接利用
第７回XcalableMPワークショップ
24

2019/11/05

Cygnus ‒ multi-hybrid accelerated
supercomputer at U. Tsukuba

第７回XcalableMPワークショップ
25

2019/11/05

Single node configuration (Albireo)• Each node is equipped withboth IB EDR and FPGAdirect network• Some nodes are equippedwith both FPGAs and GPUs,
and other nodes are withGPUs only

Network switch
SINGLE(100Gbps x2)
NODE(with FPGA)CPUHCA26

2019/11/05

CPU

HCA

HCA

HCA

PCIe network (switch)

PCIe network (switch)

GP
U

GP
U

GPU

FPGA

Inter-FPGA
direct network(100Gbps x4)

第７回XcalableMPワークショップ

Network switch(100Gbps x2)

GP
U

FPGA

Inter-FPGA
direct network(100Gbps x4)

Two types of interconnection network
Inter-FPGA direct network(only for Albireo nodes)
FPGA

FPGA

FPGA

FPGA

FPGAFPGA

FPGA

FPGA

FPGA

InfiniBand HDR100/200 network for parallel processing
communication and shared file system access from all nodes

IB HDR100/200 Network (100Gbps x4/node)comp. … comp.
node
node
Deneb nodes

64 of FPGAs on Albireo nodes(2 FPGAS/node) areconnected by 8x8 2D torus
network without switch

27

第７回XcalableMPワークショップ
2019/11/05

comp.
node

comp.
node

Albireo nodes

…

comp.
node…

comp.
node…

Inter-FPGA torus network

For all computation nodes (Albireo and Deneb) areconnected by full-bisection Fat Tree network with 4channels of InfiniBand HDR100 (combined to HDR200switch) for parallel processing communication such asMPI, and also used to access to Lustre shared filesystem.

おわりにXcalableMPの規格は2.0へ、特にtask処理のノード内・ノード間並列実行への展開に重点
l XcalableACCをXcalableMPの正式なブランチへ⇨ GPU, FPGAへの対応
l 従来の OpenACC + MPI のコードをより簡単に記述可能に
l FPGAの特性を活かし、GPUとの協働による高性能処理へ
l28