武蔵野 Advent Calendar 2018 18日目の記事です。（投稿日は気にしないように！）
概要
Gentoo上でROCmというAMDのGPGPUソフトウェアスタックを使って、OpenCLを動かしたり、hccというヘテロジニアスコンパイラを使って、
GPGPUコードをCPUコードと同じ様に記述できる機能を紹介します。
その後、R9 Nano / RavenRidge で性能測定してみます。
ROCm 2.0.0
ROCmとはAMDのGPUを計算用途に使うためのソフトウェアスタックですかね？
Next Horizonというイベントでアナウンスされていましたが、つい先日に2.0.0のtagが打たれていました。
https://github.com/RadeonOpenCompute/ROCm/blob/master/README.md
今回はROCmを構成するソフトウェアにうち、HSAランタイム(ROCT-Thunk-Interface/ROCR-Runtime), hccをGentoo上に構築して、Radeon R9 Nano(Fiji, dGPU)とRyzen 7 2700U(RavenRidge, APU, 公式にはROCmはAPU非対応)で動かしてみたいと思います。
インストール方法
overlayを公開しているのでこれを追加するか、適当にebuildをコピーしてきて
emerge hsa-rocr rocm-opencl hcc
とするだけです。多分インストールできます。カーネルは最新版を使いましょう。
clinfo / hsainfo
R9 Nano / RavenRidge 上で clinfo / hsainfo を実行した結果です。
https://gist.github.com/kazuki/1620a990984468ffc2387ea839a77b92
ROCm 1.9.2からの差分としてはRavenRidge上でも搭載メモリ全体を利用できるように報告されるようになった点ぐらいでしょうか。
(1.9.2では1GBしか利用できなかった)
ROCm 2.0.0からOpenCL2.0対応とリリースノートには書いてありましたが、1.9.x時代からランタイムはOpenCL 1.2でカーネルはOpenCL 2.0でしたので、
ココらへんの出力に差異はありませんでした。
hcc
OpenCLの場合、CPU上で実行するホストコードとは別にGPU上で実行するコードをOpenCL Cで用意する必要がありました。
そのため、外部のOpenCL Cソースコードを実行時に読み込むのが面倒くさいとか、
デバッグがだるいとか、CPU上のコードとコード共有しづらいという問題がありました。
hccは同じソースコードからCPU側のコードとGPU側のコードを両方とも生成することで、
より簡単にGPU側に計算をオフロードさせることが出来る様になる、コンパイラのことです。
例えば、y = ax + y (SAXPY: Single-Precision A·X Plus Y) を実行する場合、OpenCLでは以下のようになります
#include <cl/cl2.hpp>
int main() {
const float a = 100.f;
const int N = 1024 * 1024;
float *x = new float[N];
float *y = new float[N];
// OpenCL Cのコードを実行時にコンパイルして、関数を取得
std::string code{R"CLC(
__kernel void saxpy(__global float *src, __global float *dst, float factor) {
long i = get_global_id(0);
dst[i] += src[i] * factor;
}
)CLC"};
saxpy_program = cl::Program(code);
saxpy_program.build("-cl-std=CL2.0");
cl::KernelFunctor<cl::Buffer, cl::Buffer, float> kernel(saxpy_program, "saxpy");
// GPU側がCPU側のメモリを参照できるようにしてあげる
cl::Buffer src(CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR | CL_MEM_HOST_NO_ACCESS, N * sizeof(float), x);
cl::Buffer dst(CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR | CL_MEM_HOST_READ_ONLY, N * sizeof(float), y);
// GPU側でSAXPYを計算
kernel(cl::EnqueueArgs(cl::NDRange(N)), src, dst, a).wait();
...
}
しかし hcc を使うと以下のようになります。
#include <hc.hpp>
int main() {
const float a = 100.f;
const int N = 1024 * 1024;
float *x = new float[N];
float *y = new float[N];
hc::array_view<float, 1> av_x(N, x);
hc::array_view<float, 1> av_y(N, y);
hc::parallel_for_each(hc::extent<1>(N), [a, av_x, av_y](hc::index<1> i) [[hc]] {
// GPU上で実行される
av_y[i] += a * av_x[i];
});
}
CPU側とコード共有したいときは共有したい関数等に__HC__と__CPU__を指定するだけです。
(__HC__だけだとCPUでは実行できない関数になる)
void saxpy(float a, const float *x, float *y, int i) __CPU__ __HC__ {
y[i] += a * x[i];
}
int main() {
const float a = 100.f;
const int N = 1024 * 1024;
float *x = new float[N];
float *y = new float[N];
hc::array_view<float, 1> av_x(N, x);
hc::array_view<float, 1> av_y(N, y);
hc::parallel_for_each(hc::extent<1>(N), [a, av_x, av_y](hc::index<1> i) [[hc]] {
// GPU上で実行される
saxpy(a, av_x.data(), av_y.data(), i[0]);
});
// CPU上で実行される
float x = 1.f, y = 2.f;
saxpy(10.f, &x, &y, 0);
}
hccを使えば使い慣れたC++でそのままGPU側で動作するコードが書けるし、CPU側とコードを共有することも簡単ですね。
(ちなみに仕組みとしてはCPU用とGPU用コンパイラを計2回走らせているようで、コードにエラーがあるときは同じコンパイルエラーが二回出力されます)
HSAの夢
2012年頃にAMDらが発表したHSAでは、CPU-GPUでメモリ空間を利用し、
メモリコヒーレンシをとることで、GPU等をより簡単に使う仕組みだった。
もし、CPU-GPU間で同じメモリ空間を利用しコヒーレンシがとれていれば、hccを使った場合に
#include <hc.hpp>
int main() {
const float a = 100.f;
const int N = 1024 * 1024;
float *x = new float[N];
float *y = new float[N];
hc::parallel_for_each(hc::extent<1>(N), [a, x, y](hc::index<1> i) [[hc]] {
y[i] += a * x[i];
});
}
と記述することが可能となり、hc::array_viewといったメモリ管理をする必要がなくなるはずだが...が、
最近のAMDはAPUよりもdGPUの方に重きを置いているので、現在の仕組みではdGPUとキャッシュコヒーレンシをとるのは難しく、
メモリ管理用のコードを書く必要がある...Zen2世代でPCIe 4.0 + CCIXとかでdGPUでもhUMAが実現してくれたらなぁ...
RavenRidge (Ryzen 7 2700U, APU)
HSAを実現したAPUとしてKaveri/Carrizo...などがあったが、これらのCPUが出た当時は(一般に入手可能な)ソフトウェアスタックが不十分だったし、
メモリコヒーレンシをとったメモリアクセスはあまり帯域が出なかった。
RavenRidgeはCPU-GPU間のインターコネクトもInfinity Fabricに統合されており(参考)、
より高いメモリアクセス性能が期待できそうである。
ROCmのHSA Finalizerはgfx902(RavenRidge)向けに提供されておらず、またhccも同様だが、なぜかOpenCL経由ではgfx902であっても実行できる。
ベンチマーク (SAXPY)
ソースコード: https://github.com/kazuki/rocm-hcc-opencl-saxpy-sample
256* 2^20個の要素(1GB)のsaxpyを計算するベンチマークを実行してみた。メモリ的には2GBロードして1GB書き込む感じかな。
cpuはsaxpyをcpuで計算した場合の時間、initは計算元データのCPU側メモリへのコピー時間、gpuはGPUへのメモリ転送と計算時間, verifyはCPUの計算結果とGPUの計算結果の比較時間。3回実行し平均を表示している。括弧内は分散。
- R9 Nano
[HCC ] cpu:0.1263(0.0000) init:0.4421(0.0000) gpu:0.5980(0.0001) verify:0.0761(0.0000)
[HOST_PTR ] cpu:0.1254(0.0000) init:0.4441(0.0000) gpu:0.2466(0.0000) verify:0.0766(0.0000)
[SVM Coarse-grained] cpu:0.1284(0.0000) init:1.0990(0.0000) gpu:0.7645(0.0069) verify:0.0802(0.0000)
[SVM Fine-grained ] cpu:0.1318(0.0000) init:0.7073(0.0000) gpu:0.2002(0.0000) verify:0.0866(0.0000)
[HOST_PTR ] cpu:0.1647(0.0000) init:0.4429(0.0000) gpu:0.2472(0.0000) verify:0.2173(0.0000) # gcc
[SVM Coarse-grained] cpu:0.1619(0.0000) init:1.1017(0.0000) gpu:0.7087(0.0000) verify:0.2190(0.0000) # gcc
[SVM Fine-grained ] cpu:0.1684(0.0000) init:0.7086(0.0001) gpu:0.2010(0.0000) verify:0.2296(0.0001) # gcc
- RavenRidge
[HOST_PTR ] cpu:0.1335(0.0000) init:0.5201(0.0000) gpu:0.1279(0.0000) verify:0.3639(0.0354)
[SVM Coarse-grained] cpu:0.1367(0.0000) init:0.6080(0.0000) gpu:0.1236(0.0000) verify:0.3598(0.0208)
[SVM Fine-grained ] cpu:0.1382(0.0000) init:0.6063(0.0000) gpu:0.1234(0.0000) verify:0.5351(0.0014)
[SVM System ] cpu:0.1348(0.0000) init:0.5252(0.0001) gpu:0.1282(0.0000) verify:0.4239(0.0438)
[HOST_PTR ] cpu:0.1770(0.0000) init:0.5168(0.0000) gpu:0.1275(0.0000) verify:1.0153(0.0242) # gcc
[SVM Coarse-grained] cpu:0.1790(0.0000) init:0.5985(0.0000) gpu:0.1235(0.0000) verify:0.9307(0.0135) # gcc
[SVM Fine-grained ] cpu:0.1796(0.0000) init:0.5988(0.0000) gpu:0.1235(0.0000) verify:0.9006(0.0258) # gcc
[SVM System ] cpu:0.1784(0.0000) init:0.5216(0.0000) gpu:0.1276(0.0000) verify:0.8043(0.0012) # gcc
- (参考) Core i7-7600U(Kabylake) beignet 1.4
[HOST_PTR ] cpu:0.1357(0.0000) init:0.4919(0.0000) gpu:0.7762(0.0000) verify:0.1360(0.0000) errors:268435323
[SVM Coarse-grained] cpu:0.1384(0.0000) init:0.4823(0.0001) gpu:0.3126(0.0000) verify:0.1403(0.0000)
dGPUではCoarse-grainedがFine-grainedよりも遅いという結果になったが、SVMを使わずに従来のホストのメモリアドレスをBufferに渡してあげるのが一番速かった。hccはちょっと遅め...
RavenRidge上ではどれを使ってもあまり速度はかわらないため(init+gpu+verifyの合計値が実際の処理オフロードにかかった時間)、バッファ管理が不要でCPU側のポインタをそのまま渡せるSVM fine-grained systemがとても楽で良い感じ。APUではverifyに時間がかかっているので、GPU側で更新したメモリ領域は実際に触りに行ったりしないと、CPU側からは参照できるようにはなっていないのかな？
メモリ帯域が律速になりそうだけど、dGPU版の方がinit+gpu+verifyの時間が短い。RavenRidgeでInfinity Fabricに統合されたんだけどな...
おわりに
ネタがなかったため、まとまりがない、作業ログ的な記事になってしまいましたが、とりあえず25日までには間に合わせられたような...