Help us understand the problem. What is going on with this article?

CUDAより高速で簡単なOpenACCのランタイム関数で実行中のPGI Unified Binaryの演算デバイスを切り替えられるのか検証してみた

この記事は何?

CUDA123 に代わる新しい GPU の 標準API 、OpenACC が最近 CUDA よりも高速という結果4が出てきて調子が良いので、現状のヘテロジニアス環境でのマルチデバイス実行(複数のデバイス:CPU、GPU、FPGA、メニーコアCPUを組み合わせた演算)の対応状況を調査してみました。

そもそもOpenACCとは?

NVIDIAが中心となって本腰を入れているGPGPUなどのハードウェアアクセラレーションの標準APIです。その特徴として最も特筆すべき点はCUDAと比較してプログラミングが超簡単なことです。以前なら超は言いすぎだったかも知れませんが、現在ではCUDAと色褪せない性能が出せる5のでこれくらい言ってもいいと思います。

今からGPGPUを始めるならCUDAではなくOpenACCをおすすめします。

NVIDIAが本格的な布教活動を世界各地で無料で行っているので(当然日本でもやってます)、今後確実に普及すると思います。トレンドとして抑えておくべきです。
過去の例:
- https://developer.nvidia.com/gpubootcamp/RIKEN-CCS
- https://www.ccs.tsukuba.ac.jp/bootcamp-20191202/
今後の講習の開催については「GPU Bootcamp」でググると出てきます。

ターゲットデバイスは NVIDIA GPU に限らず、AMD GPU やメニーコアCPUに対応しています。さらにFPGAにも対応したコンパイラも開発されています6
https://ja.wikipedia.org/wiki/OpenACC
https://www.openacc.org/
これまでは、高性能なコンパイラが無料で公開されていなかったこと7、OpenACC自体がCUDAとの演算速度に水を開けてしまっていたこと、などを理由として普及してきませんでした。

しかし、

  • 2年ほど前にNVIDIA傘下で開発されている無料版の PGI Compiler が公開されたこと
    (詳細は https://www.softek.co.jp/SPG/Pgi/pgi_community.html
  • OpenACCコンパイラによる最適化が形になってきて(特に2019年現在最新のVoltaアーキテクチャ8で)CUDAよりも高速な結果4も出ていること

により、注目を浴びてきています910

記法も少しだけ説明します。OpenACCはOpenMPのようなディレクティブ形式(C、C++なら#pragma ~)の規格です。C++、C、Fortranに対応しています。CUDAみたいな独立言語ではないため、既存のCPUコードも簡単にGPU化出来ます。
簡単に書けると言いましたが、基本的には以下の3つの構文だけ入れればよいという感じです。あとは、CUDA Unified Memory を有効にしてコンパイルすると 何も考えなくても かなりの高速化が得られます11。勿論、最高性能を出すにはもう少し最適化が必要ですが、それでもCUDAに比べればプログラミングエフォートは低いです。

openacc_sample.c
// ホストとGPU間でやり取りするデータの指定(無くても動く)
#pramga acc data copy(a[:N], b[:N], result[:N])
{
// GPUでの演算対象コードの指定
#pragma acc kernels
// それぞれのループに対する指定(無くても動く、independentは独立した配列の意)
#pragma acc loop independent
    for (i=0; i<N; i++) {
        result[i] = a[i] * b[i];
    }
}

また、一時期、OpenMPにオフロードが追加された12のでそっちに流れたほうが良いという記事がQiita13でも書かれていましたが、それはもう過去のことです。PGI Compiler の OpenACC → 最適な GPU バイナリの生成能力はNVIDIA自身が力を入れて何年も掛けてようやく形になりました。果たして NVIDIA が直接且つ意欲的に OpenMP コンパイラの開発を行うか考えると、NVIDIA GPU の性能を OpenACC と同程度に OpenMP で得ることはほぼ不可能でしょう14

ヘテロジニアス環境でのマルチデバイス実行とは?

ここの説明は抜かしてしまっても良かったんですが、話がぶっ飛んでしまうのでなんでこんなことを検証しようとしているのかを説明します。
これ、実は一歩先の技術の検証だったりします。というのもみなさん多分GPUだけとかもしくはクラスタでCPUだけ使って演算速度を極めてる人多いと思うんです。

でも、GPUで処理している間にCPUで処理するみたいなこと(非同期実行)したらもっと早くなると思いませんか?

実は現状でも出来ます。ただし、そこまで大きな効果が得られない割に面倒なのでやられることは少ないです。

でもGPUを使ったことがある人ならわかると思いますが、GPUで全ての計算が早くなるわけではありません。とはいえCPUが遅いからGPUで演算加速しているわけで、CPUに戻すのは解せません。

そこで、専用のプロセッサ(ASIC:エイシックと言います。Domain Specific Architecture:ドメイン指向アーキテクチャとも)を組み合わせて更に高速化できるようにしようということが構想されているのです。それがマルチデバイス実行です。ちなみにマルチ(多数)なので、CPUとGPUだけじゃなくて3つ以上の組み合わせのことを暗に含んでいます。ヘテロジニアス環境というのはCPU以外のデバイスがハードウェアアクセラレータとして搭載されているマシンのことです。GPU搭載マシンも含みます。

ASICが具体的に思い浮かばない人もいると思いますが、TPUもその一つです。ただ、TPUはGPUと競合しているので組み合わせても面白くないと思います。

あと、ASICではないですが、FPGAもアクセラレータとして利用されています。実例として、東工大でのFPGAによる機械学習(ディープラーニング)の高速化があります15。その他、Bingやニコニコ動画にFPGAが採用されたというニュースも記憶に新しいと思います。FPGAは自由に回路を書き込めるので、GPUと完全に競合しないデバイスです。しかも、専用の回路を構築できるためCPUより高速なことが多いです。

まだマルチデバイス実行は研究段階ではあります。スーパーコンピューターでは今や上位の殆どがアクセラレータ搭載マシンとなっているのですが、マルチデバイスなマシンは世界的にも今年になって筑波大学のCygnus(GPUとFPGAを搭載したスパコン)16が出てきたくらいです。

ですが、各所は対応準備をしています。例えば、Intelが今年発表予定のoneAPIはマルチデバイス実行には直接言及していないもののCygnusのようなマシンを狙っています。もちろん、OpenACCでもマルチデバイス実行に対応させようと議論がされています。その一つが筑波大でやられているのですが、これは研究中なので、今回はPGIの実装について検証します。

PGI Unified Binary と OpenACC Runtime

前置きが長かった割に、興味を持っていただけなかったら大変悲しいのですが、本題に入ります。

OpenACCでは、通常一つのデバイスしかハードウェアアクセラレータとして使えません。なぜなら、OpenACCのディレクティブには現状マルチデバイス実行に対応するものは無いからです(2019/11/23追記:3.0から追加されたようですがまだ実装されているコンパイラはありません)。そもそもOpenACCが出始めの頃はマルチデバイス実行なんてまだ検討されていませんでした。ASICやFPGAがアクセラレータとして注目されるようになったのは本当にここ3年くらいのことです。

また、複数のデバイスを使用するにはそれぞれのデバイス用に異なるバイナリコードを生成する必要があります。これについてはOpenACCの仕様17によって異なるデバイス用のバイナリコードを含んだファイルを生成するコンパイラ実装が示唆されています。しかし、マルチデバイス実行など考慮されていない既存のコンパイラではコンパイルする際に使用するデバイスを指定するようになっています。

ただし、OpenACCにはランタイム関数が用意されています。これを使うことで、コード中で使用するデバイスを指定できるのです。これは、通常環境変数で指定するようになっているOpenACCディレクティブ挿入部(オフロード部)の演算デバイスをコード内で定めてしまうために用意されたもので、当然一回決めたら変更しないものとして策定されていたはずです。ところが、これに限らずOpenACCではデバイスの設定などに関する振る舞いは仕様で決められておらず、実装依存となっています。もし、コンパイラが吐き出したバイナリファイルが複数のデバイス用コードを含んでいて、且つそのファイルがちゃんとオフロード先のデバイスの変更を受け付けるようになっていれば、実行中に切り替えができるはずです。

PGI Unified Binary はまさに前者を実現したPGI Compilerの機能です。ホストCPUとGPUしか選択できませんが、少なくとも複数のプロセッサで演算するバイナリコードが一つのファイルに含まれています。

というわけで、後者の「演算デバイスが実行途中で切替可能か」を検証してみました。
実行環境は、PGI Compiler 18.10、CUDA 9.2.148、GPUの世代はPascalです。

pgi_unified_binary_test.c
#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>

int isHost;
int isNvidia;
void funcHOST(float* a, float* b, float* c, int size);
void funcGPU(float *a, float *b, int size);

int main(int argc, char** argv) {

    int size = 256 * 16;

    float* A = (float*) malloc(size * sizeof(float));
    float* B = (float*) malloc(size * sizeof(float));
    float* C = (float*) malloc(size * sizeof(float));
    float* D = (float*) malloc(size * sizeof(float));

    int i, error=0;

    for (i = 0; i < size; i++) {
        A[i] = (float) i;
        B[i] = (float) i * 100;
    }

printf("acc_get_device_type(default): %d\n", acc_get_device_type());
acc_set_device_type(acc_device_host);
printf("acc_get_device_type: %d\n", acc_get_device_type());

    funcHOST(A, B, C, size);
printf("isHost: %d\n", isHost);

    for (i = 0; i < size; i++) {
        if (C[i] != (float) i + (float) i * 100) error++;
    }

    printf("errorHOST:%d\n", error);

acc_set_device_type(acc_device_nvidia);
printf("acc_get_device_type: %d\n", acc_get_device_type());

    funcGPU(C, D, size);
printf("isNvidia: %d\n", isNvidia);

    for (i = 0; i < size; i++) {
        if (D[i] != (float) i + (float) i * 100 + 1.0) error++;
    }

    printf("errorGPU:%d\n", error);

    return 0;
}

void funcHOST(float* a, float* b, float* c, int size) {
    int j;
#pragma acc data copyin(a[0:size], b[0:size]) copyout(c[0:size]) copy(isHost)
{
#pragma acc kernels
{
#pragma loop independent
        for (j = 0; j < size; j++) {
            c[j] = a[j] + b[j];
        }
isHost = acc_on_device(acc_device_host);
}
}
}

void funcGPU(float *a, float *b, int size) {
    int j;
#pragma acc data copyin(a[0:size]) copyout(b[0:size]) copy(isNvidia)
{
#pragma acc kernels
{
#pragma acc loop independent
    for (j = 0; j < size; j++) {
        b[j] = a[j] + 1.0;
    }
isNvidia = acc_on_device(acc_device_nvidia);
}
}
}

コンパイルコマンド

pgcc -Minfo -acc -ta=tesla,host pgi_unified_binary_test.c

実行結果

acc_get_device_type(default): 4
acc_get_device_type: 2
isHost: -1
errorHOST:0
acc_get_device_type: 4
isNvidia: 1
errorGPU:0

結果ですが、isHost、isNvidiaは、Cなので非ゼロでTRUEです。
なお、acc_on_deviceの引数を入れ替えた場合、どちらも0(FALSE)になりました。

結論

PGI Unified Binary では、コードの実行中にプロセッサの切り替えが可能

尤もホストCPUで実行するなら普通にOpenMP使うよって感じかもしれません。しかし、今回の検証の議論の範疇ではないのでそれは置いときましょう。
一応実行中に演算デバイスを変更できるバイナリがOpenACCコンパイラでは実装可能だということなので、マルチデバイス実行に対しての一つの希望ということになります。

発展的な話

現状のOpenACCでマルチデバイス実行を行うには課題があります。

  1. ランタイム関数では変なところで実行されると意図しないデバイスの切り替えが発生してしまうこと
  2. このバイナリーは無駄が多くてファイルがでかいということ
  3. マルチデバイスを使うために別バージョンの用意が必要なこと

この1つ目ですが、結構深刻な問題です。仮に、複数のファイルに分割してプログラムを記述した場合、メイン関数外でランタイム関数acc_set_device_type()を呼び出すことがあるかも知れません。しかし、関数だとどこにでも書けてしまうため、デバイスがどこで切り替わるのかわかりにくいです。また、同じデバイスを使っているうちはこの関数をプログラマはわざわざ書かないため、後々の修正でバグを発生させる原因になります。

2つ目ですが、これはOpenACCでオフロードしようとしている2箇所とも、CPU用とGPU用のバイナリをそれぞれ生成しているからです。一箇所に付き片方のデバイス用のみで十分です。
なぜなら、OpenACCでいくら簡単にコードが書けるとは言え、性能を引き出すにはプロセッサごとに異なる記述が必要だからです。また、GPUで高速に実行できる計算を他のプロセッサでやってもASICでもない限りは優位な差は出ません。同じソースコードから生成してもどちらかしか速く実行できないうえに、GPUとCPUならどっちが高速か実験するまでもなく分かることも多いので、あまり需要が無いのです。device_type 節でそれぞれのデバイスに異なる並列度を指定すればある程度解決できる可能性はありますが、デバイス実行部分のコードを書き換えたほうが速くなる場合もあるため、完全な解決は難しいです。

3つ目はデバイスごとの最適化記述もあるため、最高性能を出したい時はあまり重要ではないかも知れませんが、OpenACCが簡単にデバイスを演算させられるが故に、GPU以外のデバイスも搭載しているマシンのユーザーなら少しでも資源を有効活用したいと思って試しにマルチデバイス実行したくなるかも知れません。しかし、ランタイム関数ではディレクティブのように通常のコンパイルで無視されないので気軽に試してみることは難しくなります。

私は、特に1つ目の問題からどのデバイスで実行するかはディレクティブでオフロードするコード範囲ごとに指定できるようにしたほうが良いと思います。

2019/11/23追記:OpenACC 3.0からランタイム関数に対応するディレクティブが追加されたようですが、オフロードするコードブロックごとの指定ではなくてランタイム関数のようにどこにでも挿入可能なようです。


  1. CUDAを使ったGPUプログラミング超入門 

  2. CUDAプログラミング第一歩 

  3. 機械学習?ディープラーニング?知識不要!TensorFlow/Kerasのパワーを最も簡単に体感する方法 

  4. OpenACCがCUDAより高速な例 https://dl.acm.org/citation.cfm?id=3218228 

  5. CUDAと並ぶ性能の例(ABCI) https://waccpd.org/wp-content/uploads/2019/11/ws_waccpd_yamaguchi.pdf 

  6. OpneARC by 米国ORNL https://pdfs.semanticscholar.org/9712/65b3150c5743f9033b5e06ed50cc40cd404d.pdf 

  7. GCCのOpenACCを試してみた 

  8. Turing アーキテクチャの方が新しいですが、Volta を異なるターゲット向け(HPC向けではない)にチューニングしたバージョンで、基本設計は同じです。 

  9. PGIコンパイラを用いてOpenACCを利用 

  10. OpenACCでCUDA-awareなMPI_Send,Recvをする 

  11. これ有効にしてればぶっちゃけ#pragma acc kernelsだけでも高速化できちゃうと思う 

  12. OpenMP と OpenACC の関係 

  13. CUDAとOpenCLどっちがいいの? 

  14. GPUのコンパイラはアーキテクチャ設計者しか最適化出来ないと言っても過言ではなく、GNUのOpenACCの実装はあまりにも大変なためか半ば心が折れてる感じすらします。PGI以外のOpenACC実装も殆どがCUDA変換によるソース to ソースコンパイラです。また、後発で機能的な重複のある OpenMP オフロードの最適化にNVIDIA自身が意欲的になることは考えにくいです。AMDがFrontier(米国ORNLの世界一速いスパコンSummitの次のマシン)の開発に合わせて、OpenMP のオフロードを実装するみたいですが、OpenACCが5年以上かかったものを最適化するのは相当苦労が必要だと思います。ついでに個人的な意見ですが、Intel側(OpenMP陣営)に迎合してもIntelの独壇場になってしまいそうなので、GPUしか競合してないNVIDIA側のOpenACCに対応させるほうが良かったんじゃないかなと思います。 

  15. 東工大での機械学習向けFPGA活用 https://ainow.ai/2018/12/13/158542/ 

  16. 筑波大のCygnus https://news.mynavi.jp/article/20190403-801024/ 

  17. OpenACC 3.0 https://www.openacc.org/sites/default/files/inline-images/Specification/OpenACC.3.0.pdf 

takashift
IT土方にはなりたくない学生エンジニア。最近は髭剃ってます。 将来はhage散らかす予定。
https://higechira.hatenablog.com/
Why not register and get more from Qiita?
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
Comments
No comments
Sign up for free and join this conversation.
If you already have a Qiita account
Why do not you register as a user and use Qiita more conveniently?
You need to log in to use this function. Qiita can be used more conveniently after logging in.
You seem to be reading articles frequently this month. Qiita can be used more conveniently after logging in.
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
ユーザーは見つかりませんでした