はじめに
MSYS2でOpenCLのプログラムをコンパイル&実行したかった．
pacmanでOpenCL関連のヘッダを導入することはできるが，OpenCLのインポートライブラリは導入することはできない．
ここでは，MSYS2でOpenCLの環境を導入する一連の手順を紹介する．
pacmanを用いると，OpenCLのヘッダを導入できる．
x64ならば，
$ pacman -S mingw-w64-x86_64-opencl-headers
x86ならば，
$ pacman -S mingw-w64-i686-opencl-headers
とするとよい．
OpenCLのインポートライブラリを作成する
上記の手順では，ヘッダファイルしか導入できない．
インポートライブラリは自分で作成する必要がある．
この手順は，ここを参考にした．
$ 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
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
}
static inline void
alignedFree(void* ptr) noexcept
{
#if defined(_MSC_VER) || defined(__MINGW32__)
_aligned_free(ptr);
#else
std::free(ptr);
#endif
}
struct AlignedDeleter
{
void
operator()(void* p) const noexcept
{
alignedFree(p);
}
};
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;
}
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;
}
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...);
}
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からインポートライブラリを作成する必要がある．
参考文献
