第11回 配信講義 計算科学技術特論A（2025）
881 Views
June 17, 25
スライド概要
第11回 6 月 26日 ベンダーニュートラルGPUコンピューティング
GPUコンピューティングの基礎と，ベンダーニュートラルな（特定のベンダーに依存しない）GPU化手法について紹介する．
R-CCS 計算科学研究推進室
関連スライド
各ページのテキスト
ベンダーニュートラル GPUコンピューティング入門 三木 洋平 （東京大学 情報基盤センター） 計算科学技術特論A（2025） 第11回
自己紹介 2005-2009 2009-2011 2011-2014 2011-2013 2014-2017 2017-2024 2024- 2 筑波大学 第一学群 自然学類 筑波大学大学院 数理物質科学研究科 物理学専攻（博士前期課程） 筑波大学大学院 数理物質科学研究科 物理学専攻（博士後期課程） 筑波大学大学院 システム情報工学研究科 コンピュータサイエンス専攻（博士前期課程） 筑波大学 計算科学研究センター 研究員 東京大学 情報基盤センター 助教 東京大学 情報基盤センター 准教授 • 専門分野：宇宙物理学，高性能計算 • 銀河の形成・進化（銀河考古学，銀河衝突） • 銀河と中心ブラックホールの共進化 • GPUを用いた演算加速 • 重力多体計算，数値流体力学 • （準）力学平衡状態生成コードの作成 2025/6/26 計算科学技術特論A（2025） 第11回
Contents 3 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
Contents 4 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
GPU（Graphics Processing Units） • 元々は画像処理を行うために特化していた専用プロセッサを汎用計算にも 使えるように拡張 • 最近のGPUには（主に深層学習用に）行列積向けユニットも • 高い演算性能，太いメモリバンド幅，消費電力あたり性能も高い • 太古の昔には安かったが，最近はとても高い • 多数のコア（最新GPUは１万越え）を搭載した並列計算機 • NVIDIA A100 (SXM, PCIe): 64 × 108 SMs = 6912 cores • NVIDIA H100 (SXM): 128 × 132 SMs = 16896 cores • NVIDIA H100 (PCIe): 128 × 114 SMs = 14592 cores • AMD MI250X: 64 × 110 CUs × 2 GCDs = 14080 cores • AMD MI250: 64 × 104 CUs × 2 GCDs = 13312 cores • ざっくり言うと，コア数の10倍以上の並列度があると性能を発揮しやすい • スレッド並列（e.g., OpenMP）的考えが分かっていると良い 2025/6/26 計算科学技術特論A（2025） 第11回 5
Green500 Ranking (June 2025) TOP 500 System 1 259 JEDI, EuroHPC/FZJ, Germany 2 148 3 484 Adastra 2, GENCI-CINES, France 4 5 255 6 66 Capella, TU Dresden, ZIH, Germany 7 304 8 85 Helios GPU, Cyfronet, Poland 9 Accelerator Cores 6 https://www.top500.org/lists/green500 HPL Rmax [PFlop/s] Power GFLOPS/W [kW] NVIDIA GH200 Superchip 19,584 4.50 67 72.733 NVIDIA GH200 Superchip 47,328 9.86 160 70.912 AMD Instinct MI300A 16,128 2.53 37 69.098 183 Isambard-AI phase 1, University of Bristol, UK NVIDIA GH200 Superchip 34,272 7.42 117 68.835 Otus (GPU only), Universitaet Paderborn PC2, Germany NVIDIA H100 SXM5 80GB 19,440 4.66 NVIDIA H100 SXM5 94GB 85,248 24.06 445 68.053 SSC-24 Energy Module, Samsung Electronics, NVIDIA H100 SXM5 80GB South Korea 11,200 3.82 69 67.251 NVIDIA GH200 Superchip 89,760 19.14 317 66.948 399 AMD Ouranos, Atos, France AMD Instinct MI300A 16,632 2.99 48 66.464 10 412 Henri, Flatiron Institute, USA NVIDIA H100 80GB PCIe 8,288 2.88 44 65.396 99 131 184,320 11.16 674 16.575 119 126 Carpenter, ERDC DSRC, USA 276,480 11.62 1,100 10.561 ROMEO-2025, ROMEO HPC Center Champagne-Ardenne, France PRIMEHPC FX1000, Central Weather Administration, Taiwan (CPU-only: Fujitsu A64FX) (x86 CPU-only: AMD EPYC 9654) 68.177
（NVIDIAの）GPUの構成（1/2） NVIDIA H100 Tensor Core GPU Architecture 2025/6/26 計算科学技術特論A（2025） 第11回 7
（NVIDIAの）GPUの構成（2/2） • SM：Streaming Multiprocessor • AMDの場合には Compute Unit (CU) • Intelの場合には Execution Unit (EU) • SMの中での構造は気にしなくても 性能が出せる • 注：32スレッドのグループ（ワープ）内で 同じ動作をするように意識しておく • SM内ではL1キャッシュ/シェアードメモリ を共有（H100では256KB） • 配分は（ある程度）調節可能 • シェアードメモリはCUDAでは使えるが OpenACCでは（明示的には）使えない機能 • グローバルメモリよりも圧倒的に速い 2025/6/26 NVIDIA H100 Tensor Core GPU Architecture 計算科学技術特論A（2025） 第11回 8
GPUプログラミングでの基本思想 • このスライドではCUDA用語で説明 • 演算器の数よりも多数のスレッドを 立てて計算 • 各種レイテンシを隠蔽するため • スレッドブロックが基本単位 • ブロックあたりのスレッド数は32の倍 数，できれば128以上が推奨 • SMに複数ブロックを割り当てるのが 普通 • ブロック内では32スレッド単位で動作 （この組をワープと呼ぶ） • スレッドブロックの集合をグリッドと 呼ぶが，意識しなくてOK 2025/6/26 https://developer.nvidia.com/blog/cudarefresher-cuda-programming-model/ 計算科学技術特論A（2025） 第11回 9
Contents 10 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
GPUスパコンの近況 11 • GPUは最近のスパコンでのメジャーな演算加速器（特に上位システムで顕著） • 「GPU = NVIDIAのGPU」と言って良いぐらいの独占状態 • 国内では，現在HPCIに資源提供されている全てのGPUスパコンはNVIDIA製 • ２024年度稼働開始のTSUBAME4.0，玄界，Miyabi もNVIDIA製GPUを搭載 • 2025年度はQST/NIFS，筑波大CCSがそれぞれAMD MI300Aを搭載したシステムを運用開始 • 海外のハイエンドスパコンではAMD，Intel製GPUも採用 • Frontier などAMD製GPUを搭載したシステム，AuroraなどIntel製GPUを搭載したシステム • TOP５００の上位にAMD，Intel，NVIDIA製GPUがランクイン • GPUベンダー間の競争が活性化 • NVIDIA製GPUの性能向上: P100, V100, A100の間は各世代 2倍➔H100では約3倍 • 発散するプログラミング環境への対応も必要 • ポスト富岳（富岳NEXT）は演算加速器を搭載する • どのベンダー製の演算加速器なのか（= どうプログラミングすれば良いか）が不明 • HPCI構成機関が今後どのようなシステムを入れてくるかも不明，拠点ごとの多様性も？ • ➔ ベンダーニュートラルな実装手法を採用したいが，実現可能か？ 十分な性能は出せるか？ 2025/6/26 計算科学技術特論A（2025） 第11回
GPU向けのプログラミング環境 12 2024年2月のPCCC AI/HPC OSS活用WSでの講演資料 （https://www.pccluster.org/ja/event/data/240205_pccc_wsAI-HPC-OSS_06_hanawa-miki.pdf） OpenACC (Cray/HPE compiler) 指示文 OpenACC OpenMP target (NVHPC) OpenMP target (ROCm) OpenMP target (oneAPI) OpenMP target (Cray/HPE compiler) Standard Parallelism Fortran 2018/202X, C++17 標準言語仕様 Std Par., SYCL (AdaptiveCpp) (plugins by Codeplay) 低レベル 2025/6/26 SYCL (DPC++) CUDA HIP NVIDIA GPU AMD GPU 計算科学技術特論A（2025） 第11回 (chipStar) Intel GPU
GPUプログラミング手法の比較表 13 • 独断と偏見に基づく不完全な比較表であることに注意 • 特に Fortran はケアできていない（Fortranを読み書きできない人が作った資料） • Fortran対応で「OK」というのは動作するという意味．C/C++と同程度の性能が出るかは別問題 手法 ベース言語など 強み CUDA C++ C++ 詳細な最適化可能 記述量が多い CUDA 最新機能が使える GPU専用コード Fortran NVIDIA限定 HIP C++ 詳細な最適化可能 記述量が多い GPUFORT？ ほぼCUDA C++ GPU専用コード （開発停滞中？） AMD, NVIDIA （Intel: chipStar?） SYCL C++ 詳細な最適化可能 記述量が多い ラムダ式 ラムダ式 N/A Intel, NVIDIA, AMD OpenACC 指示文 移植コスト低 CUDAより遅い OK CPUコードと共通 （メモリ律速なら 化可能 それなり？） ほぼNVIDIA限定 （HPE Cray コンパイラ はAMDも対応） OpenMP 指示文 （target指示文） 移植コスト低 CUDAより遅い OK CPUコードと共通 （メモリ律速なら 化可能 それなり？） NVIDIA, AMD, Intel 言語の標準規格 C++17以降 CPUでも同じ 2025/6/26 Fortran 2008 コードが動く 弱み Fortran対応 多くの制約あり Fortran CUDAより遅い 2008 対象GPU NVIDIA，Intel （AMD: roc-stdpar?）
2018年にあったやり取り 14 • 木構築なども含めて全てGPU上で動作する重力ツリーコードGOTHICを CUDA Cで開発（YM & Umemura 2017, New Astronomy, 52, 65-81） • 論文を見た海外の研究者からコンタクト ➢加速度に使っている定式化が特殊だが，GOTHICに組み込めるか？ ➢（論文を見る限りは，おそらく）可能 ➢自分の環境でどのぐらいの実行時間になるか見積もってくれませんか？ ➢アクセスできるGPUの情報を教えてください ➢AMD Vega 64です ➢CUDA実装なので，AMD製GPU上では動きません．NVIDIA製GPUへのアクセス はありませんか？ ➢以降返信なし • 興味を持った人がアクセスできる・動かしたい環境と開発者が想定する環境 が一致する保証はない（OSS開発者はこの視点を持っておくことも重要） • ベンダーロックインは自分が困ることも多々あるが，共同研究の機会を逃すことにも 2025/6/26 計算科学技術特論A（2025） 第11回
FortranベースのGPU化とベンダーロックイン 15 • AMD/IntelのFortranサポートは，基本的に指示文（OpenMP target） • 指示文だけでGPU化でき，かつ性能に満足できるならば ベンダーニュートラル化は可能 • HIPやSYCLはFortranを直接サポートしていない • NVIDIAの場合にはCUDA Fortranがある • 指示文だけでGPU化できない and/or 十分な性能が 出せない場合には 1. 一部をC/C++に書き換えて， 2. HIPやSYCLでCUDAレベルの実装を作り， 3. Fortran側から呼び出す すべてのループ GPU化で高速化可能なループ （CUDA/HIP/SYCLなど） 指示文による GPU化で高速化 可能なループ • ➔ Fortran だけで完結しない = C/C++ のコードも読み書きできないといけない • Fortran で完結，かつ性能も犠牲にしたくないと思うと何が起こる？ • ベンダーニュートラルな実装は，（少なくとも現時点では）実質的に不可能 • 市場原理が働くなるため，調達価格は高止まりする（競合相手がいれば安くなる） • 結果的に，導入されるシステム規模が小さくなる = 実はFortranユーザも損している • C/C++ユーザからすれば，自分たちとは関係ないところで勝手にシステム規模を小さくされた話
Contents 16 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
指示文ベースでGPU化したい場合の選択肢 • OpenACC • GPU向けのメジャーな指示文 • PGIがNVIDIAに買収された結果，NVIDIA色が強くなってしまった • AMD，Intelは（きっと）サポートしない • HPE Crayコンパイラであれば，AMD GPU向けのOpenACCもサポート • IntelはOpenACCからOpenMP targetへの変換ツールを開発中 • ➔GPUベンダー（AMD，Intel）による直接支援が受けられない • OpenMPのtarget指示文 • OpenMP 4.0以降でアクセラレータへのオフロードがサポート • OpenMP 5.0で loop 指示節が追加，OpenACC的実装も可能に • NVIDIA，AMD，Intel 全てのGPU向けにサポートされる • 現時点ではOpenACCの全ての機能に対応できていない • 非同期実行の（細やかな）制御など • 実装時には，使う指示文についても選択する必要がある 2025/6/26 計算科学技術特論A（2025） 第11回 17
指示文（OpenACC/OpenMP）でのGPU化方針 18 1. Unified/Managed Memory を使って実装 • NVIDIA GH200の場合: CPU/GPU側のメモリは相互に読み書き可能 • -gpu=mem:unified を指定（-gpu=mem:unified:nomanagedalloc の方がおすすめ） 注：この書き方はNVIDIA HPC SDK用の記法 • 通常環境の場合: GPU上のメモリ確保，CPU-GPU 間のデータ転送は全てお任せ • -gpu=mem:managed を指定 • まずは演算部分のGPU実装に注力 • （マルチコアCPU向けの）OpenMP実装されていれば， #pragma omp parallel for をGPU向けの指示文に置き換えていく 2. （Unified/Managed Memory 実装での性能に満足できない場合） データ指示文を使ってコードをアップデート • Managed Memory では，一旦メモリを読んで，ページフォルトがあればハードウェ アレベルでページ単位で転送という仕組み • 必要なデータ転送は自分で指示した方が必然的に速くなる • Unified/Managed Memory ではCPU上のアドレスとGPU上のアドレスを同一 視してしまうので，GPUDirect系の機能を活用する際に不利
OpenACCでの実装（演算部分） • #pragma omp parallel for を置き換えていく • #pragma acc kernels はGPU化するかコンパイラが判断（GPU化しないことも） • #pragma acc parallel はGPU化できるとユーザが保証（GPU化される） • 性能的に特に重要なものについては，スレッド数の調整も • vector_length（スレッド数） をkernels/parallel指示文に付与して示唆 • vector（スレッド数） をloop指示文に付与しても同じことができる #pragma omp parallel for for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma acc kernels #pragma acc loop independent for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma omp parallel for for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma acc kernels vector_length(NTHREADS) #pragma acc loop independent for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } 2025/6/26 計算科学技術特論A（2025） 第11回 19
OpenMPでの実装（演算部分：distribute） • #pragma omp parallel for を置き換えていく • 性能的に特に重要なものについては，スレッド数の調整も • thread_limit（スレッド数） としてコンパイラに示唆（強制はできない） • num_teams(チーム数) という方法もある（が，あまり使いたくはない） • 全体の問題サイズが決まっている際にはスレッド数の指定と等価になるが，実行時にパラメータ ファイルを読み込んで問題サイズを設定するような実装には不向き #pragma omp parallel for for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma omp target teams distribute parallel for simd for (type::int_idx ii = 0U; ii < Ni; ii++) { //ループ内の実装は省略 } #pragma omp parallel for for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma omp target teams distribute parallel for simd thread_limit(NTHREADS) for (type::int_idx ii = 0U; ii < Ni; ii++) { //ループ内の実装は省略 } 2025/6/26 計算科学技術特論A（2025） 第11回 20
OpenMPでの実装（演算部分：loop） 21 • #pragma omp parallel for を置き換えていく • （先述の）distributeよりも，コンパイラに多くを委ねる実装法 • OpenMP 5.0 で導入された記法 • 性能的に特に重要なものについては，スレッド数の調整も • thread_limit（スレッド数） としてコンパイラに示唆（強制はできない） • num_teams(チーム数) という方法もある（が，あまり使いたくはない） #pragma omp parallel for for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma omp target teams loop for (type::int_idx ii = 0U; ii < Ni; ii++) { //ループ内の実装は省略 } #pragma omp parallel for for (type::int_idx ii = 0U; ii < Ni; ii++) { // ループ内の実装は省略 } #pragma omp target teams loop thread_limit(NTHREADS) for (type::int_idx ii = 0U; ii < Ni; ii++) { //ループ内の実装は省略 } 2025/6/26 計算科学技術特論A（2025） 第11回
OpenACCでの実装（データ転送部分） • Unified/Managed Memoryを使わない場合のみ必要 • GPU上に置くべきデータ，必要なデータ転送を指定 // メモリ確保 #pragma acc enter data create(pos_ptr [0:num], vel_ptr [0:num], acc_ptr [0:num]) // CPU → GPU のデータ転送 #pragma acc update device(pos_ptr [0:num], vel_ptr [0:num]) // GPU → CPU のデータ転送 #pragma acc update host(acc_ptr [0:num]) // メモリ解放 #pragma acc exit data delete (pos_ptr [0:num], vel_ptr [0:num], acc_ptr [0:num]) 2025/6/26 計算科学技術特論A（2025） 第11回 22
OpenMPでの実装（データ転送部分） • Unified/Managed Memoryを使わない場合のみ必要 • GPU上に置くべきデータ，必要なデータ転送を指定 // メモリ確保 #pragma omp target enter data map(alloc : pos_ptr [0:num], vel_ptr [0:num], acc_ptr [0:num]) // CPU → GPU のデータ転送 #pragma omp target update to(pos_ptr [0:num], vel_ptr [0:num]) // GPU → CPU のデータ転送 #pragma omp target update from(acc_ptr [0:num]) // メモリ解放 #pragma omp target exit data map(delete : pos_ptr [0:num], vel_ptr [0:num], acc_ptr [0:num]) 2025/6/26 計算科学技術特論A（2025） 第11回 23
Contents 24 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
指示文によるGPU化が意味するものは何か？ 25 • 具体的にどうGPU化するかはコンパイラ任せ • 細部まで指定したい人は CUDA/HIP/SYCL を使う • つまり，実装の詳細はブラックボックスであってもGPU化されていればOK • CUDA/HIP/SYCL で全力最適化した場合よりも性能が低いことは受け入 れているユーザ層 • 最大性能よりも，CPUコードとの互換性や移植工数削減の方を優先 • OpenACC と OpenMP targetの性能差は CUDA/HIP/SYCL からの性能差に 比べれば小さいはず • 本質的には OpenACC or OpenMP target？ はどちらでも良い（はず） • 何か指示文的なものを書いたらGPU化してくれればOK • 新しい指示文を作るようなことは検討しない • 新しい移行コストが発生するだけ • 新しいコンパイラを開発するようなことは検討しない • コンパイラの開発・メンテナンスが止まった時にユーザーを道連れにしてしまう • 大変なので決してやりたくない，プロの方々にお任せする
マクロを用いた指示文のブラックボックス化 • Solomon（Simple Off-LOading Macros Orchestrating multiple Notations）を実装 • Miki & Hanawa (2024, IEEE Access) • バックエンドで OpenACC or OpenMP target に展開 • Fallback mode （マルチコアCPU向けのOpenMPに展開）も実装済み • ユーザ的にはオーバーオールのフラグ制御だけで OpenACC or OpenMP target の切り替えが可能 • NVIDIA GPU 上では OpenACC で， AMD/Intel GPU 上では OpenMP target で動かす，ということが可能になる • 最適化レベルを揃えた上で OpenACC と OpenMP target の性能比較 • コンパイラではないのでベンダー製のコンパイラ性能をそのまま利用できる • （GPU向けプログラミングに詳しい人は）HIP の指示文版とイメージすると良い • 自作コンパイラの場合には，最新機能への追随のためのコストが継続的に生じる • 開発者が更新をさぼっても，自分でマクロを付け足すことも簡単 • 新コンパイラ/新指示文の実装であれば，一般ユーザはほぼ手出しできない 2025/6/26 計算科学技術特論A（2025） 第11回 26
Solomonを用いた実装例
27
• OpenACCとOpenMP両方に対応した指示文の追加例（右側）
• 場合分け（ACC for GPU, OMP for CPUなど）がかなり煩雑
• $ nvc++ -acc=multicore -mp=gpu … も（見かけないが）実は可能
• プリプロセッサマクロを用いてインターフェースを統合するライブラリを開発
• https://github.com/ymiki-repo/solomon で公開
• Miki & Hanawa (2024, IEEE Access)
• 手品の種: _Pragma()形式で指示文を記述
• 対応しているバックエンド:
• OpenACC, OpenMP target, OpenMP
• やる気が出るのはどちらの手法?
• 通常の（煩雑な）実装方法
•
Solomon を用いて簡易化した手法
OFFLOAD(AS_INDEPENDENT, NUM_THREADS(NTHREADS))
for (int32_t i = 0; i < N; i++) {
2025/6/26
#ifdef OFFLOAD_BY_OPENACC
#pragma acc kernels vector_length(NTHREADS)
#pragma acc loop independent
#endif // OFFLOAD_BY_OPENACC
#ifdef OFFLOAD_BY_OPENMP_TARGET
#ifdef OFFLOAD_BY_OPENMP_TARGET_LOOP
#pragma omp target teams loop
thread_limit(NTHREADS)
#else // OFFLOAD_BY_OPENMP_TARGET_LOOP
#pragma omp target teams distribute parallel
for simd thread_limit(NTHREADS)
#endif // OFFLOAD_BY_OPENMP_TARGET_LOOP
#endif // OFFLOAD_BY_OPENMP_TARGET
for (int32_t i = 0; i < N; i++) {
Solomon のインターフェース（指示文） • 簡易記法，OpenACC的記法，OpenMP的記法の3種を提供 28
Solomon のインターフェース（指示節） • 簡易記法，OpenACC的記法，OpenMP的記法の3種を提供 • 1つの指示文に対して記述の中で複数の記法を混ぜてもOK 29
バックエンドの切替方法 30 • コンパイル時にフラグを渡すだけでOK • 各コンパイラに対する OpenACC / OpenMP target を有効化するため のフラグは別途必要 • OpenACC 無効化時には， -DOFFLOAD_BY_OPENACCも自動的に無効化される • -DOFFLOAD_BY_OPENACCと-DOFFLOAD_BY_OPENMP_TARGETを両方指定した 場合 • OpenACC を用いてGPU化 • -DOFFLOAD_BY_OPENACCと-DOFFLOAD_BY_OPENMP_TARGETをどちらも指定 しなかった場合 • マルチコアCPU向けにOpenMPでスレッド並列化 2025/6/26 計算科学技術特論A（2025） 第11回
Solomon 使用にあたっての注意点，推奨事項など 31 • 指示節などはカンマ区切りで入力する • 各指示節が適用可能かを（Solomon側で）判定し，OKなもののみを有効化する実装 • OpenMP的記法をOpenACCバックエンドで動かすために実装した機能 • _Pragma(“omp target teams loop collapse(3) thread_limit(128)”) ➔ _Pragma(“acc kernels vector_length(128)”) _Pragma(“acc loop collapse(3)”) • 指示節の適切な振り分けは，Solomon 側で対応すべき機能の一つ • OpenACCでもOFFLOAD(…), PRAGMA_ACC_[KERNELS PARALLEL]_LOOP(…) の 使用を推奨 • PRAGMA_ACC_[KERNELS PARALLEL](…)とPRAGMA_ACC_LOOP(…)に分けて書くと， OpenMP target への適切な変換が困難になる（LOOP側に入力した指示節が渡らなくなる） • AS_INDEPENDENT は（複数の）指示節の先頭に置いておく（必須） • バックエンドを OpenMP にした時には simd に変換されるが，simd は構文の一部分である ため，中間に他の指示節が紛れ込むとエラーとなってしまう • 指示節候補をソートした後で判定・有効化するよう実装すれば回避できるはず（未対応） • PRAGMA_OMP_TARGET_DATA(…) は非推奨 • OpenACC における data, host_data と OpenMP target の data が対応 • バックエンドを OpenACC にした時にエラーとなる場合がある • 推奨:PRAGMA_ACC_[DATA HOST_DATA](…), DATA_ACCESS_BY_[DEVICE HOST](…)
Contents 32 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
N体計算（重力多体計算） 33 • 粒子どうしに働く自己重力による系の時間進化を，運動方程式に基づいて計 算 • データ量: 𝒪(𝑁) • 重力計算: 𝒪(𝑁 2 ) • 時間積分: 𝒪(𝑁) • N体業界的によく使う用語 i j j • i-粒子: 重力を受ける粒子 • j-粒子: 重力を及ぼす粒子 j j • 直接法のソルバーはツリーコードのバックエンドとして使えるので，高速化し ておく価値が高い（また，衝突系N体計算であれば直接法を用いる） 2025/6/26 計算科学技術特論A（2025） 第11回
実装例（N体計算，簡易記法）
34
• CPU上で初期条件を生成，GPUに粒子データ転送後に重力計算
set_uniform_sphere(num, pos, vel, Mtot, rad, virial, newton);
MEMCPY_H2D(pos [0:num], vel [0:num])
calc_acc(num, pos, acc, num, pos, eps);
• MEMCPY_H2D() でデータ転送
• 重力計算関数（ほぼ省略版）
void calc_acc(…) {
OFFLOAD(AS_INDEPENDENT, NUM_THREADS(NTHREADS))
for (std::remove_const_t<decltype(Ni)> i = 0; i < Ni; i++) {
// 初期化（省略）
PRAGMA_ACC_LOOP(ACC_CLAUSE_SEQ)
for (std::remove_const_t<decltype(Nj)> j = 0; j < Nj; j++) {
// ループ内は省略
}
iacc[i] = ai;
}
}
2025/6/26
計算科学技術特論A（2025） 第11回
• i-ループを並列化
• スレッド数はNTHREADS
• j-ループが並列化されると
性能低下の要因となるため，
並列化を抑止
（OpenACC）
実装例（N体計算，OpenACC的記法）
• CPU上で初期条件を生成，GPUに粒子データ転送後に重力計算
set_uniform_sphere(num, pos, vel, Mtot, rad, virial, newton);
PRAGMA_ACC_UPDATE_DEVICE(pos [0:num], vel [0:num])
calc_acc(num, pos, acc, num, pos, eps);
• PRAGMA_ACC_UPDATE_DEVICE() でデータ転送
• 重力計算関数（ほぼ省略版）
void calc_acc(…) {
PRAGMA_ACC_KERNELS_LOOP(ACC_CLAUSE_INDEPENDENT, ACC_CLAUSE_VECTOR_LENGTH(NTHREADS))
for (std::remove_const_t<decltype(Ni)> i = 0; i < Ni; i++) {
// 初期化（省略）
PRAGMA_ACC_LOOP(ACC_CLAUSE_SEQ)
for (std::remove_const_t<decltype(Nj)> j = 0; j < Nj; j++) {
// ループ内は省略
}
iacc[i] = ai;
}
}
2025/6/26
計算科学技術特論A（2025） 第11回
35
実装例（N体計算，OpenMP的記法）
• CPU上で初期条件を生成，GPUに粒子データ転送後に重力計算
set_uniform_sphere(num, pos, vel, Mtot, rad, virial, newton);
PRAGMA_OMP_TARGET_UPDATE_TO(pos [0:num], vel [0:num])
calc_acc(num, pos, acc, num, pos, eps);
• PRAGMA_OMP_TARGET_UPDATE_TO() でデータ転送
• 重力計算関数（ほぼ省略版）
void calc_acc(…) {
PRAGMA_OMP_TARGET_TEAMS_LOOP(OMP_TARGET_CLAUSE_SIMD, OMP_TARGET_CLAUSE_THREAD_LIMIT(NTHREADS))
for (std::remove_const_t<decltype(Ni)> i = 0; i < Ni; i++) {
// 初期化（省略）
PRAGMA_ACC_LOOP(ACC_CLAUSE_SEQ)
for (std::remove_const_t<decltype(Nj)> j = 0; j < Nj; j++) {
// ループ内は省略
}
iacc[i] = ai;
}
}
2025/6/26
計算科学技術特論A（2025） 第11回
36
計算機環境 37 NVIDIA H100 SXM 80GB NVIDIA GH200 480GB AMD Instinct MI210 Intel Data Center GPU Max 1100 66.9 TFlop/s 66.9 TFlop/s 22.6 TFlop/s 22.2 TFlop/s 並列度（FP32） 16896 16896 6656 7168 動作周波数 1980 MHz 1980 MHz 1700 MHz 1550 MHz メモリ容量 HBM3 80GB HBM3 96GB HBM2e 64GB HBM2e 48GB メモリバンド幅 3.36 TB/s 4.02 TB/s 1.64 TB/s 1.23 TB/s TDP/TBP 700 W 1000 W (total) 300 W 300 W ホストCPU Intel Xeon Platinum NVIDIA Grace 8468 AMD EPYC 7713 Intel Xeon Platinum 8468 48 cores × 2 sockets 72 cores 64 cores × 2 sockets 48 cores × 2 sockets 2.1 GHz 3.0 GHz 2.0 GHz 2.1 GHz CUDA 12.3 CUDA 12.4 ROCm 6.0.2 Intel oneAPI 2024.1.0 NVIDIA HPC SDK 24.3-0 NVIDIA HPC SDK AdaptiveCpp 24.5-1 24.02.0 FP32性能 コンパイラ環境 Intel oneAPI 2025/6/262024.1.0 LLVM 18.1.7 計算科学技術特論A（2025） 第11回
N体計算（適切なコンパイルオプションの指定時） 2025/6/26 38
N体計算（-Ofast -gpu=cc90 相当の場合） 2025/6/26 39
実装例（3次元拡散方程式）
• OpenACCコード
をSolomon化
• 元コードは星野さん
（名古屋大）が作った
OpenACCのサン
プルコード
40
init(nx, ny, nz, dx, dy, dz, f);
PRAGMA_ACC_DATA(ACC_CLAUSE_COPY(f [0:n]), ACC_CLAUSE_CREATE(fn [0:n])) {
for (; icnt < nt && time + 0.5 * dt < 0.1; icnt++) {
flop += diffusion3d(nx, ny, nz, dx, dy, dz, dt, kappa, f, fn);
swap(&f, &fn);
time += dt;
}
}
OFFLOAD(AS_INDEPENDENT, COLLAPSE(3), ACC_CLAUSE_PRESENT(f, fn))
for (int i = 0; i < nx; i++) {
for (int j = 0; j < ny; j++) {
for (int k = 0; k < nz; k++) {
const int ix = INDEX(nx, ny, nz, i, j, k);
const int ip = INDEX(nx, ny, nz, IMIN(i + 1, nx - 1), j, k);
const int im = INDEX(nx, ny, nz, IMAX(i - 1, 0), j, k);
const int jp = INDEX(nx, ny, nz, i, IMIN(j + 1, ny - 1), k);
const int jm = INDEX(nx, ny, nz, i, IMAX(j - 1, 0), k);
const int kp = INDEX(nx, ny, nz, i, j, IMIN(k + 1, nz - 1));
const int km = INDEX(nx, ny, nz, i, j, IMAX(k - 1, 0));
fn[ix] = cc * f[ix] + ce * f[ip] + cw * f[im] + cn * f[jp] + cs * f[jm] + ct * f[kp] + cb * f[km];
}}} 2025/6/26
計算科学技術特論A（2025） 第11回
3次元拡散方程式 • NVIDIA GPU上では，OpenMP (distribute) のみ遅い • AMD/Intel GPUs上では，loop と distribute に性能差なし • B/F=2.5 であり，メモリ律速なアプリケーション（キャッシュ律速） • NVIDIA: 4.51 TB/s (H100), 4.58 TB/s (GH200) • AMD MI210: 2.43 TB/s • Intel Data Center GPU Max 1100: 1.82 TB/s 41
小まとめ 42 • GPU向け指示文は複数あり，ユーザーはどれか1つを選択する必要がある （諸悪の根源は各社の政治的な思惑？） • OpenACC: 機能・資料が充実しているが，ほぼNVIDIA GPU向け • OpenMP target: NVIDIA/AMD/Intel 全社に対応，機能はキャッチアップ中 • GPU向け指示文統合マクロSolomonを開発（Miki & Hanawa 2024） • Simple Off-LOading Macros Orchestrating multiple Notations • プリプロセッサマクロ経由で指示文を記載するためのマクロ集 • NVIDIA GPU上ではOpenACCを，AMD/Intel GPU上ではOpenMP target を選択，ということができる • 簡易記法，OpenACC的記法，OpenMP的記法があるため，学習コストを低減 • GPU提供ベンダー製のコンパイラをそのまま使える • OpenACC と OpenMP target の性能比較も簡単にできる • GitHub で公開: https://github.com/ymiki-repo/solomon • 今は C/C++ のみ対応 2025/6/26 計算科学技術特論A（2025） 第11回
Contents 43 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
GPU向けのプログラミング環境 44 2月のPCCC AI/HPC OSS活用WSでの講演資料 （https://www.pccluster.org/ja/event/data/240205_pccc_wsAI-HPC-OSS_06_hanawa-miki.pdf） OpenACC (Cray/HPE compiler) 指示文 OpenACC OpenMP target (NVHPC) OpenMP target (ROCm) OpenMP target (oneAPI) OpenMP target (Cray/HPE compiler) Standard Parallelism Fortran 2018/202X, C++17 標準言語仕様 Std Par., SYCL (AdaptiveCpp) (plugins by Codeplay) 低レベル 2025/6/26 SYCL (DPC++) CUDA HIP NVIDIA GPU AMD GPU 計算科学技術特論A（2025） 第11回 (chipStar) Intel GPU
SYCL環境の構築 46 • Intel oneAPI （コンパイラは icpx） • Codeplay 社提供のプラグインをインストール (for NVIDIA/AMD GPUs) （https://codeplay.com/solutions/oneapi/plugins/） • AdaptiveCpp （コンパイラは acpp） （https://github.com/AdaptiveCpp/AdaptiveCpp） • ドキュメントにしたがってインストールすれば特に苦労なく使えた • Intel GPU 向けのドキュメントが見当たらなかったので，NVIDIA/AMD GPUs 限定の話 • NVIDIA GH200 向けにもインストールできた（Arm CPU なので，oneAPI が使えない） • ただし，nvclangをバックエンドに取ることはできなかった（llvm-tblgenが不足） • いくつかのインストール方法が提示されているが，試したのは以下の手順 1. LLVM を手でインストール（NVPTX or AMDGPU を有効にしておく） • ドキュメントでは LLVM は dnf install することを推奨されていたが，管理者権限なしでインストールで きる方法を取ることにした（スパコンに一般ユーザとしてインストールして使うことを想定した予行演習） 2. AdaptiveCpp のソースを取ってきて，cmake してコンパイル • レジスタを大量に消費する場合に計算がこける場合がある • スレッド数が512以上かつILP数8以上というかなり極端な条件なので，普通は出くわさないはず 2025/6/26 計算科学技術特論A（2025） 第11回
コードの実装手順・勉強手順 47 • CUDA C++版 • 簡単なのでスクラッチから実装 • HIP C++版 • CUDA版のうちいくつか簡単なものを hipify-clang を用いて変換 （シェアードメモリの使い方など，ドキュメントを眺めるよりも簡単に使い方が分かる） • 感触をつかんだ後は，普通にスクラッチから実装できるようになる • SYCL版 • CUDA版のうちいくつか簡単なものを dpct （今はSYCLomatic）を用いて変換 • デフォルトでは各queueが Out-of-Order 実行されるので適宜 wait() をかける • sycl::property::queue::in_order を指定して queue を作成すると， CUDAのdefault streamを使ったときと同じ振る舞い • CUDAに慣れている人にとってはこちらの使い方のほうが馴染みやすいと思われる • 感触をつかんだ後は，普通にスクラッチから実装できるようになる • 注：「簡単」，「普通に」などはあくまでも個人の感想なので保証はしません 2025/6/26 計算科学技術特論A（2025） 第11回
Contents 48 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
CUDA C++で実装する際の記述項目 • GPUの起動（cudaSetDevice()） • これが必要なのは複数GPUを触りに行く場合（多くの場合で省略できる） • 1 GPU / MPI プロセスであれば，CUDA_VISIBLE_DEVICESによる制御が可能 • デバイスメモリの確保（cudaMalloc()） • データ転送用メモリの確保（cudaMallocHost()） • CPU⇔GPU間のデータ転送（cudaMemcpy()） • __global__関数の実装（CPUから起動するGPU関数） • __device__関数の実装（GPU関数から呼び出すGPU関数） • カーネル立ち上げ命令の追加（kernel<<<blck, thrd>>>()） • 確保したメモリの開放（cudaFree(), cudaFreeHost() ） 2025/6/26 計算科学技術特論A（2025） 第11回 49
CUDA C++での実装（お手軽版：1/5）
50
• まずはGPU上で動作させたい関数をGPU化
• 関数定義の先頭に__global__をつける
• GPU上の関数から呼ぶ関数については，__device__をつける
• 一番外側のfor文を削除し，代わりに自動設定されるスレッドIDと紐づけ
• if(ii < Ni){…} をつけないですむように，スレッド数の定数倍のメモリを確保
• 余分に確保したメモリ領域には，質量0の粒子を置くなどの工夫を施す
• （ツリー法などもう一工夫必要な場合もあるが，細かい話なのでここでは省略）
void calc_acc(const type::int_idx Ni, const
type::position *const ipos, …) {
#pragma omp parallel for
for (type::int_idx ii = 0U; ii < Ni; ii++) {
// 関数の中身は省略
}
}
2025/6/26
__global__ void calc_acc_device(const
type::position *const ipos, …) {
const type::int_idx ii = blockIdx.x * blockDim.x
+ threadIdx.x;
// 関数の中身は省略
}
計算科学技術特論A（2025） 第11回
CUDA C++での実装（お手軽版：2/5）
• GPU化した関数をCPUから起動する
• スレッド数，問題サイズから必要なブロック数を設定（マクロ関数が便利）
• スレッド数： NTHREADS
• ブロック数： マクロ関数 BLOCKSIZE を使用
• <<<ブロック数，スレッド数，動的確保するシェアードメモリ容量，ストリーム>>>
• 後ろ2つは省略されることが多い（デフォルト設定をそのまま使用）
constexpr auto BLOCKSIZE(const type::int_idx num, const type::int_idx thread)
{ return (1U + ((num - 1U) / thread)); }
static inline void calc_acc(const type::int_idx Ni, const type::position *const ipos,
type::acceleration *__restrict iacc, const type::int_idx Nj, const type::position
*const jpos, const type::flt_pos eps2) {
calc_acc_device<<<BLOCKSIZE(Ni, NTHREADS), NTHREADS>>>(ipos, iacc, Nj, jpos, eps2);
}
2025/6/26
計算科学技術特論A（2025） 第11回
51
CUDA C++での実装（お手軽版：3/5）
• Managed Memory を使用する場合
• メモリの確保・解放だけ記述すればOK
• GH200でUnified Memoryを使用する際には，malloc/new などでメモリを確保するだけ
• （CPU-GPU間のデータ転送は自分では何もしない）
// 配列サイズを NTHREADS の整数倍にするための細工
auto size = static_cast<size_t>(num);
if ((num % NTHREADS) != 0U) {
size += static_cast<size_t>(NTHREADS - (num % NTHREADS));
}
// マネージドメモリの確保
cudaMallocManaged((void **)pos, size * sizeof(type::position));
// マネージドメモリの解放
cudaFree(pos);
2025/6/26
計算科学技術特論A（2025） 第11回
52
CUDA C++での実装（お手軽版：4/5） • Unified/Managed Memoryを使わない場合は，データ転送も記述 // GPU上のメモリ確保 cudaMalloc((void **)pos_dev, size * sizeof(type::position)); // CPU上の（pinned）メモリ確保（CPU・GPU間のデータ転送高速化のため） cudaMallocHost((void **)pos_hst, size * sizeof(type::position)); // 上記で確保したメモリの解放 cudaFree(pos_dev); cudaFreeHost(pos_hst); // CPU → GPUのデータ転送 cudaMemcpy(pos_dev, pos_hst, num * sizeof(type::position), cudaMemcpyHostToDevice); // GPU → CPUのデータ転送 cudaMemcpy(acc_hst, acc_dev, num * sizeof(type::acceleration), cudaMemcpyDeviceToHost); 2025/6/26 計算科学技術特論A（2025） 第11回 53
CUDA C++での実装（お手軽版：5/5） • もし必要があれば，cudaDeviceSynchronize()を追加 • GPU上で関数を起動すると，関数の終了を待たずにCPUに処理が帰る • 複数のCUDAストリームを使った場合には，どこかで同期が必要 • 性能測定時など，GPU上の関数の状態を把握すべき場合（下の例） • Unified Memoryを使用した際に，CPUから読み出したデータが不正だった場合 auto timer = util::timer(); cudaDeviceSynchronize(); timer.start(); calc_acc(num, pos_dev, acc_dev, num, pos_dev, eps2); cudaDeviceSynchronize(); timer.stop(); 2025/6/26 計算科学技術特論A（2025） 第11回 54
Contents 55 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
CUDAコードをHIP化するにはどうすれば良いか？ 56 • HIPIFYをかけると，CUDAコードをHIP化してくれる • $ hipify-clang *.cu *.cuh --cuda-path=/cuda/path -I/include/path • これだけで大体のことは完了してしまう • 他に hipify-perl というツールもある • hipify-clang の方がよりコンパイラ寄りのツールかつ推奨ツールとのこと • 以前（数年前に）使った際に失敗した例 • 注：今は挙動が変わっている可能性があるので，まずは試してみてください • #ifdef … #else … #endif のようなマクロスイッチがある時 • #ifdef などを無視して変換するようで，const を付けた同名変数を複数回定義している，などと いうエラーが出たりする • 変換後のコードで有効にするフラグを切り替えることを想定すると仕方ない気もする • CUDA版ライブラリとROCm版ライブラリで不整合がある場合 • cuRAND → rocRAND で，ヘッダの名前やライブラリのマクロ名などが違っていた （主に引っかかったのは Mersenne Twister まわり） • （AMD GPUで動かすには）変換後のコードを自分で手直しする必要があった 2025/6/26 計算科学技術特論A（2025） 第11回
とても簡単なコードであれば，区別がつかないことも
• N体計算の重力計算部分を実装した例（あまり最適化はしていない）
__global__ void calc_acc(const int Ni, float4 *ipos, float4 *iacc, const int Nj, float4
*jpos, const float eps){
const int ii = blockIdx.x * blockDim.x + threadIdx.x;
float4 pi = ipos[ii]; pi.w = eps * eps;
float4 ai = {0.0F, 0.0F, 0.0F, 0.0F};
for(int jj = 0; jj < Nj; jj++){
const float4 pj = jpos[jj];
float4 rji;
rji.x = pj.x - pi.x;
rji.y = pj.y - pi.y;
rji.z = pj.z - pi.z;
const float r2 = fmaf(rji.z, rji.z, fmaf(rji.y, rji.y, fmaf(rji.x, rji.x, pi.w)));
rji.w = 1.0F / sqrtf(r2);
rji.w *= rji.w * rji.w;
rji.w *= pj.w;
ai.x = fmaf(rji.x, rji.w, ai.x);
ai.y = fmaf(rji.y, rji.w, ai.y);
ai.z = fmaf(rji.z, rji.w, ai.z);
}
iacc[ii] = ai;
}
2025/6/26
計算科学技術特論A（2025） 第11回
57
HIPとCUDAでの実装の違いは何か？ • 大雑把な違い： • ホストで呼ぶ cuda* という関数・変数には，hip* が対応 • 新しいCUDA関数については，対応するHIP版がないこともある • cudaMallocHost ➔ hipHostMalloc のような例外もある（HIPIFYに任せればOK） • GPU上で動かすカーネル関数の起動方法 • func<<<blck, thrd, shmem, stream>>>(var0, var1, …); • shmem, streamは省略することも多い（デフォルトで0が入る） • hipLaunchKernelGGL(func, blck, thrd, shmem, stream, var0, var1, …); • shmem, stream は省略できない（HIPIFY すると自動的に入れてくれていた） • （ドキュメントによると，CUDAと同じ記法でも良いらしい） • 特にデバイス関数の中身については，HIPとCUDAはほぼ同じ（前ページ） • ハードウェア側の違いを吸収しておく方が重要 • NVIDIA GPUでは，warpSize = 32 を基本単位として考える • （CDNA系の）AMD GPUでは，waveSize = 64 を基本単位として考える • AMD GPUでも，RDNA系であれば waveSize = 32, 64 で切り替え可能 2025/6/26 計算科学技術特論A（2025） 第11回 58
CUDAコードをSYCL化するにはどうすれば良いか？ 59 • SYCLomatic（旧dpct） • コマンドは c2s • （試したのがかなり昔なので，$ c2s --help で使い方を表示して試してください） • 変換前のコードがコメントアウトされた状態で残されており， また（必要に応じて）追加メッセージが書き込まれていることも • CUDAとSYCLを見比べやすいという意味で親切な設計と言える • CUDAとHIPはよく似ていたが，SYCLはそれなりに雰囲気が異なる • C++のラムダ式を活用した実装になる • SYCLはむしろKokkosに近い（差分もかなりあるので，似ているとは言い難いが） • （HIPIFYも同様だが）自分が中身を理解しているコードを別言語で実装し直 したコードが出力される ➔自分向けにカスタマイズされたサンプルコードを作ってくれるツール • ドキュメントだけを眺めているよりも短時間でHIPやSYCLを学習できる 2025/6/26 計算科学技術特論A（2025） 第11回
先程のN体コードに似せて実装したSYCLコード
#include <sycl/sycl.hpp>
static constexpr auto BLOCKSIZE = [](const auto num, const auto num_threads) {
return (1 + ((num - 1) / num_threads));
};
static constexpr auto KERNEL_PARAM_1D = [](const size_t num, const size_t num_threads) {
const auto grid = sycl::range<1>{BLOCKSIZE(num, num_threads)};
const auto block = sycl::range<1>{num_threads};
return (sycl::nd_range<1>{grid * block, block});
};
void calc_acc_kernel(sycl::nd_item<1> nd, const int Ni, float4 *ipos, float4 *iacc, const int Nj, float4 *jpos, const float eps) {
const int ii = nd.get_global_id(0);
float4 pi = ipos[ii]; pi.w = eps * eps;
float4 ai = {0.0F, 0.0F, 0.0F, 0.0F};
for (int jj = 0; jj < Nj; jj++) {
const float4 pj = jpos[jj];
float4 rji;
rji.x = pj.x - pi.x;
rji.y = pj.y - pi.y;
rji.z = pj.z - pi.z;
const float r2 = sycl::fma(rji.z, rji.z, sycl::fma(rji.y, rji.y, sycl::fma(rji.x, rji.x, pi.w)));
rji.w = 1.0F / sycl::sqrt(r2);
rji.w *= rji.w * rji.w;
rji.w *= pj.w;
ai.x = sycl::fma(rji.x, rji.w, ai.x);
ai.y = sycl::fma(rji.y, rji.w, ai.y);
ai.z = sycl::fma(rji.z, rji.w, ai.z);
}
iacc[ii] = ai;
}
void calc_acc(sycl::queue &queue, const int Ni, float4 *ipos, float4 *iacc, const int Nj, float4 *jpos, const float eps) {
queue.parallel_for(KERNEL_PARAM_1D(Ni, NTHREADS), [=](sycl::nd_item<1> item) {
calc_acc_kernel(item, Ni, ipos, iacc, Nj, jpos, eps);
});
}
2025/6/26
計算科学技術特論A（2025） 第11回
60
Contents 61 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
チューニング・性能評価環境 62 GPU NVIDIA H100 SXM 80GB NVIDIA GH200 480GB AMD Instinct MI210 Intel Data Center GPU Max 1100 FP32 peak 66.9 TFlop/s 66.9 TFlop/s 22.6 TFlop/s 22.2 TFlop/s # of units 132 SMs 132 SMs 104 CUs 448 EUs FP32 parallelism 16896 16896 6656 7168 Clock frequency 1980 MHz 1980 MHz 1700 MHz 1550 MHz TDP 700 W （全体で）1000 W 300 W 300 W Host CPU Intel Xeon Platinum 8468 NVIDIA Grace AMD EPYC 7713 Intel Xeon Platinum 8468 48 cores × 2 sockets 72 cores 64 cores × 2 sockets 48 cores × 2 sockets 2.1 GHz 3.0 GHz 2.0 GHz 2.1 GHz CUDA 12.3 CUDA 12.4 ROCm 6.0.2, 5.4.3 Intel oneAPI 2024.1.0 LLVM 18.1.7 18.1.7 18.1.7 AdaptiveCpp 24.02.0 24.02.0 24.02.0 2025/6/26 2024.1.0 計算科学技術特論A（2025） 第11回 2024.1.0
最適な逆数平方根演算命令の特定 63 • N=4Mにおいて，1秒あたりに計算できた相互作用ペア数を比較 GPU Compiler 1.0F/std::sqrt() rsqrtf() __frsqrt_rn() sycl::rsqrt() sycl::native:: rsqrt() NVIDIA H100 nvcc 8.06 × 1011 NVIDIA H100 icpx 1.15 × 1012 1.77 × 1012 1.77 × 1012 NVIDIA H100 acpp 7.90 × 1011 8.08 × 1011 8.08 × 1011 NVIDIA GH200 nvcc 8.15 × 1011 NVIDIA GH200 acpp 8.01 × 1011 8.08 × 1011 8.08 × 1011 AMD MI210 hipcc 2.36 × 1011 AMD MI210 icpx 7.13 × 1011 7.13 × 1011 7.13 × 1011 AMD MI210 acpp 2.35 × 1011 7.06 × 1011 7.06 × 1011 Intel 1100 icpx 5.52 × 1011 5.52 × 1011 5.52 × 1011 2025/6/26 1.51 × 1012 1.51 × 1012 5.28 × 1011 7.68 × 1011 7.68 × 1011 7.05 × 1011 計算科学技術特論A（2025） 第11回
NVIDIA製GPU向けの実装 64 • CUDA, SYCL 実装を測定 • HIP版はCUDA版と同性能が得られると分かっているので，今回は割愛 （管理者権限なしで（= rpmパッケージを使わずに）HIP環境を構築するのは面倒） • 逆数平方根，ブロックあたりのスレッド数，ループアンローリング段数， Instruction Level Parallelism（ILP）数については，パラメータ探査の 結果最高性能だったものを採用 • シェアードメモリの使い方（single or double buffer，ブロック単位 or ワープ単 位）やmemcpy_async()の使用についても同様 • ただしmemcpy_async()の使用はCUDA版のみで実装 • 非正規化数についてFTZ（flush-to-zero）を有効化するかどうかも比較 • パラメータ探査時の粒子数はN=4M • 逆数平方根はrsqrtf()を使用 （acppでは__hipsycl_if_target_cuda()経由， icpxではsycl::rsqrt()） • CUDA: memcpy_async()を使ったほうが速くなった（FTZ有効時限定） • 最適なパラメータはCUDA, SYCL (icpx), SYCL (acpp) それぞれで異なる 2025/6/26 計算科学技術特論A（2025） 第11回
AMD製GPU向けの実装 65 • HIP, SYCL 実装を測定 • 逆数平方根，ブロックあたりのスレッド数，ループアンローリング段数， Instruction Level Parallelism（ILP）数については，パラメータ探査の 結果最高性能だったものを採用 • シェアードメモリの使い方（single or double buffer，ブロック単位 or ウェーブ単 位）の使用についても同様 • Packed FP32命令（FP32の加算・乗算・積和算の性能が2倍）の使用・不使用 • CDNA 2世代のGPUで導入された • 非正規化数についてFTZ（flush-to-zero）を有効化するかどうかも比較 • パラメータ探査時の粒子数はN=4M • 逆数平方根は__frsqrt_rn()を使用 （acppでは__hipsycl_if_target_hip()経由， icpxではsycl::rsqrt()） • スレッド数は256，シェアードメモリは不使用 2025/6/26 計算科学技術特論A（2025） 第11回
Intel製GPU向けの実装 66 • SYCL 実装を測定，SYCL も Intel oneAPI のみ • 逆数平方根，ブロックあたりのスレッド数，ループアンローリング段数， Instruction Level Parallelism（ILP）数については，パラメータ探査の 結果最高性能だったものを採用 • シェアードメモリの使い方（single or double buffer，ブロック単位 or ワープ単 位）の使用についても同様 • 非正規化数についてFTZ（flush-to-zero）を有効化するかどうかも比較 • パラメータ探査時の粒子数はN=4M • 逆数平方根はsycl::rsqrt()を使用 • スレッド数は1024 • シェアードメモリを使用（single buffer） • ループアンローリングは2段 • ILPは使用せず（= 通常の実装） 2025/6/26 計算科学技術特論A（2025） 第11回
GPU状態の監視 67 • GPUの温度，動作周波数，消費電力を 計算中に取得 • GPUの温度が下がった状態から測定するため， 毎回の測定前に10分間スリープ • NVIDIA GPU: NVML • AMD GPU: ROCm SMI • Intel GPU: Level-Zero Sysman API • 管理者権限が要求されるメトリックあり • 電力測定方法はNVIDIA，AMDと少し違う • 本計算に用いていない余剰コアを用い， 0.1秒間隔で状態を取得 • OpenMPのタスク指示文を活用 • 本計算側のコードについては一切変更なし 2025/6/26 static bool repeat; repeat = true; #pragma omp parallel num_threads(2) { #pragma omp single { #pragma omp task { while (repeat) { observe_GPU(); sleep(); } } #pragma omp task { run_simulation(); repeat = false; } #pragma omp taskwait } } 計算科学技術特論A（2025） 第11回
NVIDIA/AMD/Intel製GPUの性能比較 2025/6/26 68
GPUの状態 • GPUの温度，消 費電力，動作周 波数を監視 • GPU温度上昇に よる動作周波数 低下はなかった • AMD MI210 では，供給電力 不足による動作 周波数低下 （packed FP32命令の使 用頻度に依存） 2025/6/26 69
各世代の最上位製品どうしの比較 70 • AMD MI210, Intel Data Center GPU Max 1100 は最上位製品で はないので，最上位製品を用いた際の性能を推定 • 理論ピーク性能に基づく推定：実際には供給電力が不足するので上限値（下矢印つき） • AMD MI250X: MI210の2.12倍の理論ピーク性能，1.87倍のTDP（注：MI210でも不足） • Intel Data Center GPU Max 1550: 1100の2.36倍の理論ピーク性能，2倍のTDP • 消費電力に基づく推定：電力性能が変わらなければこちらの方が正解に近いはず • 電力性能は動作周波数に依存するため，（動作周波数が下がると）推定値より上がることもある 2025/6/26
Kokkos との性能比較 • 欧米では性能可搬フレームワーク Kokkosがよく使われている • 来週，似鳥さんが解説してくれます • CMakeの使い方に一癖あり， （CMakeにはそれなりに慣れてい たが）CMakeまわりの調整に一番 時間がかかった • Miyabi-G 上での測定 • NVIDIA GH200 120GB • CUDA 12.6 • AdaptiveCpp 24.10.0 • LLVM 19.1.7 • Kokkos 4.5.99 • nvcc_wrapper は手直しした • （なぜか）SYCL (acpp) が 一番速い • Kokkos も良い勝負 2025/6/26 計算科学技術特論A（2025） 第11回 71
小まとめ 72 • CUDA/HIP/SYCL を用いてN体計算コードを実装・最適化し， NVIDIA/AMD/Intel製GPU上での性能を比較した • Kokkos についてもMiyabi-Gで実験開始．今のところ良好な成功を確認 • SYCL実装の性能が良好 • 全ベンダー製GPUに対応でき，なおかつきちんと性能を出せる • Intel oneAPI と AdaptiveCpp を使い分けられることも大きい • NVIDIA GH200 上では AdaptiveCpp が使える • NVIDIA H100, Miyabi-G 上では CUDA よりも速かった（!!） • AMD MI210 上では HIP とほぼ同性能 • NVIDIA Hopper, AMD CDNA 2, Intel PVC 世代においては NVIDIA H100 (SYCL, icpx) が最高性能となった • 電力性能についても同様（CPU，冷却などの寄与は入っていない点に注意） • 使用電力は600 Wを越えており，こちらも最大 （電力性能が一番高く，かつ供給電力も最大なので，性能も最高値となる） • 次（or 最新）の世代ではどうなる？（AMD MI300, NVIDIA B200, …） 2025/6/26 計算科学技術特論A（2025） 第11回
Contents 73 • GPUとは？ なぜGPU搭載型スパコンが増えるのか？ • GPU向けの多様な開発手法とベンダー依存性 • 指示文ベースのGPU化手法 • OpenACCやOpenMP targetによる実装方針 • OpenACC/OpenMP target統合マクロSolomon • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • CUDA/HIP/SYCLを用いたGPU化 • CUDAでの実装概要 • CUDA実装からHIPやSYCLへの変換方法（あるいは教科書作成方法） • N体計算（直接法）を例題としたNVIDIA/AMD/Intel GPUs上での性能比較 • （もし時間があれば）Miyabi@JCAHPCの紹介 2025/6/26 計算科学技術特論A（2025） 第11回
JCAHPC第二世代システム Miyabi 74 • JCAHPC: 最先端共同HPC基盤施設（2013年～） • 筑波大学計算科学研究センターと東京大学情報基盤センターが共同で調達・運用 • 導入の経過 • 2022年11月より調達プロセスを開始 • 2022年6月には事前性能評価によりGPUアーキテクチャを決定 • 2023年11月に開札，富士通が落札 • 準備期間を1年以上確保 • システムの特徴 • システム全体性能の飛躍的な向上のため，演算加速装置としてGPUを主体とするシス テムへ • 消費電力も削減➔電力あたり性能の劇的な向上 • CPU-GPU間が高速リンクで密結合され，既存アプリのGPU化が容易に • GPU化が困難なアプリケーションのために，汎用CPUのみの計算ノードも導入 • ストレージの高性能化に向けてAll Flashを導入 2025/6/26 計算科学技術特論A（2025） 第11回
Miyabiの外観 2025/6/26 75 計算科学技術特論A（2025） 第11回
Miyabi (OFP-II) (1/2) 76 • Miyabi-G: CPU+GPU: NVIDIA GH200 • Node: NVIDIA GH200 Grace-Hopper Superchip • Grace: 72c, 3.456 TF, 120 GB, 512 GB/sec (LPDDR5X) • H100: 66.9 TF DP-Tensor Core, 96 GB, 4,022 GB/sec (HBM3) • Cache Coherent between CPU-GPU • NVMe SSD for each GPU: 1.9TB, 8.0GB/sec, GPUDirect Storage • Total (Aggregated Performance: CPU+GPU) N VIDIA GH200 Grace Hopper Superchip LPDDR5X 120 GB • 1,120 nodes, 78.8 PF, 5.07 PB/sec, IB-NDR 200 • Miyabi-C: CPU Only: Intel Xeon Max 9480 (SPR) • Node: Intel Xeon Max 9480 (1.9 GHz, 56c) x 2 • 6.8 TF, 128 GiB, 3,200 GB/sec (HBM2e only) • Total • 190 nodes, 1.3 PF, IB-NDR 200 • 372 TB/sec for STREAM Triad (Peak: 608 TB/sec) HBM 3 96 GB 512 GB/s GRACE CPU 72c, 2.6 GHz PCIe Gen4 x4 NVM e SSD 1.92 TB 4.022 TB/s N VLink C2C Hopper GPU 450 GB/s PCIe Gen5 x8 IB NDR HCA ConnectX-7 IB N DR200 (200 Gbps)
Miyabi (OFP-II) (2/2) 77 • ファイルシステム: DDN EXA Scalar, Lustre FS • 11.3 PB (NVMe SSD) 1.0TB/sec, “Ipomoea-01” (26 PB) も利用可能 • Miyabi-G/C の全ノードはフルバイセクションバンド幅で接続 • (400Gbps/8)×(32×20＋16×1) = 32.8 TB/sec • 2025年1月運用開始 • Miyabi-G/C間の通信はh3-Open-SYS/WaitIO により実現 IB-NDR（400Gbps） IB-HDR（200） IB-NDR200（200） Miyabi-G NVIDIA GH200 1,120 78.8 PF, 5.07 PB/sec 2025/6/26 Miyabi-C Intel Xeon Max (HBM2e) 2 x 190 1.3 PF, 608 TB/sec File System DDN EXA Scaler 11.3 PB, 1.0TB/sec 計算科学技術特論A（2025） 第11回 Ipomoea-01 Common Shared Storage 26 PB
NVIDIA GH200の特性 • Grace-Hopper相互のメモリ空間を直接参照可能，NUMA的な扱い • CPU-GPU間： コヒーレントインタフェース（NVLink-C2C） • PCIe Gen 5の7倍以上の帯域（450GB/sec/dir） • CPU・GPUの効率的使い分けも可能 • 従来はデータ転送がボトルネック • プログラミングも用意 • AMD MI300Aも同じ方向性 • 小規模問題，GPUが不得意な計算をCPUが柔軟に処理することも可能 2025/6/26 計算科学技術特論A（2025） 第11回 78
Graceからのメモリビュー • NUMAとして見える • malloc()ではFirst Touchが大事 NUMA Domain 0 79 • 従来のCUDAのメモリモデルも使えて 普通に動く（コード改変不要）＋転送が速い • cudaMalloc() + cudaMemcpy() • cudaMalloc()した領域はGraceからアクセス不可 NUMA Domain 1 NVIDIA GH200 Grace Hopper Superchip GRACE CPU LPDDR5X 120 GB 72c, 2.6 GHz Hopper GPU NVLink C2C HBM3 96 GB 450 GB/s 512 GB/s 4.02 TB/s PCIe Gen4 x4 NVMe SSD 1.92 TB 2025/6/26 PCIe Gen5 x8 IB NDR HCA ConnectX-7 IB NDR200 (200 Gbps)
Hopperからのメモリビュー 80 • Graceのアドレスを使って直接アクセス可能 • 従来のCUDAのコードはH100とほぼ挙動が変わらない （メモリバンド幅比相当の性能向上） NVIDIA GH200 Grace Hopper Superchip GRACE CPU LPDDR5X 120 GB 72c, 2.6 GHz Hopper GPU NVLink C2C HBM3 96 GB 450 GB/s 512 GB/s 4.02 TB/s PCIe Gen4 x4 NVMe SSD 1.92 TB 2025/6/26 PCIe Gen5 x8 IB NDR HCA ConnectX-7 IB NDR200 (200 Gbps)
NVIDIA GH200のメモリアクセス 81 • CUDA メモリ メモリモード 確保される場所 Access-based Migration CPUからアクセス GPUからアクセス System-allocated (malloc, new) Unified相当 First-touch (GPU または CPU) ◯ ◯ ◯ CUDA managed (cudaMallocManaged) Managed相当 First-touch (GPU または CPU) ◯ ◯ ◯ CUDA device memory (cudaMalloc) Separate相当 (Device) GPU CUDA host memory (cudaMallocHost) CPU ◯ ◯ ◯ • OpenACC, OpenMP target, stdpar メモリモード コンパイルフラグ デフォルトとなる環境 Separate -gpu=mem:separate OpenACC OpenMP target GPU上のデータはGPUからのみアクセス可能 GPU-CPU間の明示的なデータ移動が必要 Managed -gpu=mem:managed stdpar (Managed Memory のみの環境） 動的メモリ確保されたデータはGPU, CPU どちらから もアクセス可能 Unified -gpu=mem:unified stdpar (Unified Memory 全てのデータはGPU, CPU どちらからもアクセス可能 計算科学技術特論A（2025） 第11回 対応の環境） 2025/6/26
GPU移行プラットフォームとしてのMiyabi 82 • 日本国内でも，GPU搭載スパコンが続々と増えている • TSUBAME4.0@科学大，玄界@九大，Miyabi@JCAHPC，ABCI 3.0@AIST，⋯ • 科学技術計算用のGPUスパコンとしてはMiyabi-Gが国内最大のシステム • GPU初心者にとって移行が一番簡単なシステムはMiyabi-G • GH200では，GPU初心者がはまりやすい罠が大幅に軽減されている • CPUからもGPUからも相互にメモリ空間が参照できる • CPU-GPU間のデータ転送が（x86-Hopperに比べて）性能ボトルネックになりづらい • コードを部分的にGPU移植しながら動作テストするのも比較的容易 • CPU-GPU間のデータ転送コストが（通常のPCIe接続に比べて）大幅に軽減されている （コードを移植途中に「遅くなった」と思って熱が冷めるリスクが減る） • GPU移植しづらい部分をCPUに置き去りにしても，性能面で問題になりづらい • 東大情報基盤センターとしてGPUを主体とするシステムはMiyabiが初 • したがって，ユーザコードの移植支援にも力を入れている状態 • GPU関係の講習会，GPU移行相談会，ポータルサイトでの情報提供など • ユーザアカウントを持っていなくてもOK，すべて無料 2025/6/26 計算科学技術特論A（2025） 第11回
東大スパコンの利用制度（https://www.cc.u-tokyo.ac.jp/guide） • 一般利用 • 大学・公共機関に在籍の方（大学院生は代表者としては申し込めません） • 電気代相当料金の利用負担金支払いが必要 • 企業利用 • 企業に在籍の方 • 書面・ヒアリング審査あり • 成果公開型： 利用成果報告書を公開，利用負担金は一般利用の約1.2倍 • 成果非公開型： 報告書・テーマ・社名など非公開，利用負担金は一般利用の約4倍 • 若手・女性利用 • 大学・公共機関に在籍の方 • 4月1日現在40歳以下の若手，または女性，または学生 • 利用負担金なし • 書類審査あり，成果報告義務あり • 学際大規模情報基盤共同利用・共同研究拠点（JHPCN）への課題申請 • HPCI課題への課題申請 2025/6/26 計算科学技術特論A（2025） 第11回 83
最後に 84 • （今後も）GPU搭載型スパコンが増えていくと考えられる • 汎用CPUよりも消費電力あたりの性能が圧倒的に高い • TSUBAME4.0, 玄界, Miyabi, Grand Chariot 2, QST/NIFSの新システム, … • GPU向けのプログラミング手法は数多くあるので，自分にあった手法を選択 • コードの書き換えコストはどの程度まで許容できる？ • 達成したい性能はどの程度？ • ベンダーロックインをどこまで気にするか？ • コードの寿命は最低どのぐらい？ • 指示文の場合: OpenACC or OpenMP target を選択する必要あり ➔ Solomon を使ってインターフェースを統一することもできる • GPU提供ベンダーが直接サポートする開発環境: CUDA, HIP, SYCL ➔ Direct N-body: SYCLはNVIDIA/AMD/Intel製GPU上で高性能 • 今日は紹介しなかったが，世の中的にはKokkosを使った実装も（主に欧米） • 来週，似鳥さん@理研R-CCSが紹介してくれます