MIOpen HIP backend 入門

MIOpen(NVIDIA cuDNN の AMD 版)を HIP backend で使います.

OpenCL 版はこちら

MIOpen(OpenCL backend) API の入門
https://qiita.com/syoyo/items/2a2038d665499f1f6df1

HIP とは?

基本的には AMD 版の NVIDIA CUDA っぽいものです.
CUDA っぽい HIP API(?) が提供されおり(cudaMalloc に対応して, hipMalloc とか), この API を使ってプログラムを記述しておきます.
HIP のコードは, バックエンドを hcc(AMD GPU 用), nvcc(NVIDIA CUDA 用) で切り替えて使うことができます.

ただ, HIP を使って CUDA と ROCm でコードを統一するというのはなかなかのいばらの道かなと思います.
(CUDA でバージョンが上がったら HIP も追随していかなくてはならない).

ついでに, CUDA で動かそうと思っても, 開発の時点では CUDA SDK と HIP SDK(rocm-dev)の両方をインストールしなければならず, 面倒です.

NVIDIA GPU で動かすものは CUDA で開発し,
AMD GPU 用には HIP で開発と, 割り切って開発したほうが良い気がします.
(ROCm PyTorch が CUDA コードを HIP にスクリプトで置き換えてコンパイルするスタイルであるが, いろいろ大変なことをしている印象があります)

hcc とは

HIP が hcc backend の場合, C++/HIP コード は hcc の言語モードで変換され, バイナリ(ELF)になります. (clang で -x hc とかでコンパイルされる)
ちなみに, HCC 自体は deprecated になっています.

正確には, HCC を構成する要素のうち, C++AMP だけのようではあります: https://github.com/RadeonOpenCompute/hcc/wiki

OpenCL kernel

MIOpen では, カーネル自体は OpenCL で記述されています.
HIP backend では, MIOpen のライブラリビルド時は OpenCL コードは clang-ocl を経由してカーネルをコンパイルしています.
(/opt/rocm/bin/clang-ocl も, これは bash スクリプトで内部的には /opt/rocm/opencl の clang を呼んでいる)

コンパイラについて

以下は hip(AMD GPU)バックエンドを使う前提です.

ROCm SDK には hipcc と hcc などいろいろコンパイラがありますが, hipcc を使います.
(hcc だとエラーになるかも).

hipcc では, ファイルの拡張子が合っていないと, hcc の設定がうまく効かないです.
コンパイルがうまくいかなかったら, 明示的に HIP のコードであることを示す -x hip を指定するとよいかもしれません.

ちなみに, hipcc は Perl(!) のスクリプトで, 中身としては hc 拡張された clang を呼んでいます.
(hcc はまた別物っぽい模様)

hipcc でうまくコンパイルできないときは, このスクリプトの中身を見たりして引数がどうなっているかデバッグするのがよいかもです.

MIOpen HIP と MIOpen OpenCL どちらを使う?

とりあえず MIOpen 動かしたい, 複数 GPU との collective 通信ライブラリ RCCL など, HIP(HCC)しか提供されていないライブラリや機能を使いたいなら HIP,
将来を見据え機械学習アプリをポータブルに開発して動かしていきたいなら OpenCL, というところでしょうか.
(OpenCL 版だと, いろいろ頑張れば MIOpen が macOS とか Windows, Intel GPU でも動くはず?)

セットアップ

Ubuntu 16.04 or 18.04 を想定します.

を参考にし,
ROCm 開発環境(sudo apt install rocm-dev) が入っているものとします.

MIOpen HIP backend が入っていなかったら, インストールしときます.

$ sudo apt install miopen-hip

サンプルコード

メモリ転送あたりを hipMalloc, hipMemcpy で行うところが異なるくらいで, 基本的には OpenCL 版と同じです.

# include <array>
# include <cstdio>
# include <cstdlib>
# include <iostream>
# include <numeric>
# include <vector>

# include <miopen/miopen.h>
# include <miopen/version.h>

# define CHECK_MIOPEN(cmd)                                                      \
  {                                                                            \
    miopenStatus_t err = (cmd);                                                \
    if (err != miopenStatusSuccess) {                                          \
      fprintf(stderr, "error: '%s'(%d) at %s:%d\n", miopenGetErrorString(err), \
              err, __FILE__, __LINE__);                                        \
      exit(EXIT_FAILURE);                                                      \
    }                                                                          \
  }

# define CHECK_HIP(cmd)                                                      \
  {                                                                         \
    hipError_t hip_error = (cmd);                                           \
    if (hip_error != hipSuccess) {                                          \
      fprintf(stderr, "error: '%s'(%d) at %s:%d\n",                         \
              hipGetErrorString(hip_error), hip_error, __FILE__, __LINE__); \
      exit(EXIT_FAILURE);                                                   \
    }                                                                       \
  }

int main(int argc, char **argv) {
  std::cout << "MIOPEN_VERSION_MAJOR:" << MIOPEN_VERSION_MAJOR << std::endl;
  std::cout << "MIOPEN_VERSION_MINOR:" << MIOPEN_VERSION_MINOR << std::endl;
  std::cout << "MIOPEN_VERSION_PATCH:" << MIOPEN_VERSION_PATCH << std::endl;

  miopenHandle_t handle;
  CHECK_MIOPEN(miopenCreate(&handle));

  CHECK_MIOPEN(miopenEnableProfiling(handle, true));

  // input tensor
  const int in_shape[4] = {1, 1, 5, 5};  // NCHW

  miopenTensorDescriptor_t in_desc;
  CHECK_MIOPEN(miopenCreateTensorDescriptor(&in_desc));
  CHECK_MIOPEN(miopenSet4dTensorDescriptor(in_desc, miopenFloat, in_shape[0],
                                           in_shape[1], in_shape[2],
                                           in_shape[3]));  // NCHW

  // filter tensor
  const int filt_shape[4] = {1, 1, 2, 2};  // KCHW

  miopenTensorDescriptor_t filt_desc;
  CHECK_MIOPEN(miopenCreateTensorDescriptor(&filt_desc));
  CHECK_MIOPEN(miopenSet4dTensorDescriptor(filt_desc, miopenFloat,
                                           filt_shape[0], filt_shape[1],
                                           filt_shape[2], filt_shape[3]));

  miopenConvolutionDescriptor_t conv_desc;
  CHECK_MIOPEN(miopenCreateConvolutionDescriptor(&conv_desc));

  const int pad_h = 1;
  const int pad_w = 1;
  const int stride_h = 1;
  const int stride_w = 1;
  const int dilation_h = 1;
  const int dilation_w = 1;

  CHECK_MIOPEN(miopenInitConvolutionDescriptor(conv_desc, miopenConvolution,
                                               pad_h, pad_w, stride_h, stride_w,
                                               dilation_h, dilation_w));

  // output
  int out_shape[4];  // NCHW

  CHECK_MIOPEN(miopenGetConvolutionForwardOutputDim(
      conv_desc, in_desc, filt_desc, &out_shape[0], &out_shape[1],
      &out_shape[2], &out_shape[3]));

  miopenTensorDescriptor_t out_desc;
  CHECK_MIOPEN(miopenCreateTensorDescriptor(&out_desc));
  CHECK_MIOPEN(miopenSet4dTensorDescriptor(out_desc, miopenFloat, out_shape[0],
                                           out_shape[1], out_shape[2],
                                           out_shape[3]));

  // workspace
  size_t ws_size = 0;
  CHECK_MIOPEN(miopenConvolutionForwardGetWorkSpaceSize(
      handle, /* w */ filt_desc, /* x */ in_desc, conv_desc, /* y */ out_desc,
      &ws_size));
  std::cout << "ws_size = " << ws_size << "\n";

  size_t in_data_size = in_shape[0] * in_shape[1] * in_shape[2] * in_shape[3];
  size_t filt_data_size =
      filt_shape[0] * filt_shape[1] * filt_shape[2] * filt_shape[3];
  size_t out_data_size =
      out_shape[0] * out_shape[1] * out_shape[2] * out_shape[3];

  float *in_data = nullptr;
  CHECK_HIP(hipMalloc(&in_data, in_data_size * sizeof(float)));

  float *filt_data = nullptr;
  CHECK_HIP(hipMalloc(&filt_data, filt_data_size * sizeof(float)));

  float *out_data = nullptr;
  CHECK_HIP(hipMalloc(&out_data, out_data_size * sizeof(float)));

  float *ws_data = nullptr;
  CHECK_HIP(hipMalloc(&ws_data, ws_size));

  // fill with dummy data.
  std::vector<float> in_buf(in_data_size);
  std::iota(in_buf.begin(), in_buf.end(), 0);
  CHECK_HIP(hipMemcpy(in_data, in_buf.data(), in_data_size * sizeof(float),
                      hipMemcpyHostToDevice));

  std::vector<float> filt_buf(filt_data_size);
  std::iota(filt_buf.begin(), filt_buf.end(), 0);
  CHECK_HIP(hipMemcpy(filt_data, filt_buf.data(),
                      filt_data_size * sizeof(float), hipMemcpyHostToDevice));

  miopenConvAlgoPerf_t perf{};
  int algo_count = 0;
  bool exhaustive_search = false;
  std::cout << "find conv algo" << std::endl;
  CHECK_MIOPEN(miopenFindConvolutionForwardAlgorithm(
      handle, in_desc,
      /* ptr */ in_data, filt_desc, /* ptr */ filt_data, conv_desc, out_desc,
      /* ptr */ out_data, /* req algos*/ 1, &algo_count, &perf,
      /* ptr */ ws_data, ws_size, exhaustive_search));

  float alpha = 1.0f;
  float beta = 0.0f;
  CHECK_MIOPEN(miopenConvolutionForward(handle, &alpha, in_desc,
                                        /* ptr */ in_data, filt_desc,
                                        /* ptr */ filt_data, conv_desc,
                                        perf.fwd_algo, &beta, out_desc,
                                        out_data, ws_data, ws_size));

  float time = -1.0f;
  CHECK_MIOPEN(miopenGetKernelTime(handle, &time));
  std::cout << "time : " << time << "\n";

  std::vector<float> out_buf(out_data_size);
  CHECK_HIP(hipMemcpy(out_buf.data(), out_data, out_data_size * sizeof(float),
                      hipMemcpyDeviceToHost));

  for (size_t i = 0; i < out_buf.size(); i++) {
    std::cout << "[" << i << "] = " << out_buf[i] << "\n";
  }

  CHECK_HIP(hipFree(ws_data));
  CHECK_HIP(hipFree(out_data));
  CHECK_HIP(hipFree(filt_data));
  CHECK_HIP(hipFree(in_data));

  CHECK_MIOPEN(miopenDestroyTensorDescriptor(out_desc));
  CHECK_MIOPEN(miopenDestroyTensorDescriptor(in_desc));
  CHECK_MIOPEN(miopenDestroyTensorDescriptor(filt_desc));

  CHECK_MIOPEN(miopenDestroyConvolutionDescriptor(conv_desc));

  CHECK_MIOPEN(miopenDestroy(handle));

  return EXIT_SUCCESS;
}

$ /opt/rocm/bin/hipcc -I/opt/rocm/miopen/include main.cc -L/opt/rocm/miopen/lib -lMIOpen

としてコンパイルし, 実行します.

MIOPEN_VERSION_MAJOR:2
MIOPEN_VERSION_MINOR:0
MIOPEN_VERSION_PATCH:1
ws_size = 576
find conv algo
time : 0.010371
[0] = 0
[1] = 3
[2] = 8
[3] = 13
[4] = 18
[5] = 8
[6] = 15
[7] = 29
[8] = 35
[9] = 41
[10] = 47
[11] = 18
[12] = 35
[13] = 59
[14] = 65
[15] = 71
[16] = 77
[17] = 28
[18] = 55
[19] = 89
[20] = 95
[21] = 101
[22] = 107
[23] = 38
[24] = 75
[25] = 119
[26] = 125
[27] = 131
[28] = 137
[29] = 48
[30] = 20
[31] = 21
[32] = 22
[33] = 23
[34] = 24
[35] = 0

Voila!

懸念点

HIP は, NVCC 同様コンパイルが劇重になるので, HIP で大きめのプロジェクトを開発するとコンパイルの遅さに耐えられないかもしれません
(NVRTC みたいに動的 JIT コンパイルして事前コンパイルの時間を減らす, というは HIP には無いっぽい模様?)

TODO

• hipcc のコンパイル時間を減らす(ターゲットアーキテクチャを絞れば早くなりそう?)
• 複数 GPU で動かす方法をさがす(miopenSetStream?)