GDEP Solutions, Inc.
  • Home
    • 新着情報一覧
    • NVIDIA NEWS
    • GPU2021 開催概要 >
      • Day1 AI & GPU セッション
      • Day2 イメージングAI
      • Day3 GPUスパコン
      • 講演レポート | Denso
      • 過去開催 >
        • GPU2020
        • GPU2019
  • GPU製品
    • 即納モデル
    • NVIDIA GPU 一覧 >
      • NVIDIA H100
      • NVIDIA A100
      • NVIDIA RTX スペック比較 >
        • NVIDIA RTX A6000
        • NVIDIA RTX A5000
        • NVIDIA RTX A4500
        • NVIDIA RTX A4000
        • NVIDIA RTX A2000
      • Quadro GV100
    • GPUレンタル一覧
    • GPU年定額プラン
    • AI・データサイエンスおすすめ一覧 >
      • NVIDIA DGX H100
      • NVIDIA DGX A100
      • DGX STATION A100
      • HP Z8 G4 Workstation
      • HP ZBook Fury17G7 Mobile Workstation
      • Dell Precision 7920 Tower
      • DeepLearning BOXⅡ
      • DeepLearning STATION
      • NVIDIA EGX サーバー
      • HITACHI SR24000
    • 数値計算・解析おすすめ一覧 >
      • HP Z4 G4 Workstation
      • HP Z8 G4 解析ソフトウェア動作確認済みモデル
      • Supermicro 7049GP
    • NVIDIA A100搭載 Supermicro 740GP
  • ストレージ
    • Synology
    • DDN STORAGE
    • PURESTORAGE
  • HPC
    • HPC Workstation
    • HPCおすすめGPUサーバー
    • HPC SIサービス
    • AXXE-L by XTREME-D
    • NVIDIA HPC SDK
    • プログラム高速化サービス
  • クラウド製品
    • セキュリティ >
      • KernelCare
    • 仮想化 >
      • Login VSI
      • Login PI
    • リモートアクセス >
      • FastX
      • NiceDCV
      • NoMachine
    • ハイブリッドクラウドNAS >
      • Morro Data
    • クラウドストレージ >
      • クラウドストレージ Wasabi
  • GPUコラム / 導入事例
    • GPUプログラミング >
      • 初級編
      • 中級編 >
        • 第4回:MPI+OpenACC実装における計算と通信のオーバーラップ
        • 第3回:拡散現象シミュレーションのおさらい
        • 第2回:簡単なOpenACC + MPI コードで考える
        • 第1回:複数のGPUを使う方法とは?
    • GPU Technology for CG/AI >
      • 深層学習を利用した画像処理・必要なGPU性能
      • トランスフォーマー 最近流行のニューラルネットワーク
      • GAN Inversion による写実的画像生成の制御
      • 深層学習におけるアノテーションコストを抑えるための取り組み Active Learning
      • 深層学習に基づく人物画像の再照明
      • GPUの起源と進化
      • AlphaGo とその後
      • CUDAを用いたシンプルなパストレーシング
      • 流体シミュレーションの応用
      • GPUを用いた高速レンダリング
      • GPUを基盤としたCG/AIの技術進化
    • シリコンバレー発信 New Technology Report >
      • 最新記事から
    • ツブ子が聞く・見る・行く! >
      • 見る!NVIDIA RTX A6000
      • 聞く!NVIDIA DGX A100
      • 見る!NVIDIA A100 Tensor Core GPU
    • 導入事例 >
      • DGX SYSTEMS >
        • NVIDIA DGX A100 | 金沢大学
      • GPU Computing
  • 会社情報
    • ごあいさつ
    • 会社概要
    • アクセスマップ
    • 採用情報
  • お問い合わせ

OpenACCではじめるGPUプログラミング

中級編 第2回

画像

第2回:簡単なOpenACC + MPI コードで考える

9/30/2021

 
中級編の続きです。
​前回はMPI + OpenMPによるHello Worldの解説でしたが、今回はもう少し計算っぽいことをしてみましょう。

​コードは全てgithub上で公開されています。https://github.com/hoshino-UTokyo/lecture_openacc_mpi.git

​簡単なOpenACC + MPI コードで考える

今回はopenacc_mpi_basic/以下のコードを使います。解説はCで行いますが、Fortranも用意されています。
画像
図1: lecture_openacc_mpi/C/openacc_mpi_basic/02_kernelsのコードの一部。
OpenACCによりループを並列化、MPIにより配列の一部を転送している。
画像
​図2: 図1の計算内容を図示したもの。
Rank1のプロセスが、Rank0に一部分だけ計算内容を送っている。
図1は今回の計算内容、図2はその計算内容を図示したものです。
このプログラムは、


  1. Rank 1 が配列aをa = 3.0*ny で初期化。
  2. Rank 1 の配列aの一部分(w * nxの部分行列)をRank 0 のbにコピー。
  3. コピーが成功していれば、Rank 0 の配列bの総和 sum = (3.0*ny) * (w*nx)となる。
  4. 最終的にはRank 0 の配列bの平均値を出力する。nx = 1,000, ny = 1,000, w = 10で初期化されているので、結局3.0*w = 30.0が出力されればOK。

という動きをします。

特に意味のある計算ではありませんが、このプログラムの高速化について考えてみましょう。
まず、OpenACCプログラムの高速化の基本として、CPU-GPU間のデータ転送を削減するために、data指示文を付ける必要があるのでした。

​慣れてきた皆様であれば、図1の46行目あたりに#pragma acc data create(a[0:n],b[0:n])と挿入して、68行目までをdata指示文で囲ってしまいたくなると思います。

​しかしその場合、真ん中のMPI関数は期待通り動くのでしょうか?

GPUにあるデータをMPIで送るためには?

実はこのプログラム、単にdata指示文で囲っただけでは期待通りに動きません!

#pragma acc data create(a[0:n],b[0:n]) という指示文で作った配列aとbは、実際にはa_cpu, a_gpuとb_cpu, b_gpuというペアになっているのでした。

問題はこのペアのうちどちらが使われるのかということです。

kernels指示文で囲まれた領域内では、ペアのうちGPU側が使われるのですが、kernels指示文で囲まれていないMPI関数内ではペアのうちCPU側が使われるのです。

​なお、kernels指示文の内側でMPI関数を呼ぶことはできません。それではどうしましょうか?方法は3つあります。

1. update指示文を使う

​update指示文を使って、配列ペアのCPU側を更新し、CPU同士で通信する方法です。
通信が終わったら、配列ペアのGPU側を更新します。


図3のコード例では、Rank 1はMPI_Sendをする前に64行目のupdate指示文で配列aのCPU側を更新。
​Rank 0 は、MPI_Recvの後に62行目のupdate指示文でGPU側を更新しています。転送経路を図示すると、図4の様になります。
画像
図3: lecture_openacc_mpi/C/openacc_mpi_basic/03_updateのコードの一部。
この外側でdata指示文を用いている。
画像
図4: CPUのメモリを経由して別のノードにMPI通信する際の転送経路

2. CUDA aware MPIとhost_data指示文を使う(推奨)

MPIのライブラリには、CUDA aware MPIという機能を持つものがあります。
これは、MPI関数の引数にGPUのアドレスを受け付けるというものです。
GPU側の配列をMPI関数に受け渡すことができれば、期待通りに通信してくれます。


図5のコード例では、66, 69行目でhost_data指示文を使っています。
host_data指示文が適用された範囲内では、use_device()内に指定された配列について、GPU側を使ってくださいという指示文です。

図5の様に関数の上に書けばその関数内で、{}で囲めばその範囲がhost_data指示文の適用範囲となります。
図5では、MPI_Send, MPI_Recvの第一引数にそれぞれGPU側のアドレスが渡ります。

このCUDA aware MPIを使うときに限り、GPU Direct RDMAという通信方式を有効化することができます。
​有効化した場合、図6のような経路でデータ転送が行われます。


​MPIの構築が難しいという点を除けば、実装も簡単でかつCPUに戻さない分速いので、この方法が推奨です。
画像
図5: lecture_openacc_mpi/C/openacc_mpi_basic/04_cuda_awareのコードの一部。
この外側でdata指示文を用いている。
画像
図6: CPUのメモリを経由せず、別のノードにMPI通信する際の転送経路

3. Unified memory 機能に任せる

最近のGPUには、Unified memoryという、CPUとGPUの間の転送を勝手にやってくれるという機能がついているのでした。
Nvidiaコンパイラ(PGIコンパイラから名前が変わりました。機能的にはあまり変更はなさそうです。)では、コンパイルオプションに-ta=tesla,
managedと付けることにより、Unified memoryを有効化できます。

以前に説明した通り、Unified memory機能を有効化すると、data指示文によるデータ転送を書かなくても、必要になったタイミングでデータ転送を行ってくれます。

ですので、図1の実装からコンパイルオプションを変更するだけでOKという点でお手軽です。

​ただし、Unified memory機能を有効にした場合、CPU側の配列とGPU側の配列という区別がなくなってしまうので、MPI関数を呼んだ場合にはCPUへのデータ転送が発生し、図4のような転送経路を辿ることになります。
​

それでは、図1, 3, 5の実装を、Wisteria/BDEC-01のGPUノードで動かしてみましょう。
画像
きちんと速くなってますね。

< 1ヵ月間有効のスパコンお試しアカウント >
​

東京大学情報基盤センターでは、教育の一環として、制限はあるものの一ヵ月の間有効なスパコンアカウントを提供しています。
現在3つのスパコンが運用されていますが、そのうちReedbushと呼ばれるスパコンには、一世代前のものではありますがGPUが搭載されていて、OpenACCを使える環境も整っています。
自分でどんどん自習したい場合は、ご利用を考えてみてください。

トライアルアカウント申し込みページ
https://www.cc.u-tokyo.ac.jp/guide/trial/free_trial.php

< 過去の講習会の資料やプログラム公開中 >
東大センターが行った過去のOpenACCに関する講習会の資料やプログラムも公開されていますので、自習する場合にはぜひご利用ください。

講習会ページ ※オンライン講習会 定期開催中!
https://www.cc.u-tokyo.ac.jp/events/lectures/

講習会で用いているプログラム
https://www.dropbox.com/s/z4fmc4ibdggdi0y/openacc_samples.tar.gz?dl=0​

コメントはクローズされています。

    著者

    東京大学
    助教 星野 哲也 先生
    <略歴>
    ​​2016年~現在:東京大学
    情報基盤センター
    スーパーコンピューティング研究部門 助教
    2018年:東京工業大学
    ​大学院情報理工学研究科 数理・計算科学専攻 博士課程修了 博士(理学)

    アーカイブ
    中級編
    第4回(2022年2月)
    第3回(2021年12月)
    ​第2回(2021年9月)
    ​第1回(2021年7月)
    ​初級編 もくじ ≫
    ​中級編 もくじ ≫

    RSS フィード

Picture
GDEPソリューションズ株式会社
東京都文京区本郷三丁目34番3号 本郷第一ビル8階
TEL:03-5802-7050
・NVIDIA認定 Elite Partner [最上位レベル]
・NVIDIA Advanced Technology Program 達成  [DGX 販売資格]
・東京都公安委員会 古物営業許可番号
 第305471905562号
≫ 新着情報
≫ GPU2021開催概要
​
GPU製品
​≫ 即納モデル
≫ NVIDIA GPU一覧
 ≫NVDIA RTX スペック比較

≫ GPUレンタル一覧
≫ GPU年定額プラン
≫ AI・データサイエンスおすすめ
≫ 数値計算・解析おすすめ一覧
≫ NVIDIA A100搭載 SM740GP


​ストレージ
​≫ Synology
≫ DDN STORAGE
≫ PURESTORAGE
​
HPC
≫ HPC Workstation
≫ HPCおすすめGPUサーバー
≫ AXXE-L by XTREME-D
≫ NVIDIA HPC SDK
≫ プログラム高速化サービス
クラウド製品
​≫ セキュリティ
≫ 仮想化
≫ リモートアクセス

≫ ハイブリッドクラウドNAS
≫ クラウドストレージ
​
​GPUコラム
≫ GPUプログラミング入門
≫ GPU Technology for CG/AI
≫ シリコンバレー発信 New Technology Repot
≫ ツブ子が聞く・見る・行く
​

導入事例
≫ DGX SYSTEMS
≫ GPU Computing
会社情報
≫ ごあいさつ
≫ 会社概要
≫ アクセスマップ
≫ 採用情報
​
≫ お問い合わせ
​
≫ 個人情報の取扱いについて
≫ 利用規約

​
≫ メルマガを購読する
グループ会社
Prometech Softwareサイト
©2021 GDEP Solutions,Inc.
  • Home
    • 新着情報一覧
    • NVIDIA NEWS
    • GPU2021 開催概要 >
      • Day1 AI & GPU セッション
      • Day2 イメージングAI
      • Day3 GPUスパコン
      • 講演レポート | Denso
      • 過去開催 >
        • GPU2020
        • GPU2019
  • GPU製品
    • 即納モデル
    • NVIDIA GPU 一覧 >
      • NVIDIA H100
      • NVIDIA A100
      • NVIDIA RTX スペック比較 >
        • NVIDIA RTX A6000
        • NVIDIA RTX A5000
        • NVIDIA RTX A4500
        • NVIDIA RTX A4000
        • NVIDIA RTX A2000
      • Quadro GV100
    • GPUレンタル一覧
    • GPU年定額プラン
    • AI・データサイエンスおすすめ一覧 >
      • NVIDIA DGX H100
      • NVIDIA DGX A100
      • DGX STATION A100
      • HP Z8 G4 Workstation
      • HP ZBook Fury17G7 Mobile Workstation
      • Dell Precision 7920 Tower
      • DeepLearning BOXⅡ
      • DeepLearning STATION
      • NVIDIA EGX サーバー
      • HITACHI SR24000
    • 数値計算・解析おすすめ一覧 >
      • HP Z4 G4 Workstation
      • HP Z8 G4 解析ソフトウェア動作確認済みモデル
      • Supermicro 7049GP
    • NVIDIA A100搭載 Supermicro 740GP
  • ストレージ
    • Synology
    • DDN STORAGE
    • PURESTORAGE
  • HPC
    • HPC Workstation
    • HPCおすすめGPUサーバー
    • HPC SIサービス
    • AXXE-L by XTREME-D
    • NVIDIA HPC SDK
    • プログラム高速化サービス
  • クラウド製品
    • セキュリティ >
      • KernelCare
    • 仮想化 >
      • Login VSI
      • Login PI
    • リモートアクセス >
      • FastX
      • NiceDCV
      • NoMachine
    • ハイブリッドクラウドNAS >
      • Morro Data
    • クラウドストレージ >
      • クラウドストレージ Wasabi
  • GPUコラム / 導入事例
    • GPUプログラミング >
      • 初級編
      • 中級編 >
        • 第4回:MPI+OpenACC実装における計算と通信のオーバーラップ
        • 第3回:拡散現象シミュレーションのおさらい
        • 第2回:簡単なOpenACC + MPI コードで考える
        • 第1回:複数のGPUを使う方法とは?
    • GPU Technology for CG/AI >
      • 深層学習を利用した画像処理・必要なGPU性能
      • トランスフォーマー 最近流行のニューラルネットワーク
      • GAN Inversion による写実的画像生成の制御
      • 深層学習におけるアノテーションコストを抑えるための取り組み Active Learning
      • 深層学習に基づく人物画像の再照明
      • GPUの起源と進化
      • AlphaGo とその後
      • CUDAを用いたシンプルなパストレーシング
      • 流体シミュレーションの応用
      • GPUを用いた高速レンダリング
      • GPUを基盤としたCG/AIの技術進化
    • シリコンバレー発信 New Technology Report >
      • 最新記事から
    • ツブ子が聞く・見る・行く! >
      • 見る!NVIDIA RTX A6000
      • 聞く!NVIDIA DGX A100
      • 見る!NVIDIA A100 Tensor Core GPU
    • 導入事例 >
      • DGX SYSTEMS >
        • NVIDIA DGX A100 | 金沢大学
      • GPU Computing
  • 会社情報
    • ごあいさつ
    • 会社概要
    • アクセスマップ
    • 採用情報
  • お問い合わせ