この記事は何？

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

そもそもOpenACCとは？

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

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

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

しかし、

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

により、注目を浴びてきています⁸⁹。

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

// ホストと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にオフロードが追加された¹¹のでそっちに流れたほうが良いという記事がQiita¹²でも書かれていましたが、それはもう過去のことです。第一にNVIDIA自身がOpenACCコンパイラの開発に関わっている以上、OpenMPがNVIDIA GPUでOpenACCより高速化（最適化）を行える可能性はあなたが異世界転生してしまうレベルで限りなくゼロです¹³。

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

ここの説明は抜かしてしまっても良かったんですが、話がぶっ飛んでしまうのでなんでこんなことを検証しようとしているのかを説明します。
これ、実は一歩先の技術の検証だったりします。というのもみなさん多分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による機械学習（ディープラーニング）の高速化があります¹⁴。その他、Bingやニコニコ動画にFPGAが採用されたというニュースも記憶に新しいと思います。FPGAは自由に回路を書き込めるので、GPUと完全に競合しないデバイスです。しかも、専用の回路を構築できるためCPUより高速なことが多いです。

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

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

PGI Unified Binary と OpenACC Runtime

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

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

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

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

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

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

#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コンパイラでは実装可能だということなので、マルチデバイス実行に対しての一つの希望ということになります。

発展的な話

ただし、OpenMP云々よりも課題はもっと別にあります。

1. ランタイム関数では変なところで実行されると意図しないデバイスの切り替えが発生してしまうこと
2. このバイナリーは無駄が多くてファイルがでかいということ
3. 移植性に影響すること

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

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

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

私は、特に1つ目の問題からどのデバイスで実行するかはディレクティブでオフロードするコードごとに指定できるようにしたほうが良いと思います。なお、ディレクティブでのオフロード先デバイス指定は先程の筑波大¹⁷でやられているのですが、独自に拡張しての実装のようなので、今後OpenACCの仕様に追加されることを期待しています。