OpenCLのオフラインコンパイル

はじめに

OpenCLといえば，カーネルのコードに以下の2つのコンパイル方式がある．

• オンラインコンパイル
• オフラインコンパイル

オンラインコンパイルは，実行時にOpenCLのカーネルコードを文字列として関数に渡し，プログラムオブジェクトを構築する手法である． 反対に，オフラインコンパイルは事前にOpenCLのカーネルコードをコンパイルし，コンパイル結果のバイナリを生成しておく． そして，生成したバイナリを実行時に読み込んで，プログラムオブジェクトを構築する手法である． オフラインコンパイルは事前にコンパイルを行う分，実行時のコンパイル時間を削減することができるわけだ．

オフラインコンパイル

オフラインコンパイルには，大別して2つの方針がある．

1. OpenCLのSDK付属のオフラインコンパイラを利用する
2. OpenCLのAPIを用いて，コンパイルプログラムを書く

前者はSDKが必須となるが，後者は不要となる． この記事では，OpenCLのAPIを用いて，オフラインコンパイラを書くことにする．

オフラインコンパイルを行うプログラムを作成する

clCreateProgramWithSource() ， clBuildProgram() により生成した cl_program から clGetProgramInfo() を用いることで，バイナリの情報およびバイナリそのものを取得する形になる． 具体的には以下のようになる． このコードは，0番目のプラットフォームID，0番目のデバイスIDを対象に，コマンドライン引数で指定したカーネルのソースコードのオフラインコンパイルを行うものだ．

• oclc.cpp

// oclc.cpp
// g++ -gnu++11 -O3 oclc.cpp -lOpenCL -o oclc
#include <iostream>
#include <cmath>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <fstream>
#include <memory>
#include <vector>

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


static constexpr cl_uint kNDefaultPlatformEntry = 16;
static constexpr cl_uint kNDefaultDeviceEntry = 16;


/*!
 * @brief プラットフォームIDを取得
 * @param [in] nPlatformEntry  取得するプラットフォームID数の上限
 * @return  プラットフォームIDを格納した std::vector
 */
static inline std::vector<cl_platform_id>
getPlatformIds(cl_uint nPlatformEntry = kNDefaultPlatformEntry)
{
  std::vector<cl_platform_id> platformIds(nPlatformEntry);
  cl_uint nPlatform;
  if (clGetPlatformIDs(nPlatformEntry, platformIds.data(), &nPlatform) != CL_SUCCESS) {
    std::cerr << "clGetPlatformIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  platformIds.resize(nPlatform);
  return platformIds;
}


/*!
 * @brief デバイスIDを取得
 * @param [in] platformId    デバイスIDの取得元のプラットフォームのID
 * @param [in] nDeviceEntry  取得するデバイスID数の上限
 * @param [in] deviceType    取得対象とするデバイス
 * @return デバイスIDを格納した std::vector
 */
static inline std::vector<cl_device_id>
getDeviceIds(const cl_platform_id& platformId, cl_uint nDeviceEntry = kNDefaultDeviceEntry, cl_int deviceType = CL_DEVICE_TYPE_DEFAULT)
{
  std::vector<cl_device_id> deviceIds(nDeviceEntry);
  cl_uint nDevice;
  if (clGetDeviceIDs(platformId, deviceType, nDeviceEntry, deviceIds.data(), &nDevice) != CL_SUCCESS) {
    std::cerr << "clGetDeviceIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  deviceIds.resize(nDevice);
  return deviceIds;
}


/*!
 * @brief 指定されたファイル名のファイルを読み込み，std::stringに格納して返却する
 * @param [in] filename  読み込むファイル名
 * @return  ファイルの内容を格納したstd::string
 */
static inline std::string
readSource(const std::string& filename) noexcept
{
  std::ifstream ifs(filename.c_str());
  if (!ifs.is_open()) {
    std::cerr << "Failed to open " << filename << std::endl;
    std::exit(EXIT_FAILURE);
  }
  return std::string((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());
}


/*!
 * @brief 指定された複数のファイル名のファイルを読み込み，std::stringに格納して返却する
 * @param [in] filenames  読み込み対象のファイル名を格納したstd::vector
 * @return  ファイルの内容を格納したstd:stringを格納したstd::vector
 */
static inline std::vector<std::string>
readSource(const std::vector<std::string>& filenames)
{
  std::vector<std::string> srcs(filenames.size());
  for (decltype(srcs)::size_type i = 0; i < srcs.size(); i++) {
    srcs[i] = readSource(filenames[i]);
  }
  return srcs;
}


/*!
 * @brief 指定されたファイル名から拡張子を除いた文字列を返却する
 * @param [in] filename  拡張子を取り除きたいファイル名
 * @return  拡張子を除いたファイル名
 */
static inline std::string
removeSuffix(const std::string& filename) noexcept
{
  return filename.substr(0, filename.find_last_of("."));
}


/*!
 * @brief このプログラムのエントリポイント
 * @param [in] argc  コマンドライン引数の数
 * @param [in] argv  コマンドライン引数
 * @return 終了ステータス
 */
int
main(int argc, char* argv[])
{
  // OpenCLのコンパイラに渡すオプション文字列
  // 今回はとりあえず空
  static const char kOptStr[] = "";

  std::vector<std::string> args(argc - 1);
  if (args.size() < 1) {
    std::cerr << "Please specify only one or more source file" << std::endl;
    return EXIT_FAILURE;
  }
  for (decltype(args)::size_type i = 0; i < args.size(); i++) {
    args[i] = std::string(argv[i + 1]);
  }

  // プラットフォームを取得
  std::vector<cl_platform_id> platformIds = getPlatformIds(1);

  // デバイスを取得
  std::vector<cl_device_id> deviceIds = getDeviceIds(platformIds[0], 1, CL_DEVICE_TYPE_DEFAULT);

  // コンテキスト生成
  cl_int errCode;
  std::unique_ptr<std::remove_pointer<cl_context>::type, decltype(&clReleaseContext)> context(
      clCreateContext(nullptr, 1, &deviceIds[0], nullptr, nullptr, &errCode), clReleaseContext);
  if (errCode != CL_SUCCESS) {
    std::cerr << "clCreateContext() failed" << std::endl;
    return EXIT_FAILURE;
  }

  // ソースコード読み込み
  std::vector<std::string> kernelSources = readSource(args);
  std::pair<std::vector<const char*>, std::vector<std::string::size_type> > kernelSourcePairs;
  kernelSourcePairs.first.reserve(kernelSources.size());
  kernelSourcePairs.second.reserve(kernelSources.size());
  for (const auto& kernelSource : kernelSources) {
    kernelSourcePairs.first.emplace_back(kernelSource.c_str());
    kernelSourcePairs.second.emplace_back(kernelSource.length());
  }

  // プログラム生成
  // 複数ソースファイルに対応
  std::unique_ptr<std::remove_pointer<cl_program>::type, decltype(&clReleaseProgram)> program(
      clCreateProgramWithSource(
        context.get(),
        static_cast<cl_uint>(kernelSourcePairs.first.size()),
        kernelSourcePairs.first.data(),
        kernelSourcePairs.second.data(),
        &errCode),
      clReleaseProgram);
  if (errCode != CL_SUCCESS) {
    std::cerr << "clCreateProgramWithSource() failed" << std::endl;
    return EXIT_FAILURE;
  }

  // カーネルソースコードのコンパイル
  switch (clBuildProgram(program.get(), 1, &deviceIds[0], kOptStr, nullptr, nullptr)) {
    case CL_SUCCESS:
      break;
    case CL_BUILD_PROGRAM_FAILURE:
      {
        // コンパイルエラーを表示
        std::array<char, 2048> buildLog;
        std::size_t logSize;
        clGetProgramBuildInfo(program.get(), deviceIds[0], CL_PROGRAM_BUILD_LOG, buildLog.size(), buildLog.data(), &logSize);
        std::cerr << "Compile error:\n" << buildLog.data() << std::endl;
      }
      break;
    case CL_INVALID_BUILD_OPTIONS:
      std::cerr << "Invalid option is specified" << std::endl;
      return EXIT_FAILURE;
    default:
      std::cerr << "clBuildProgram() failed" << std::endl;
      return EXIT_FAILURE;
  }

  // デバイス数を取得 (このプログラムでは1が返却されるはず)
  cl_uint nDevice;
  if (clGetProgramInfo(program.get(), CL_PROGRAM_NUM_DEVICES, sizeof(nDevice), &nDevice, nullptr) !=  CL_SUCCESS) {
    std::cerr << "clGetProgramInfo() failed" << std::endl;
  }

  // 各デバイス向けのコンパイル後のバイナリのサイズを取得
  std::unique_ptr<std::size_t[]> binSizes(new std::size_t[nDevice]);
  if (clGetProgramInfo(program.get(), CL_PROGRAM_BINARY_SIZES, sizeof(std::size_t) * nDevice, binSizes.get(), nullptr) != CL_SUCCESS) {
    std::cerr << "clGetProgramInfo() failed" << std::endl;
  }

  // コンパイル後のバイナリをコピー
  std::vector<std::unique_ptr<char> > bins(nDevice);
  for (std::size_t i = 0; i < nDevice; i++) {
    bins[i] = std::unique_ptr<char>(binSizes[i] == 0 ? nullptr : new char[binSizes[i]]);
  }
  if (clGetProgramInfo(program.get(), CL_PROGRAM_BINARIES, sizeof(char*) * nDevice, bins.data(), nullptr) != CL_SUCCESS) {
    std::cerr << "clGetProgramInfo() failed" << std::endl;
    return EXIT_FAILURE;
  }

  // コピーしたバイナリを全てファイルに出力
  std::string basename = removeSuffix(args[0]);
  for (std::size_t i = 0; i < nDevice; i++) {
    if (bins[i] == nullptr) {
      continue;
    }
    std::string filename = basename + ".bin";
    if (nDevice > 1) {
      filename += "." + std::to_string(i);
    }
    std::ofstream ofs(filename, std::ios::binary);
    if (ofs.is_open()) {
      ofs.write(bins[i].get(), binSizes[i]);
    } else {
      std::cerr << "Failed to open: " << filename << std::endl;
    }
  }

  return EXIT_SUCCESS;
}

コンパイルは以下のようにして行う．

$ g++ -gnu++11 -O3 oclc.cpp -lOpenCL -o oclc

次に，以下のようなカーネルのコードを用意する．

• kernel.cl

// kernel.cl
__kernel void
vecAdd(__global float* z, __global const float* x, __global const 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];
  }
}

そして，以下のようにしてカーネルのソースコードをコンパイルする．

$ ./oclc kernel.cl

これでカレントディレクトリに kernel.bin が生成されていれば，オフラインコンパイルは成功である． コンパイルエラーがある場合は，エラーメッセージを表示するようにしてある．

生成したカーネルのバイナリのテストには以下のコードを用いる．

// main.cpp
// g++ -std=gnu++11 -O3 main.cpp -lOpenCL -o 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


static constexpr cl_uint kNDefaultPlatformEntry = 16;
static constexpr cl_uint kNDefaultDeviceEntry = 16;




/*!
 * @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 プラットフォームIDを取得
 * @param [in] nPlatformEntry  取得するプラットフォームID数の上限
 * @return  プラットフォームIDを格納した std::vector
 */
static inline std::vector<cl_platform_id>
getPlatformIds(cl_uint nPlatformEntry = kNDefaultPlatformEntry)
{
  std::vector<cl_platform_id> platformIds(nPlatformEntry);
  cl_uint nPlatform;
  if (clGetPlatformIDs(nPlatformEntry, platformIds.data(), &nPlatform) != CL_SUCCESS) {
    std::cerr << "clGetPlatformIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  platformIds.resize(nPlatform);
  return platformIds;
}


/*!
 * @brief デバイスIDを取得
 * @param [in] platformId    デバイスIDの取得元のプラットフォームのID
 * @param [in] nDeviceEntry  取得するデバイスID数の上限
 * @param [in] deviceType    取得対象とするデバイス
 * @return デバイスIDを格納した std::vector
 */
static inline std::vector<cl_device_id>
getDeviceIds(const cl_platform_id& platformId, cl_uint nDeviceEntry = kNDefaultDeviceEntry, cl_int deviceType = CL_DEVICE_TYPE_DEFAULT)
{
  std::vector<cl_device_id> deviceIds(nDeviceEntry);
  cl_uint nDevice;
  if (clGetDeviceIDs(platformId, deviceType, nDeviceEntry, deviceIds.data(), &nDevice) != CL_SUCCESS) {
    std::cerr << "clGetDeviceIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  deviceIds.resize(nDevice);
  return deviceIds;
}



/*!
 * @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(int argc, char* argv[])
{
  static constexpr int ALIGN = 4096;
  static constexpr std::size_t N = 65536;

  if (argc < 2) {
    std::cerr << "Please specify only one or more source file" << std::endl;
    return EXIT_FAILURE;
  }

  // ホストのメモリを確保
  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);

  // プラットフォームを取得
  std::vector<cl_platform_id> platformIds = getPlatformIds(1);

  // デバイスを取得
  std::vector<cl_device_id> deviceIds = getDeviceIds(platformIds[0], 1, CL_DEVICE_TYPE_DEFAULT);

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

  // コマンドキューを生成
  std::unique_ptr<std::remove_pointer<cl_command_queue>::type, decltype(&clReleaseCommandQueue)> cmdQueue(
      clCreateCommandQueue(context.get(), deviceIds[0], 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(argv[1], std::ios::binary);
  if (!ifs.is_open()) {
    std::cerr << "Failed to kernel binary: " << argv[1] << std::endl;
    std::exit(EXIT_FAILURE);
  }
  std::string kernelBin((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());

  // プログラムオブジェクトの生成
  const unsigned char* kbin = reinterpret_cast<const unsigned char*>(kernelBin.c_str());
  std::size_t kbinSize = kernelBin.size();
  cl_int binStatus;
  std::unique_ptr<std::remove_pointer<cl_program>::type, decltype(&clReleaseProgram)> program(
      clCreateProgramWithBinary(context.get(), 1, &deviceIds[0], &kbinSize, &kbin, &binStatus, &errCode), clReleaseProgram);

  // カーネルソースコードのコンパイル (必要な環境もあるらしい?)
  // errCode = clBuildProgram(program.get(), 1, &deviceIds[0], 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;
}

実行は以下のようにする．

$ ./main kernel.bin

これで問題なく実行できればOKである．

雑感

個別のコンパイルプログラムを準備し，事前にオフラインコンパイルを行い，カーネルのバイナリを用意するのは正直なところどうなのかと感じた． プログラムの初回起動時に，オンラインコンパイルの結果を保存し，2回目以降にその結果を再利用する形にするのがよさそうだ．

個別のコンパイルプログラムは，OpenCLのプログラムのコンパイルエラーのチェック等に用いるとよいだろう．

もう少しまともな形として作るならば，以下のような形として作るとよさそうだ．

koturn/oclc

まとめ

OpenCLのオフラインコンパイルには，オンラインコンパイルの結果を保存する手法がある．