この記事は何？

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

そもそもOpenACCとは？

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

今から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にも対応したコンパイラも開発されています⁶。
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¹³でも書かれていましたが、それは過去の状況下での話です。PGI Compiler の OpenACC → 最適な GPU バイナリの生成能力はNVIDIA自身が力を入れて何年も掛けてようやく形になりました。少なくとも、OpenACC が NVIDIA GPU 向けの最も標準なディレクティブ形式のAPIであることは、今後も揺るがないでしょう¹⁴。

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

ここの説明は抜かしてしまっても良かったんですが、話がぶっ飛んでしまうのでなんでこんなことを検証しようとしているのかを説明します。
これ、実は一歩先の技術の検証だったりします。というのもみなさん多分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のディレクティブには現状マルチデバイス実行に対応するものは無いからです（2019/11/23追記：3.0から対応可能そうなものが追加されたようですがまだ実装されているコンパイラはありません）。そもそも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コンパイラでは実装可能だということなので、マルチデバイス実行に対しての一つの希望ということになります。

発展的な話

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

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

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

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

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

2019/11/23追記：OpenACC 3.0からランタイム関数に対応するディレクティブが追加されたようですが、dataディレクティブのような指定範囲への適用ではなく、機能的に互換（どこにでも挿入可能）なようです。