最近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が制定している、ヘテロジニアス計算機向け開発環境だそうな。
なんでも、Heterogeneous-compute Interface for Portabilityの略だとかなんだとか。
文法としてはほぼほぼ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へのパスも指定しなくちゃいけないですね。
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でよいみたいです。
# 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してみると動くのかな？
試してみて動いたら記事を書こうかしらと思います。