はじめに
MSYS2でOpenCLのプログラムをコンパイル&実行したかった. pacmanでOpenCL関連のヘッダを導入することはできるが,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
- x86環境
$ 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関連のヘッダのインストールを行い,WindowsのOpenCLのDLLからインポートライブラリを作成する必要がある.