koturnの日記

普通の人です.ブログ上のコードはコピペ自由です.

MSYS2でOpenCLを使う

はじめに

MSYS2でOpenCLのプログラムをコンパイル&実行したかった. pacmanOpenCL関連のヘッダを導入することはできるが,OpenCLのインポートライブラリは導入することはできない. ここでは,MSYS2でOpenCLの環境を導入する一連の手順を紹介する.

OpenCLのヘッダを導入する

pacmanを用いると,OpenCLのヘッダを導入できる. x64ならば,

$ pacman -S mingw-w64-x86_64-opencl-headers

x86ならば,

$ pacman -S mingw-w64-i686-opencl-headers

とするとよい.

OpenCLのインポートライブラリを作成する

上記の手順では,ヘッダファイルしか導入できない. インポートライブラリは自分で作成する必要がある. この手順は,ここを参考にした.

  • x64環境
$ mkdir lib64 && gendef - /c/Windows/system32/OpenCL.dll > lib64/OpenCL.def
$ dlltool -l lib64/libOpenCL.a -d lib64/OpenCL.def -A -k
$ mkdir lib32 && gendef - /c/Windows/SysWOW64/OpenCL.dll > lib32/OpenCL.def
$ dlltool -l lib32/libOpenCL.a -d lib32/OpenCL.def -A -k

あとは必要に応じて, /usr/local/lib 下に置くなどするとよい.

コンパイルの確認

適当に以下のようなソースコードを準備する. (面倒なので, main() 関数でのエラー処理は一切行っていない.)

#include <cstdlib>
#include <fstream>
#include <iostream>
#include <memory>
#include <random>
#include <string>

#ifdef __APPLE__
#  include <OpenCL/opencl.h>
#else
#  include <CL/cl.h>
#endif


/*!
 * @brief アラインメントされたメモリを動的確保する関数
 * @param [in] size       確保するメモリサイズ (単位はbyte)
 * @param [in] alignment  アラインメント (2のべき乗を指定すること)
 * @return  アラインメントし,動的確保されたメモリ領域へのポインタ
 */
template<typename T = void*, typename std::enable_if<std::is_pointer<T>::value, std::nullptr_t>::type = nullptr>
static inline T
alignedMalloc(std::size_t size, std::size_t alignment) noexcept
{
#if defined(_MSC_VER) || defined(__MINGW32__)
  return reinterpret_cast<T>(_aligned_malloc(size, alignment));
#else
  void* p;
  return reinterpret_cast<T>(posix_memalign(&p, alignment, size) == 0 ? p : nullptr);
#endif  // defined(_MSC_VER) || defined(__MINGW32__)
}


/*!
 * @brief アラインメントされたメモリを解放する関数
 * @param [in] ptr  解放対象のメモリの先頭番地を指すポインタ
 */
static inline void
alignedFree(void* ptr) noexcept
{
#if defined(_MSC_VER) || defined(__MINGW32__)
  _aligned_free(ptr);
#else
  std::free(ptr);
#endif  // defined(_MSC_VER) || defined(__MINGW32__)
}


/*!
 * @brief std::unique_ptr で利用するアラインされたメモリ用のカスタムデリータ
 */
struct AlignedDeleter
{
  /*!
   * @brief デリート処理を行うオペレータ
   * @param [in,out] p  アラインメントされたメモリ領域へのポインタ
   */
  void
  operator()(void* p) const noexcept
  {
    alignedFree(p);
  }
};


/*!
 * @brief OpenCLのデバイス情報を表示する
 * @param [in] device  OpenCLのデバイスID
 */
static inline void
showDeviceInfo(cl_device_id device) noexcept
{
  char info[2048];
  std::size_t size;

  int err = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(info), info, &size);
  std::cout << "  Device: " << info << std::endl;
  err = clGetDeviceInfo(device, CL_DEVICE_VERSION, sizeof(info), info, &size);
  std::cout << "  CL_DEVICE_VERSION: " << info << std::endl;
}


/*!
 * @brief OpenCLのプラットフォーム情報を表示する
 * @param [in] platformId  OpenCLのプラットフォーム情報
 */
static inline void
showPlatformInfo(cl_platform_id platformId) noexcept
{
  static constexpr int MAX_DEVICE_IDS = 8;
  char info[2048];

  std::cout << "========== Platform Information ==========" << std::endl;
  std::size_t size;
  int err = clGetPlatformInfo(platformId, CL_PLATFORM_NAME, sizeof(info), info, &size);
  std::cout << "Platform: " << info << std::endl;
  err = clGetPlatformInfo(platformId, CL_PLATFORM_VERSION, sizeof(info), info, &size);
  std::cout << "CL_PLATFORM_VERSION: " << info << std::endl;

  cl_device_id devices[MAX_DEVICE_IDS]; cl_uint nDevice;
  err = clGetDeviceIDs(platformId, CL_DEVICE_TYPE_ALL, MAX_DEVICE_IDS, devices, &nDevice);
  if (err != CL_SUCCESS) {
    std::cerr << "Error (clGetDeviceIDs): " << err << std::endl;
    std::exit(EXIT_FAILURE);
  }
  for (cl_uint j = 0; j < nDevice; j++) {
    showDeviceInfo(devices[j]);
  }
  std::cout << "==========================================" << std::endl;
}


/*!
 * @brief カーネル関数へ引数をまとめてセットする関数の実態
 * @param [in] kernel  OpenCLカーネルオブジェクト
 * @param [in] idx     セットする引数のインデックス
 * @param [in] first   セットする引数.可変パラメータから1つだけ取り出したもの
 * @param [in] rest    残りの引数
 * @return OpenCLのエラーコード.エラーが出た時点でエラーコードを返却する.
 */
template<typename First, typename... Rest>
static inline cl_uint
setKernelArgsImpl(const cl_kernel& kernel, int idx, const First& first, const Rest&... rest) noexcept
{
  cl_uint errCode = clSetKernelArg(kernel, idx, sizeof(first), &first);
  return errCode == CL_SUCCESS ? setKernelArgsImpl(kernel, idx + 1, rest...) : errCode;
}


/*!
 * @brief カーネル関数へ最後の引数をセットする
 * @param [in] kernel  OpenCLカーネルオブジェクト
 * @param [in] idx     引数のインデックス
 * @param [in] last    最後の引数
 * @return  OpenCLのエラーコード
 */
template<typename Last>
static inline cl_uint
setKernelArgsImpl(const cl_kernel& kernel, int idx, const Last& last) noexcept
{
  return clSetKernelArg(kernel, idx, sizeof(last), &last);
}


/*!
 * @brief カーネル関数へ引数をまとめてセットする
 * @param [in] kernel  OpenCLカーネルオブジェクト
 * @param [in] args    セットする引数群
 * @return OpenCLのエラーコード.エラーが出た時点でエラーコードを返却する.
 */
template<typename... Args>
static inline cl_uint
setKernelArgs(const cl_kernel& kernel, const Args&... args) noexcept
{
  return setKernelArgsImpl(kernel, 0, args...);
}




/*!
 * @brief このプログラムのエントリポイント
 * @return 終了ステータス
 */
int
main()
{
  static constexpr int ALIGN = 4096;
  static constexpr std::size_t N = 65536;
  static const char KERNEL_FILENAME[] = "kernel.cl";

  // ホストのメモリを確保
  std::unique_ptr<float[], AlignedDeleter> hostX(alignedMalloc<float*>(N * sizeof(float), ALIGN));
  std::unique_ptr<float[], AlignedDeleter> hostY(alignedMalloc<float*>(N * sizeof(float), ALIGN));
  std::unique_ptr<float[], AlignedDeleter> hostZ(alignedMalloc<float*>(N * sizeof(float), ALIGN));

  // 初期化
  std::mt19937 mt((std::random_device())());
  for (std::size_t i = 0; i < N; i++) {
    hostX[i] = static_cast<float>(mt());
    hostY[i] = static_cast<float>(mt());
  }
  std::fill_n(hostZ.get(), N, 0.0f);

  // プラットフォームを取得
  cl_platform_id platformId;
  cl_uint nPlatform;
  cl_int errCode = clGetPlatformIDs(1, &platformId, &nPlatform);
  showPlatformInfo(platformId);

  // デバイス情報を取得
  cl_device_id deviceId;
  cl_uint nDevice;
  errCode = clGetDeviceIDs(platformId, CL_DEVICE_TYPE_DEFAULT, 1, &deviceId, &nDevice);

  // コンテキストを生成
  std::unique_ptr<std::remove_pointer<cl_context>::type, decltype(&clReleaseContext)> context(
      clCreateContext(nullptr, 1, &deviceId, nullptr, nullptr, &errCode), clReleaseContext);

  // コマンドキューを生成
  std::unique_ptr<std::remove_pointer<cl_command_queue>::type, decltype(&clReleaseCommandQueue)> cmdQueue(
      clCreateCommandQueue(context.get(), deviceId, 0, &errCode), clReleaseCommandQueue);

  // デバイスが用いるメモリオブジェクトの生成
  std::unique_ptr<std::remove_pointer<cl_mem>::type, decltype(&clReleaseMemObject)> deviceX(
      clCreateBuffer(context.get(), CL_MEM_READ_WRITE, N * sizeof(float), nullptr, &errCode), clReleaseMemObject);
  std::unique_ptr<std::remove_pointer<cl_mem>::type, decltype(&clReleaseMemObject)> deviceY(
      clCreateBuffer(context.get(), CL_MEM_READ_WRITE, N * sizeof(float), nullptr, &errCode), clReleaseMemObject);
  std::unique_ptr<std::remove_pointer<cl_mem>::type, decltype(&clReleaseMemObject)> deviceZ(
      clCreateBuffer(context.get(), CL_MEM_READ_WRITE, N * sizeof(float), nullptr, &errCode), clReleaseMemObject);

  // ホストのメモリをデバイスのメモリに転送
  errCode = clEnqueueWriteBuffer(cmdQueue.get(), deviceX.get(), CL_TRUE, 0, N * sizeof(float), hostX.get(), 0, nullptr, nullptr);
  errCode = clEnqueueWriteBuffer(cmdQueue.get(), deviceY.get(), CL_TRUE, 0, N * sizeof(float), hostY.get(), 0, nullptr, nullptr);
  errCode = clEnqueueWriteBuffer(cmdQueue.get(), deviceZ.get(), CL_TRUE, 0, N * sizeof(float), hostZ.get(), 0, nullptr, nullptr);

  // カーネルのソースコードを文字列として取得
  std::ifstream ifs(KERNEL_FILENAME);
  if (!ifs.is_open()) {
    std::cerr << "Failed to read kernel program: " << KERNEL_FILENAME << std::endl;
    return EXIT_FAILURE;
  }
  std::string kernelSource((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());

  // プログラムオブジェクトの生成
  const char* ksrc = kernelSource.c_str();
  std::string::size_type srcSize = kernelSource.length();
  std::unique_ptr<std::remove_pointer<cl_program>::type, decltype(&clReleaseProgram)> program(
      clCreateProgramWithSource(context.get(), 1, &ksrc, &srcSize, &errCode), clReleaseProgram);

  // カーネルソースコードのコンパイル
  errCode = clBuildProgram(program.get(), 1, &deviceId, nullptr, nullptr, nullptr);

  // カーネルオブジェクトの生成
  std::unique_ptr<std::remove_pointer<cl_kernel>::type, decltype(&clReleaseKernel)> kernel(
      clCreateKernel(program.get(), "vecAdd", &errCode), clReleaseKernel);

  // カーネル関数に引数を渡す
  errCode = setKernelArgs(kernel.get(), deviceZ.get(), deviceX.get(), deviceY.get(), static_cast<int>(N));

  // カーネルプログラムの実行
  errCode = clEnqueueTask(cmdQueue.get(), kernel.get(), 0, nullptr, nullptr);

  // 終了待機等
  errCode = clFlush(cmdQueue.get());
  errCode = clFinish(cmdQueue.get());

  // 実行結果をデバイスからホストへコピー
  errCode = clEnqueueReadBuffer(cmdQueue.get(), deviceZ.get(), CL_TRUE, 0, N * sizeof(float), hostZ.get(), 0, nullptr, nullptr);

  // 計算結果の確認
  for (std::size_t i = 0; i < N; i++) {
    if (std::abs(hostX[i] + hostY[i] - hostZ[i]) > 1.0e-5) {
      std::cerr << "Result verification failed at element " << i << "!" << std::endl;
      return EXIT_FAILURE;
    }
  }
  std::cout << "Test PASSED" << std::endl;

  return EXIT_SUCCESS;
}

そして,カーネルソースコードとして以下のものを準備する. ファイル名は kernel.cl とする.

__kernel void
vecAdd(__global float* z, __global float* x, __global float* y, int n)
{
  const int para = 4;
  const int end = (n / para) * para;

  for (int i = 0; i < end; i += para) {
    float4 vtmp = vload4(0, x + i) + vload4(0, y + i);
    vstore4(vtmp, 0, z + i);
  }

  for (int i = end; i < n; i++) {
    z[i] = x[i] + y[i];
  }
}

コンパイルと実行は以下の通り. これで問題なくコンパイル&実行できればOKである.

$ g++ -std=gnu++11 main.cpp -Llib64 -lOpenCL -o main

余談

カーネル関数へ引数を渡す関数: clSetKernelArg() は,引数のインデックスと引数へのポインタを渡す形となっており,非常にダサい. また,カーネル関数の引数の数だけ呼び出さなければならないのも大変だ.

int n = static_cast<int>(N);
clSetKernelArg(kernel.get(), 0, sizeof(deviceZ.get()), &deviceZ.get());
clSetKernelArg(kernel.get(), 1, sizeof(deviceX.get()), &deviceX.get());
clSetKernelArg(kernel.get(), 2, sizeof(deviceY.get()), &deviceY.get());
clSetKernelArg(kernel.get(), 3, sizeof(n), &n);

うまく可変引数テンプレートを用いれば,このダサさを解消できると考え,以下のような関数の実装を行った.

template<typename First, typename... Rest>
static inline cl_uint
setKernelArgsImpl(const cl_kernel& kernel, int idx, const First& first, const Rest&... rest) noexcept
{
  cl_uint errCode = clSetKernelArg(kernel, idx, sizeof(first), &first);
  return errCode == CL_SUCCESS ? setKernelArgsImpl(kernel, idx + 1, rest...) : errCode;
}


template<typename Last>
static inline cl_uint
setKernelArgsImpl(const cl_kernel& kernel, int idx, const Last& last) noexcept
{
  return clSetKernelArg(kernel, idx, sizeof(last), &last);
}


template<typename... Args>
static inline cl_uint
setKernelArgs(const cl_kernel& kernel, const Args&... args) noexcept
{
  return setKernelArgsImpl(kernel, 0, args...);
}

実装としては,ユーザが呼び出しを行う setKernelArgs() と実際の処理を担当する setKernelArgsImpl() の2つに分けている. setKernelArgsImpl() は,可変テンプレートのアンパックとインデックスのインクリメントを行う. 引数のサイズは sizeof 演算子で取得するので,引数の型はちゃんと意識する必要がある.

呼び出し側は以下のようになる. 呼び出し側の見た目としては,

  • インデックスが必要ない
  • カーネル関数の引数をポインタとして渡す必要がない
  • 引数のサイズを渡す必要がない

ため,非常にスッキリしていると思う.

setKernelArgs(kernel.get(), deviceZ.get(), deviceX.get(), deviceY.get(), static_cast<int>(N));

C++11様々である.

まとめ

MSYS2でOpenCLを使用するには,pacmanでOepnCL関連のヘッダのインストールを行い,WindowsOpenCLのDLLからインポートライブラリを作成する必要がある.

参考文献