CUDAをRadeonで動かす（導入編）

Kernel Fusion Driver

Kernel Fusion Driver(KFD)とは、ROCmの各種プロジェクトを動かすための「Linuxカーネルと一体となったドライバー」のことです。

なので、後述の手順で説明しますが、カーネルを入れ替える必要があります。
個人的にはkmodでどうにかしてほしいんですが・・・。

ともかく、このKFDを入れることで、AMD Radeonの上で様々なことができるようになります。

HCC compiler

HCC compilerは、KFDの上で動く、異種混合実行環境向けのC++コンパイラーです。
HCCとは、Heterogeneous Compute C++の略のようです。

現在のところ、バックエンドとしてはAMD独自のAMD GCN ISAを生で出力するモードと、HSA財団の策定したHSAILという中間言語を出力するモードがあります。
ただし、今回はRadeonを使うので前者で十分なことと、既定で前者が使われること、加えてそもそも後者はこの先開発はされないようなので、ここではAMD GCN ISAを使うモードを前提として話を進めます
（とは言っても、バックエンドが何であるかはアプリケーション層ではあまり気にする必要はなさそうですので、よく分からない方は飛ばしても大丈夫です）。

また、フロントエンドにはC++17 Parallel STL、C++ AMP、OpenMP offload、HC C++ APIなどが用意されています。
今回は説明しませんが、C++ AMPをLinuxで動かしたい場合は、このHCC compilerを使うと動かせるようになります。

HIP

HIPは、CUDAをHCCでコンパイルできるようにするプラットフォームです。今回の話の主役になります。何の略称かは分かりません・・・。

HIP自体は、ほぼCUDAと同じ形式のRuntime APIやカーネル言語を既定するもので、基本的にcuXXXXをhipXXXXに置き換えたものです（例えばcuMemcpyはhipMemcpyに対応します）。

HIPで書かれたものは、hipccというコンパイラによって解釈され、バックエンドであるnvccまたはhccに渡すことで、NVIDIAでもAMDでも動かせるようにできます。
このhipccは、CUDAあるいはHC言語そのものへの直接的な翻訳機に近い動きをするので、HIPで書くことにオーバーヘッドなどはなく、性能劣化することは全くないらしいです。

また、hipifyという、CUDAからHIPへ（逆）変換するツールが付属しており、これを用いることで既存のCUDAコードを簡単にHIPが解釈できるようになります。

ということで、このHIPを信じるなら、CUDAで書いたコードを、性能劣化させることなくRadeonでも動く形にすることができる、ということになります。

環境の準備

ROCmのREADMEやROCm Hardware Requirementsに情報が散乱しているのですが、ROCmを動かすには、現在のところ以下の環境が必須です。

OS

Ubuntu 14.04.04またはFedora23

（おそらく動く：Ubuntu 16.04またはFedora 22）

CPU

Intel Haswell以降またはAMD Zen以降（PCIe Gen3 Atomicsのサポート必須）

GPU

Fiji以降のAMD GPU

今回は、Ubuntu 14.04、Intel Core i7-4790、AMD Radeon R9 FuryXの環境で試しました。

また、カーネルを入れ替えるという性質上、クリーンインストールすることにしました。
既存のシステムの上に入れても動かないわけではないと思いますが、何かあった時にサルベージが大変なので・・・。

あとこれ重要なのですが既存のAMDのプロプライエタリ・ドライバーとは衝突してまともに動かなくなるっぽいです。
公式ブログのコメント欄でしれっと書かれているのですが、
KFDでドライバーも提供されているので、ROCm導入時には既存のプロプライエタリ・ドライバーは除去する必要があります。
つまり、ROCmを入れるとOpenCLは使えなくなるということです（Issue参照）。
将来的にはサポートされるかもしれませんが・・・。

ということで、既存の環境をかなり破壊することになるので、試される方は新しいストレージを持ってきて、クリーンインストールして始めることをオススメします
（これで安全にはなりますが、カーネルの入れ替え等に限らず、ここに書いてあることは全て自己責任でお願いします）。

ROCmのインストール

計算機を用意して、OSをクリーンインストールしたら、まずは普通に起動します。

起動したら、とりあえずまずインストールされているパッケージを最新に更新します。
これはセキュリティーリスクもそうですが、今回用意した環境だと後でバージョン違いでやり直しという羽目になったので、ここで一度更新しておくことをオススメします。
Ubuntuならapt-get upgrade、Fedoraならdnf updateですね。

その後、ビルド済みバイナリをAMDのレポジトリから持ってきます。
Ubuntuの場合と
Fedoraの場合の両方が用意されています。
今回はUbuntuを使ったので、以下のコマンドを実行しました。

$ wget -qO - http://packages.amd.com/rocm/apt/debian/rocm.gpg.key | sudo apt-key add -
$ sudo sh -c 'echo deb [arch=amd64] http://packages.amd.com/rocm/apt/debian/ trusty main > /etc/apt/sources.list.d/rocm.list'
$ sudo apt-get update
$ sudo apt-get install rocm

なお、AMDのサーバーがとても遅いのか、恐ろしく時間がかかります。
具体的には、一晩放置しても終わってませんでした。半日まるっとかかるぐらいは覚悟して放置しましょう。

インストールが無事に終わったら再起動します。
再起動後、uname -rとかでカーネル名が4.4.0-kfd-compute-rocm-rel-1.1.1-10などになっていれば新しいカーネルに切り替わっています。

その後、手順に従って、HSAILが動くことを確認しましょう。
ちゃんと動けばインストール成功です。
これで、ROCmの一式が全て入っているので、HCC compilerもHIPも全て/opt/rocmの下にインストールされました。

hipify-clangの導入

先のレポジトリから持ってくる手順で、CUDAからHIPに変換するツールであるhipifyは/opt/rocm/bin/hipifyにインストールされました。
が、実はこの標準のhipifyだけでは変換が不十分になります。
これは、このhipifyがただの文字列置換Perlスクリプトであり、C++の構文解析を一切していないからです。

なので、hipify-clangと言う、clangベースのツールを別に入れる必要があります。
別のツールといっても、clang-hipifyもHIPの中の1つです（↑の手順で入らないのはまだβ版だからのようですね）。

手順は、HIPのmasterレポジトリをダウンロードして、
clang-hipifyのREADMEに従ってビルド＆インストールしましょう。
今回は以下のコマンドを実行しました。

$ cd ~
$ git clone https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP.git

$ cd HIP
$ wget http://llvm.org/releases/3.8.0/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz
$ tar xvfJ clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz

$ mkdir build
$ cd build
$ cmake -DBUILD_CLANG_HIPIFY=1 -DLLVM_DIR=~/HIP/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04/ -DCMAKE_BUILD_TYPE=Release ..
$ make
$ sudo make install

$ wget http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1404/x86_64/cuda-repo-ubuntu1404_7.5-18_amd64.deb
$ sudo gdebi cuda-repo-ubuntu1404_7.5-18_amd64.deb
$ sudo apt-get update
$ sudo apt-get install cuda-minimal-build-7-5 cuda-curand-dev-7-5

これで、/opt/rocm/hip/bin/clang-hipifyがインストールされます。

元にするCUDAコードの用意

ということで準備が終わったので、実際にCUDAコードを動かす手順に移ります。

ということで手始めに元にするコードを用意します。今回は簡単なベクトルの加算を用意しました。

__global__
void kernel(double z[], const double x[], const double y[])
{
 const auto i = (blockIdx.x * blockDim.x + threadIdx.x);

 z[i] = x[i] + y[i];
}

せっかくなので、これをまずは普通にCUDAの動く環境で動かしてみてください（vectorAddに移動してmakeで動くはずです）。

CUDAをHIPに変換する

CUDAコードを用意したら、いよいよ次はCUDAをHIPに変換する時です！

ということで、とりあえず普通についてくるhipifyを使ってみましょう。
hipify [yoursource.code]とやると、標準出力に変換結果が表示されます。
これをやると、例えば

#include "kernel.hpp"

__global__
void kernel(double z[], const double x[], const double y[])
{
 const auto i = (blockIdx.x * blockDim.x + threadIdx.x);

 z[i] = x[i] + y[i];
}


void Kernel(double z[], const double x[], const double y[],
        const std::size_t block, const std::size_t thread)
{
 kernel<<<block, thread>>>(z, x, y);
}

は、

#include <hip/hip_runtime.h>
#include "kernel.hpp"

__global__
void kernel(double z[], const double x[], const double y[])
{
        const auto i = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);

        z[i] = x[i] + y[i];
}


void Kernel(double z[], const double x[], const double y[],
        const std::size_t block, const std::size_t thread)
{
        hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(block), dim3(thread), 0, 0, z, x, y);
}

と変換されます。概ね一対一になっていて対応関係はそんなに難しいものではないと思います。

ただ、これは実は不完全です。これを後述のhipccに食わせてもコンパイルできません。
やってみると分かるのですが、カーネル関数の第一引数に必要な物がない、と言われます。
実はHIPの言語仕様として、カーネル関数の最初にはhipLaunchParm lpというものが必要ということになっています。

そのため、通常のhipifyを使っただけでは、変換後に手動でhipLaunchParm lpを挿入する必要があります。面倒ですね。

そこで、先にインストールしたhipify-clangの出番です。
これを使ってhipify-clang [yoursource.code] -o [yourhip.code]とすると、無事に全部変換してくれます。

しかしhipify-clangだけだとまだ不完全なようで、ひとつは、hip_runtime.hのincludeを追加してくれないんですね・・・。
ということで、現時点では正しく自動変換させるには、hipify-clangを書けた後に、通常のhipifyをする必要があるようです。

HIPのビルド

HIPに変換できたら、あとはhipcc [yourhip.code]とするだけでビルドできます。
あとはビルドされた実行ファイルを普通に実行できるはずです。

というわけで、HIPへの変換からビルドまでを一括でやってくれるMakefileを書きました。
ご覧のとおり、このMakefileでは以下のことをやってくれます。

1. カレントディレクトリの*.cppをHIP Runtime APIを使うように変更し、obj/*.hip.cppに保存
2. カレントディレクトリの*.cuをHIP Kernel Languageを使うように変更し、obj/*.hip.cuに保存
3. 各obj/*.hip.c(pp|u)を、hipccでビルドしてobj/*.c(pp|u).oにコンパイル
4. 全てのobj/*.c(pp|u).oをリンクして実行ファイルを出力

先述のベクトル加算の例にこのMakefileを置いたものも用意しました。
makeとやると、vectorAddという実行ファイルができます。

$ git clone git@bitbucket.org:LWisteria/hipvectoradd.git
$ cd hipvectoradd/vectorAdd

$ make
/opt/rocm/hip/bin/hipify-clang main.cpp -o main.cpp.hip
hipify main.cpp.hip -inplace
rm main.cpp.hip.prehip -f
mv main.cpp.hip obj/main.hip.cpp
hipcc -g -O3 -Wall -I/opt/rocm/include/hip -std=c++11 -I. -c obj/main.hip.cpp -o obj/main.cpp.o
/opt/rocm/hip/bin/hipify-clang kernel.cu -o kernel.cu.hip
hipify kernel.cu.hip -inplace
rm kernel.cu.hip.prehip -f
mv kernel.cu.hip obj/kernel.hip.cu
hipcc -g -O3 -Wall -I/opt/rocm/include/hip -std=c++11 -I. -c obj/kernel.hip.cu -o obj/kernel.cu.o
hipcc -o vectorAdd obj/main.cpp.o obj/kernel.cu.o

$ ./vectorAdd
Device: Fiji
0: -1.38309, -1.38309
1: -1.91112, -1.91112
2: -0.682073, -0.682073
3: 1.17407, 1.17407
4: 1.95078, 1.95078
5: 0.93395, 0.93395
6: -0.941547, -0.941547
7: -1.95139, -1.95139
8: -1.16713, -1.16713
9: 0.69018, 0.69018

実行結果はあまりちゃんと書いてませんが、main.cppを読んでもらうと分かる通り、左側の数値がCPUの実行結果で、右側がGPUでの実行結果です。
ちゃんと実行できたことが分かりますね！