neoSYCLSX-Aurora TSUBASAのための軽量SYCL実装
2022年4月15日
PCCC HPC-OSS部会ワークショップ「高性能クラスタ・プログラミング最前線」
東北大学サイバーサイエンスセンター

滝沢 寛之<takizawa@tohoku.ac.jp>

1

もくじ• 背景と目的• SX-Aurora TSUBASA向けSYCL実装• 性能評価• まとめと今後の展望

5

背景 ～ ヘテロ構成のベクトルシステム ～東北大学のスパコンAOBA
Vector Host (VH)
= Standard x86 CPU

Vector Engine (VE)
= Vector Processor with HBMs

NEC SX-Aurora TSUBASAを中核とするHPCシステム$ gcc matmul.c –o matmul$ ./matmul$ ncc matmul.c –o matmul$ ./matmul 高性能クラスタ・プログラミング最前線

6

動機 ～ 適材適所の必要性 ～• Vector Host と Vector Engine は両方とも• 汎用プロセッサ (単なるアクセラレータではない)• 標準的なプログラミング言語を単独で実行可能 (C/C++/Fortran etc)

• 異なる性能特性• ベクトル向きの処理であれば VEは高性能を発揮可能
Egawa et al. (PBMS2020)
適材適所の利用によって高い実行性能を達成

高性能クラスタ・プログラミング最前線

7

Vector Engine Offloading (VEO)• カーネル実行をVEにオフロードするための低レベルな専用API
Main program for VH

kernel program for VE#include <ve_offload.h>
int main(){/* Load "vehello" on VE node 0 */struct veo_proc_handle *proc = veo_proc_create(0);uint64_t handle = veo_load_library(proc, "./libvehello.so");struct veo_thr_ctxt *ctx = veo_context_open(proc);
struct veo_args *argp = veo_args_alloc();
uint64_t id = veo_call_async_by_name(ctx, handle, "hello", argp);
uint64_t retval;
veo_call_wait_result(ctx, id, &retval);
veo_args_free(argp);
veo_context_close(ctx);
veo_proc_destroy(proc);
return 0;}#include <stdio.h>#include <stdint.h>
uint64_t hello(){
printf("Hello, world¥n");return 0;}

このコードをコンパイルして共有オブジェクト./libvehello.soとして準備

SX-Aurora TSUBASA専用のアプリになってしまうことが問題

高性能クラスタ・プログラミング最前線

8SYCL Implementations in Development (2021/04/27)
SYCL, OpenCL and SPIR-V, as open industry
standards, enable flexible integration anddeployment of multiple acceleration technologies

DPC++

Uses LLVM/ClangPart of oneAPI

Any CPUNVIDIA
GPUs

Source CodetriSYCL

ComputeCpp

SX-AURORATSUBASA

CUDA andHIP/ROCm

Any CPUNVIDIA GPUs

neoSYCL

hipSYCL

Open source
test bed

Multiple
Backends

Any CPU

SYCL enables Khronos toinfluence ISO C++ to (eventually)
support heterogeneous compute

TBB

Any CPUNVIDIA GPUs
VEO

Any CPU

Level ZeroIntel CPUs
Intel GPUsIntel FPGAsIntel GPUsIntel CPUs
Intel GPUsIntel FPGAs
AMD GPUs(depends on driver stack)

Arm Mali
IMG PowerVR
Renesas R-Car

Level Zero
XILINX FPGAs
POCL(open-source OpenCL
supporting CPUs and NVIDIA
GPUs and more)

This work is licensed under a Creative Commons Attribution 4.0 International License

AMD GPUsIntel CPUs
NEC VEs

Intel GPUs

Multiple Backends in Development
There is development on supporting SYCL oneven more low-level frameworks.
For more information: http://sycl.tech© The Khronos® Group Inc. 2021 - Page 9

Ke et al. (HPC Asia 2021)

neoSYCL• VEOをバックエンドとして利用するSYCLインタフェースの実装• Kernel Outlining

queue.submit([&](handler &cgh) {auto ka = buf_a.template get_access<access::mode::write>(cgh);auto kb = buf_b.template get_access<access::mode::read>(cgh);auto kc = buf_c.template get_access<access::mode::read>(cgh);
cgh.parallel_for<class triad_kernel>(range<1>(array_size), [=](const id<1> &i) {ka[i] = kb[i] + scalar * kc[i];});});カーネル部分 (functor) を分離して関数内のループに変換
 VEOを使うことで、カーネル部分だけをVEにオフロード実行可能

int triad_kernel(double *ka, double *kb, double *kc, int N) {for (int i = 0; i < N; i++) {ka[i] = kb[i] + 0.4 * kc[i];}コード変換器以外はヘッダファイルとして実装
return 0;= SYCLインタフェースを提供するための簡易実装}

高性能クラスタ・プログラミング最前線

10

Ke et al. (HPC Asia 2021)

コンパイルの流れ
カーネル
コード
(C++)
単一ソース
コード(SYCL)

nccコンパイラ

コード変換器(Clang libtool)

共有オブジェクト
(.so)

VEO

ホスト
コード
(C++)

gccコンパイラ

高性能クラスタ・プログラミング最前線

実行可能ファイル

11

他のSYCL実装との比較

Ke et al. (HPC Asia 2021)• GPU-STREAM (CPU実行)• 1スレッド実行での性能比較• OpenMP実装とneoSYCL実装の性能差は平均1.6%程度• triSYCLが使っているBoostの多次元配列がメモリ確保とデータコピーを多く実行
高性能クラスタ・プログラミング最前線

12N-bodyによる評価

Ke et al. (HPC Asia 2021)• 2つのループ (速度と位置)をそれぞれカーネルとしてオフロード• カーネル以外の部分はVHで実行 (VHの方が高性能)

高性能クラスタ・プログラミング最前線

13

Ke et al. (HPC Asia 2021)

N-bodyによる評価結果

データサイズと実行時間

データサイズ 10,000

データサイズ 40,000• データ転送時間を考慮してもオフロードしたほうが高性能• いずれの実装でもカーネル実行時間が支配的• VEO(低レベルAPI)とneoSYCLはほぼ同等の性能 (性能差は1.7%程度)
高性能クラスタ・プログラミング最前線

14

BFS(Rodinia)による評価結果

Ke et al. (HPC Asia 2021)• 初期化部分がベクトルに不向きな実装のためオフロードが効果的

VE-only

neoSYCLVEO

SYCLインタフェースによりコード複雑さを低減可能

高性能クラスタ・プログラミング最前線15

独自ベンチマーク集の作成

Li et al. (PDCAT2021)• VEO-SYCL-Bench https://github.com/Tohoku-University-Takizawa-Lab/veo-sycl-bench• SYCL-Bench (Lal et al.@Euro-Par 2020)が動かない(？)• DPC++でもneoSYCLでもコンパイルエラー• 抽象化によるオーバヘッドの評価にはVEO版も必要
SYCL-Bench’s vec_add

VEO-SYCL-Bench’s vec_add

各ベンチマークのカーネル部分だけコピーしてそれ以外はシンプルなコードに書き直し
高性能クラスタ・プログラミング最前線

16

Li et al. (PDCAT2021)

生産性と性能の評価• neoSYCLとVEOとの性能差は平均わずか 0.42%• 高抽象化に伴う実行時オーバヘッド (abstraction penalty) は無視できるレベル• SYCLコードの複雑さはVEOコードよりも常に低い• SYCLプログラミングにより他プラットフォームとの移植性や生産性を改善

Avg.NLOC

Avg.CCN

Avg.token

高性能クラスタ・プログラミング最前線

SYCL

69.2
10653.0

VEO

86.4
12764.5

17

VHとVEを適材適所で利用する効果
Computation

Takizawa et al. (2019)

File I/O

140
各関数呼び出しを一回ずつオフロード

Execution Time [sec]

120
100
80

0.1

60
40

66.3
41.2

200

Only VH

注意：本評価ではSYCLは未使用

91.2

Only VE

2.8

2.4

45.0

42.8

VH+VE(functionoffloading)

関数呼び出しを複数まとめてオフロード
VH+VE(loop offloading)

IXPUG at HPC Asia 202218

Li et al. (PDCAT2021)

最近の開発状況• 規格適合性テストとコードリファクタリング• DPC++ test cases (basic tests)

https://github.com/intel/llvm/tree/sycl/sycl/test

未サポートの機能-

parallel_for_work_groupparallel_for_work_item
nd_itemndrange

OpenCL由来の機能をサポート
しても性能を期待できないため*.
* Takizawa et al. (PARCO2021)

Coverage: 94.8%
未サポートの機能 (parallel_for_work_groupとparallel_for_work_item)に関する2件以外はテスト合格

IXPUG at HPC Asia 2022

19

さらに多様なアクセラレータ
FPGAVEGPU

高性能クラスタ・プログラミング最前線

20

Ongoing Project

SYCL for FPGAThe ESSPER project
FPGAクラスタを利用するためのフレームワーク
led by 佐野健太郎 博士 (RIKEN R-CCS)SYCL for Collaboration between CPU and FPGA

AS class+ SYCL

高性能クラスタ・プログラミング最前線

21まとめ• Vector Engine (VE)はベクトル向きの処理で効果を発揮• メモリ律速になりがちな科学技術計算での活躍を期待• x86とVEはそれぞれ異なる性能特性  適材適所の分担が重要• neoSYCL：SX-ATでのSYCLによるオフロードプログラミング• カーネル部分を切り出してVEO経由でカーネル呼出しする簡易SYCL実装• 実行時のややこしい処理はVEOが担当 • 性能と生産性の両面で有望• 現在も精力的に開発中！ 近日公開予定 (未定)

• 今後の展開• 他のアクセラレータへのオフロードに利用 (ESSPER)• カーネル切り出し時で適切にコード変換することでさらなる性能向上の可能性
高性能クラスタ・プログラミング最前線22