oneAPIによるGPU+FPGAマルチヘテロ環境
プログラミングとアプリケーション実行朴泰祐

筑波大学・計算科学研究センター（本発表の共同研究者：柏野隆太，小林諒平，藤田典久）

2022/04/15

1

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

HPCにおけるアクセラレータ：GPUが席巻
n

GPUだけでいいか？
ベクトル計算型の数多くのアプリケーションで有効
n 定型ループ，幅広いSIMD実行で高効率
nSIMD (STMD)アーキテクチャによる大容量並列処理
n 高バンド幅のHBM (HBM2, 2e, 3）と小容量だが高速なローカルメモリによるHPC計算
nn

GPUだけの処理による限界（局所的に）並列性が不十分な場合
n 処理が定型的でない場合（条件分岐）n 粒度が小さくノード間通信が頻発する場合（CPU制御，kernelswitch）
n2

2022/04/15

PCCC-OSS-Programmingワークショップ

NVIDIA Tesla H100
Tensor Core GPU(from NVIDIA web page)

Center for Computational Sciences, Univ. of Tsukuba

FPGAのHPC向け利用n

近年のハイエンドFPGAn

真の意味でアプリケーションとハードの codesigning

プログラミング環境： OpenCL, C++, etc.
n 高性能の光インターコネクト: 100Gb x N
n 柔軟な精度コントロールが可能
n （比較的）低電力
nn

課題
プログラマビリティ: OpenCLは不十分，性能不足
n GPUより低い標準FLOPS
n メモリテクノロジ，バンド幅: GPUより1世代程度遅いn32022/04/15

PCCC-OSS-Programmingワークショップ

BittWare 520N with Intel Stratix10 FPGA
equipped with 4x 100Gbps opticalinterconnection interfaces

Center for Computational Sciences, Univ. of Tsukuba

HPCソリューションとしてのGPUとFPGAの対照的性質
deviceGPUFPGAparallelization

SIMD (x multi-group)

pipeline (x multi-group)

standard FLOPS😃😃 (1000x cores)😃 (~100x pipeline)

conditional branch😢 (warp divergence)😃 (both direction)

memory😃😃 (HBM2e)

😢 (DDR)→😃 (HBM2)interconnect

😢 (via host facility)😃😃 (own optical links)

programming😃 (CUDA, OpenACC, OpenMP)😢 (HDL)→😏 (HLS)

self-controllability😢 (slave device of host CPU)😃 (autonomic)

HPC applications

😃 (various fields)😢 (not much)

42022/04/15

PCCC-OSS-Programmingワークショップ

Center for Computational Sciences, Univ. of Tsukuba

CHARM: Cooperative Heterogeneous Acceleration withReconfigurable Multi-devices
Basic cluster with GPUs (by InfiniBand)

multi-physics/multi-scalecomplicated problem

invoke GPU/FPGA kernsls
data transfer via PCIe(invoked from FPGA)GPUCPUFPGAcomp.

PCIecomm.

CPU

PCIe
Application oriented
FPGA-FPGA communicationGPU
FPGAcomp.
comm.

100Gbps direct optical link

Cooperative computing with GPU and FPGA
5

2022/04/15

PCCC-OSS-Programmingワークショップ

FPGA Network

Center for Computational Sciences, Univ. of Tsukuba

Cygnus: world first multi-hybrid cluster with GPU+FPGA@ CCS, Univ. of Tsukuba (deployed by NEC)

Cygnus supercomputer at Center for Computational Sciences, Univ. of Tsukuba (Apr. 2019~)
85 nodes in total including 32 “Albireo” nodes with GPU+FPGA (other “Deneb” nodes have GPU only)
62022/04/15

PCCC-OSS-Programmingワークショップ

Center for Computational Sciences, Univ. of Tsukuba

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

SINGLENODE(with FPGA)CPU

Network switch(100Gbps x2)

HCA

2022/04/15

PCCC-OSS-Programmingワークショップ

CPUHCA

HCA

HCA

PCIe network (switch)

PCIe network (switch)

GPUGPUGPUFPGAInter-FPGA
direct network(100Gbps x4)

7Network switch(100Gbps x2)

GPU

FPGA

Inter-FPGA
direct network(100Gbps x4)

Center for Computational Sciences, Univ. of Tsukuba

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

FPGA

FPGA

FPGA

FPGA

FPGA

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

FPGA

IB HDR100/200 Network (100Gbps x4/node)
…

FPGA

FPGA

comp.
node

…

comp.
node

Deneb nodes
64 of FPGAs on Albireo nodes (2FPGAS/node)
are connected by 8x8 2D torus network
without switch82022/04/15

PCCC-OSS-Programmingワークショップ

comp.
nodecomp.
nodeAlbireo nodes

comp.
node…

comp.
node…

Inter-FPGA torus network

For all computation nodes (Albireo and Deneb) are connected by full-bisection
Fat Tree network with 4 channels of InfiniBand HDR100 (combined to HDR200switch) for parallel processing communication such as MPI, and also used toaccess to Lustre shared file system.

Center for Computational Sciences, Univ. of Tsukuba

GPUG FP P
U GA

F GP PG UAGPUIB HDR100 x4⇨ HDR200 x2

CPUCPU100Gbps x4
FPGA opticalnetwork x2
IB HDR200
switch (forfull-bisection
Fat-Tree)

1.2Tbps/nodeAlbireo node
92022/04/15

PCCC-OSS-Programmingワークショップ

Center for Computational Sciences, Univ. of Tsukuba

CHARMプログラミングの2つのアプローチn

OpenACCのみによる記述
nnnn

NVIDIA GPU compiler (PGI/NVIDIA) + FPGA research compiler (OpenARC@ORNL) を使う
OpenACC single codeに対し、拡張directiveによってGPUオフロードとFPGAオフロード部をユーザが明記、コードを分離して各バックエンドコンパイラに投げるコードを解析して分離するメタコンパイラ：MHOAT (Multi-Hetero OpenACC Translater) を筑波大・ORNL・
理研R-CCSで共同開発 *1⇒ 開発途上でコンパイルはできるが性能が不十分

従来のGPU用コード、FPGA用コードをそのまま複合的に利用
nnn

GPU: CUDA（またはOpenACC）with NVIDIA compiler
FPGA: OpenCL with Intel SDK
ホストコード、CUDAコード、OpenCLコードが混在し見通しと記述性が悪い⇒ コード資産は再利用しつつ見通しの良いオフローディング記述がしたい⇒ Intel oneAPIを利用すれば可能だがGPUとFPGAの両方を使った実アプリケーション例がない*1: Ryuta Tsunashima, et.al., “OpenACC unified programming environment for GPU and FPGA multi-hybrid acceleration”, HLPP2020

122022/04/15

PCCC-OSS-Programmingワークショップ

Center for Computational Sciences, Univ. of TsukubaGPU・FPGA混合プログラムをどのように記述するかn これまでのプログラミングモデル→ CUDA + OpenCL混合プログラミング*1−CUDA（GPU用）とOpenCL（FPGA用）で記述された2つのプログラムを分割コン
パイル&リンク→異なるAPIを使い分ける必要がある
CUDA API&OpenCL API

n コードの見通しを良くするためoneAPIを利用してCHARMを実装する− GPUとFPGAを統一的なスキームで制御できる*1: Ryohei Kobayashi et al. “Accelerating Radiative Transfer Simulation with GPU-FPGA Cooperative Computation”, ASAP2020

132022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

oneAPI概要
n oneAPIとは、インテルにより提唱された統合プログラミング
モデル−目的は異なるアーキテクチャ間における開発を簡略化・統一化すること
n oneAPIより、開発言語 Data

Pallalell C++（DPC++）が提供されている−異なるアーキテクチャ間で統一的な記述が可能

142022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

oneAPIによるCHARM実装
n DPC++でのプログラミングが推奨されている−(a) アプローチ−問題：既存のGPUやFPGAのコード
は、CUDAやOpenCLなど、他の言語で書かれている−DPC++による再実装はユーザーにとって負担
n oneAPIは他の言語で書かれたモジュールも使用可能−(b)アプローチ−コードの再利用が可能

15

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

oneAPI DPC++ Compiler
n oneAPI DPC++ コンパイラには2つのタイプがある．

−Intel oneAPI Base Toolkit に付属する公式版−オープンソースとして公開されているOSS版*1
NVIDIA サポートはOSS版のみ→本研究ではOSS版を利用する
n OSS版コンパイラ*1を利用して簡単なベクタ加算プログラムが，以下の環境で動作することを確認−NVIDIA GPU (V100), Intel FPGA (Stratix10)*1: https://github.com/intel/llvm

16

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

oneAPIによるCHARM実装
CUDA + OpenCL を別々に記述・合成

oneAPIから CUDA+OpenCL カーネルを呼び出す

cudaMemcpy(…);
CUDA API// Call the CUDA kernel
gpu_kernel<<<grid, block>>>(…);
cudaMemcpy(…);

clEnqueueWriteBuffer(…);
OpenCL API
clSetKernelArg(…);// Call the OpenCL kernel
clEnqueueTask(…, fpga_kernel,…);clEnqueueReadBuffer(…);

gpu_queue.memcpy(…);
gpu_queue.submit([&](handler& h) {h.interop_task([=](interop_handler ih) {// Call the CUDA kernel
gpu_kernel<<<grid, block>>>(…);});});
gpu_queue.memcpy(…);
fpga_queue.memcpy(…);
fpga_queue.submit([&](sycl::handler &h) {// Call the OpenCL kernelh.set_args(…);
h.single_task(fpga_kernel);});
fpga_queue.memcpy(…);

oneAPIを使うことで、統一的APIによりGPUとFPGAの制御が可能
詳細：Ryuta Kashino, et.al., “Multi-hetero Acceleration by GPU and FPGA for Astrophysics Simulation on Intel
oneAPI Environment”, HPC Asia 202217

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukubaターゲットコード: ARGOTn ARGOT (Accelerated Radiativetransfer on Grids using Oct-Tree)−初期宇宙における天体形成をシミュレーション
筑波大学計算科学研究センターにより開発−２つの部分計算からなる               

ARGOT法点光源の輻射輸送
n GPUが得意とする計算
nART法

空間に分散した光源の輻射輸送
n GPUが不得意とする計算
n FPGAによる高速化を行う
n182022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

Cosmic
RadiativeTransferinSimulation
Twocomputation
elements
ARGOT code: ARGOT method and ART method*ARGOT• ARGOT method: Point Source processing• ART method (Authentic Radiation Transfer): Diffused Photon processing

Point Source

19 2022/04/15 PCCC-OSS-Programmingワークショップ

Diffuse PhotonCosmic
Radiative
TransferinSimulation
Twocomputationelements
ARGOT code: ARGOT method and ART method*ARGOT• ARGOT method: Point Source processing• ART method (Authentic Radiation Transfer): Diffused Photon processing

ARGOT (Accelerated Radiative transfer on Grids using Oct-Tree) code
ARGOT scheme

ART scheme

for radiative transfer (RT)from point source

for RT from mattersspatially spreading outGPUGPU

GPU

CHARM
Point Source
GPU acceleration

Diffuse Photon

FPGA

20 2022/04/15 PCCC-OSS-Programmingワークショップ

FPGAFPGAFPGA acceleration

ARGOT法(CUDAによる実装)

n ARGOT法は点光源の輻射輸送を計算

target mesh gridn 八分木を用いて3次元空間に分散する点光源を表す−遠距離にある点光源の集合を単一の光源とみなせる
計算量: O(N^2) → O(NlogN)n 重力計算における Tree-Code に似た手法−GPU実装に適している

21 2022/04/15

21
PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

ART法(OpenCLによる実装)
n ART法は空間に広がる光源からの輻射輸送を計算−問題空間の端からレイが並行に直進ART法の実行時間は全体の90%を占める

n ART法はGPUに不向きな計算−レイの進行方向によってメッシュデータのメモリアクセスパターンが変化
ランダムアクセスに近いアクセスパターン−Atomic演算が必要−計算中に含まれる演算数が不十分
n FPGAを用いて高速化−FPGAのon-chipメモリを活用−一定サイズのベクトル処理が独立に大量に必要GPUの大規模SIMDが有効活用できない
詳細：Norihisa Fujita, et.al., “Accelerating Space Radiative Transfer on FPGA using OpenCL”, HEART2018

222022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

DPC++からCUDA/OpenCLカーネルをどのように呼び出すか
Example: CUDAカーネルの呼び出し
sycl::gpu_selector gpu_selector;
sycl::queue queue(gpu_selector);

デバイスの選択

dev_mem= malloc_device<>(…, sycl_queue);queue.submit([&](handler& h) {h.memcpy(dev_mem, hmem, …);})

Example: OpenCLカーネルの呼び出し
clGetDeviceIDs(…,&dev,…)
ocl_ctx = clCreateContext(…,&dev,…);
prg = clCreateProgramWithBinary(ocl_ctx, …);
ocl_kernel = clCreateKernel(prg, …);
ocl_queue = clCreateCommandQueue(ocl_ctx, dev,…);

データ転送

queue.submit([&](handler& h) {h.interop_task([=](interop_handler ih) {// Call the CUDA kernel
cuda_kernel<<<grid, block>>>(dev_mem,…);});});カーネル実行

sycl::context sycl_ctx(ocl_ctx);
sycl::kernel sycl_kernle(ocl_kernel, sycl_ctx);
sycl::queue queue(ocl_queue, sycl_ctx);
dev_mem= malloc_device<>(…, sycl_queue);queue.submit([&](sycl::handler &h) {h.memcpy(dev_mem, hmem, …);});
queue.submit([&](sycl::handler &h) {// Call the OpenCL kernelh.set_args(dev_mem,…);h.single_task(sycl_kernel);});

23

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

DPC++からCUDA/OpenCLカーネルをどのように呼び出すか
Example: CUDAカーネルの呼び出し

Example: OpenCLカーネルの呼び出し
OpenCLオブ

sycl::gpu_selector gpu_selector;
sycl::queue queue(gpu_selector);
dev_mem= malloc_device<>(…, sycl_queue);queue.submit([&](handler& h) {h.memcpy(dev_mem, hmem, …);})
queue.submit([&](handler& h) {h.interop_task([=](interop_handler ih) {// Call the CUDA kernel
cuda_kernel<<<grid, block>>>(dev_mem,…);});});clGetDeviceIDs(…,&dev,…)
ジェクトの生成ocl_ctx = clCreateContext(…,&dev,…);
prg = clCreateProgramWithBinary(ocl_ctx, …);
ocl_kernel = clCreateKernel(prg, …);
ocl_queue = clCreateCommandQueue(ocl_ctx, dev,…);
sycl::context sycl_ctx(ocl_ctx);
sycl::kernel sycl_kernle(ocl_kernel, sycl_ctx);
sycl::queue queue(ocl_queue, sycl_ctx);

dev_mem= malloc_device<>(…, sycl_queue);queue.submit([&](sycl::handler &h) {データ転送
h.memcpy(dev_mem, hmem, …);});
queue.submit([&](sycl::handler &h) {// Call the OpenCL kernelh.set_args(dev_mem,…);h.single_task(sycl_kernel);});

24

2022/04/15

DPC++オブジェクトへ変換

PCCC-OSS-Programmingワークショップ

カーネル実行oneAPIによるARGOT法(CUDA kernel)DPC++ code on CPU

n ARGOT法はCUDAで実装−DPC++からCUDAを呼び出す
n “oneAPI for CUDA”を利用−Codeplayにより開発
OSS版DPC++コンパイラでのみ利用可能−メモリ転送やカーネル実行は右図の通り

252022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

oneAPIによるART法(OpenCL kernel)
n ART法はOpenCLで実装// Construct SYCL version of the context
sycl::context sycl_context(ocl_context);−OpenCL実装の構成// Construct SYCL version of the kernel
sycl::kernel sycl_mesh_rw_kernel(ocl_mesh_rw_kernel, sycl_context);
sycl::kernel sycl_pe_kernel_0(ocl_pe_kernels[0], sycl_context);
sycl::kernel sycl_pe_kernel_1(ocl_pe_kernels[1], sycl_context);...
sycl::kernel sycl_pe_kernel_7(ocl_pe_kernels[7], sycl_context);// Construct SYCL version of the queue
sycl::queue sycl_mesh_rw_queue(ocl_mesh_rw_queue, sycl_context);
sycl::queue sycl_pe_queue_0(ocl_pe_queues[0], sycl_context);
sycl::queue sycl_pe_queue_1(ocl_pe_queues[1], sycl_context);...
sycl::queue sycl_pe_queue_7(ocl_pe_queues[7], sycl_context);

n OpenCLオブジェクトをDPC++オブジェクトに変換している

26

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

oneAPIによるART法(OpenCL kernel)
n 変換したDPC++オブジェクトに対してアクティベート操作を行う−Queueに対してカーネル実行を行っている−同様にQueueに対してメモリ転送を行っている// OpenCL kernel launched/// mesh reader & writer
sycl_mesh_rw_queue.submit([&](sycl::handler &h) {h.set_args(d_rmesh);h.single_task(sycl_mesh_reader_kernel);});/// ray tracing (rt) processing elements/// rt_x0_y0_z0
sycl_pe_queue_0.submit([&](sycl::handler &h) {h.set_args(...);
h.single_task(sycl_pe_kernel_0);});...
/// rt_x1_y1_z1
sycl_pe_queue_7.submit([&](sycl::handler &h) {h.set_args(...);
h.single_task(sycl_pe_kernel_7);
});

}

27

2022/04/15// memcpy (FPGA -> CPU)
sycl_mesh_rw_queue.submit([&](sycl::handler &h) {h.memcpy(h_rmesh, d_rmesh, ….);});
sycl_mesh_rw_queue.wait();

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

実験環境
n PPX (Pre PACS version X)−GPU・FPGA混載ノードをもつ
ヘテロジニアス並列計算機−筑波大学計算科学研究センター
により運用されている

29CPU0

CPU1

PCIe Switch

PCIe Switch

GPUFPGA

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

GPU-only vs GPU-FPGA 協調計算（ARGOT全体）

R. Kobayashi, et. al., “Accelerating Radiative Transfer Simulation with GPU-FPGA Cooperative Computation”, ASAP2020, Jul. 2020.

30

2021/06/21

HEART2021 Special Session

性能評価
n 問題サイズ: 32!

n Others−ベクトル並列化が可能
GPUによる高速化が効果的
n ART

Better−ART法のGPU実装はCPUと同程度の性能−FPGAにより高速化
n GPU:CUDA+FPGA:OpenCL

実行時間は1.5%だけ増加−GPUのみの場合と比べて10xの性能向上
n oneAPI vs CUDA+OpenCL−oneAPI版の実行時間が1.5%増加→oneAPI処理のオーバヘッドは無視できる

312022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

ARGOT法およびART法の非同期実行
n 本来ARGOT法とART法は2種類の物理現象を独立に計算し、最後に結果を足し合わせている

0.7

n ART法の実行時間は ARGOT法の実行時間に隠蔽される−全体として1.38xの高速化

Lower is better−OpenMPを利用して，各Queueにactionを非同期に投入−Queueに投入されたactionは非同期実行される

Execution time [s]n 各Queueは非同期に実行可能

Others

0.6

ARGOT

0.5

ART

0.4
0.3
0.2
0.1
0
ARGOT(GPU)ART(FPGA) / CUDA + OpenCL

ARGOT(GPU)ART(FPGA) / CUDA +OpenCL over oneAPI

ARGOT(GPU)ART(FPGA) / CUDA +OpenCL over oneAPI async*柏野隆太他, “oneAPIを用いたGPU・FPGA混載ノードにおける宇宙物理シミュレーションコードARGOTの実装”, 第183回HPC研究会

32

2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

まとめn 筑波大学計算科学研究センターではGPU-FPGA協調計算である CHARM (Cooperative

Heterogeneous Acceleration with Reconfigurable Multidevices)コンセプトを実現する様々な方法を研究している−プログラミングが最も大きい課題である
n oneAPIはGPU-FPGA協調計算を実現するためのプログラミングモデルの一つであると考えられる
n 既存のCUDA+OpenCLで記述されたARGOTコードをoneAPIに移植し，追加オーバーヘッドが無視できることを確認−また，Queueの非同期実行により性能向上が図れることを確認−この手法は，既存のHPCアプリケーションにCHARMのコンセプトを適応させる際に役立つ
n 将来的に世界初のヘテロジニアス（GPU＋FPGA）スーパーコンピュータCygnusでも積極的に活用する予定
CygnusでのoneAPI実行がいよいよ利用可能になった
OpenCLベースのFPGA間直接通信ツールCIRCUSとGPU側のMPIを組み合わせ、並列マルチヘテロ演算加速による実アプリケーション実行が可能に
33
2022/04/15

PCCC-OSS-Programmingワークショップ
Center for Computational Sciences, Univ. of Tsukuba

研究チームnn

HPC Research Division at CCS, Univ. of Tsukuba
Architecture Team in High Performance Computing System Lab., Dept. of Computer Science, Univ.
of Tsukuba

Ryohei Kobayashi
nNorihisa Fujita

Yoshiki Yamaguchi

Ryuta Kashino

Ryuta Tsunashima

Astrophysics Research Division at CCS, Univ. of Tsukuba

Masayuki Umemura
34

2022/04/15

Kohji Yoshikawa

PCCC-OSS-Programmingワークショップ

Center for Computational Sciences, Univ. of Tsukuba