HIPを使ってAMDのMI50を動かしてみた

最近AMDのGPU( https://www.amd.com/ja/products/professional-graphics/instinct-mi50 )を手に入れたので、
ちょっくらopenCLでも動かしてみるかと思ったら、なぜかlibOpenCLが見つからんと言われてさっぱりだったので、
AMDの開発ツールROCm( https://rocmdocs.amd.com/en/latest/index.html )についでに入ってきたHIPとかいう言語(？)を使って動かしてみる事にしました。

HIPはAMDが制定している、ヘテロジニアス計算機向け開発環境だそうな。
なんでも、**H**eterogeneous-compute **I**nterface for **P**ortabilityの略だとかなんだとか。
文法としてはほぼほぼCUDAで、実際CUDAコードをhip化するhipifyというperlコードが付属していたりします。
実際、CUDAとの主な違いは、

 - cudaやcuから始まる関数はhipから始まるものに置き換える
    - ```cudaMemcpy```は```hipMemcpy```にするなど
 - デバイス関数実行時は```hipLaunchKernelGGL```という関数を使う

という2点くらいです~~じゃあもうCUDAコードそのままコンパイルしてくれればいいのに…~~。
さて、実際CUDAコードを書いて、をれをhipfyしてもいいんですが、どうも大して変わらないように見えるので、はじめからHIPで書いちゃって動かす事にしてみました。

# コード
cmakeとかでMakefile作ってもいいんですが、別にそこまでの物でも無いので、HIPのインストールパスからコンパイラの位置指定。
僕の環境では```/opt/rocm-4.0.0/hip```にありました。
実際にはclangが動いているのかな？
OpenMPとかも問題なく使えました。
場合によってはincludeへのパスも指定しなくちゃいけないですね。

```makefile
HIP_PATH = /opt/rocm-4.0.0/hip
HIPCC = $(HIP_PATH)/bin/hipcc

CXXSRC = main.cpp
CXXFLAGS += -O3 -fopenmp

all:$(CXXSRC)
    $(HIPCC) $(CXXFLAGS) $(CXXSRC)

clean:
    rm -f *.o *.out
```

で、コードですが、拡張子は普通に```.cpp```でよいみたいです。

```cpp
#include <hip/hip_runtime.h>
#include <omp.h>
#include <iostream>
#include <cstdlib>

template <typename real> __global__ void Add(real const * const a, real const * const b, double * const c, const std::size_t N){
	int i = blockDim.x * blockIdx.x + threadIdx.x;
	c[i] = a[i] + b[i];
}

template <typename fp> class on_device{
	public:
	fp* host;
	fp* dev;
	const std::size_t N;
	on_device(const std::size_t N_):N(N_){
		host = new fp[N];
		hipMalloc(&dev, N * sizeof(fp));
	}
	~on_device(){
		delete []host;
		hipFree(dev);
	}
	void sendDataToDevice(){
		hipMemcpy(dev, host, N * sizeof(fp), hipMemcpyHostToDevice);
	}
	void recvDataFromDevice(){
		hipMemcpy(host, dev, N * sizeof(fp), hipMemcpyDeviceToHost);
	}
	fp* getDevicePtr() const{
		return dev;
	}
	fp& operator[](const std::size_t i){
		return host[i];
	}
	std::size_t getSize() const{
		return N;
	}
};

int main(void){
	const int N = 1024;
	on_device<double> a(N);
	on_device<double> b(N);
	on_device<double> c(N);

	for(int i = 0 ; i < N ; ++ i){
		a[i] = (double)rand() / (double)RAND_MAX;
		b[i] = (double)rand() / (double)RAND_MAX;
	}
	a.sendDataToDevice();
	b.sendDataToDevice();
	int threadPerBlock = 256;
	int blockPerGrid = (N + threadPerBlock - 1) / threadPerBlock;
	hipLaunchKernelGGL(Add, dim3(blockPerGrid), dim3(threadPerBlock), 0, 0, a.getDevicePtr(), b.getDevicePtr(), c.getDevicePtr(), N);
	c.recvDataFromDevice();
	for(int i = 0 ; i < N ; ++ i){
		std::cout << a[i] << " + " << b[i] << " = " << c[i] << " / " << a[i] + b[i] << std::endl;
	}

	return 0;
}
```

頭に```hip/hip_runtime.h```へのincludeがある事に注意してください、これがないと```blockDim.x```などが未定義になります。
HIP的には```hipThreadIdx_x```という名前があって、それが```hip/hip_runtime.h```で定義されているようです。

見るとほとんどCUDAで、```cudaMemcpy```が```hipMemcpy```になっているくらいです。
```hipLaunchKernelGGL```に関しては、最初の引数が実行関数名で、2〜5つめの引数がCUDAで言うところの```<<< ... >>>```で指定している部分です。
~~この部分あまり好きじゃないんだけどなあ~~
6個目以降の引数は1個目の引数で指定した関数の引数になっています。
実際にこれを起動してみるとAMD GPU上で配列の足し算がなされました。

ただし、こちらの環境では、なぜか**suになってコンパイル及び実行**をしなければデバイスが見つからないと言われ実行できませんでした。
何かシステム的な関係かもしれません(そもそもlibOpenCLが見つからんと言われるのも何か関係あるかしら？)。

# 結論
というわけで、動いちゃいました。
文法としてはほぼほぼCUDAで、CUDA慣れしていればはじめからHIPで書いてもまあそんなストレスにはならないのではないでしょうか。

hipifyは使っていないのでどの程度の精度で変換してくれるのかはよくわかりません。
~~だからCUDAをそのままコンパイルしてくれればいいのに~~。
自分で書いたSPHコードとか、FDPS( https://github.com/FDPS/FDPS )に入っている$N$体サンプルコードのCUDA版を、hipifyしてみると動くのかな？
試してみて動いたら記事を書こうかしらと思います。