はじめに
OpenCLといえば，カーネルのコードに以下の2つのコンパイル方式がある．
オンラインコンパイルは，実行時にOpenCLのカーネルコードを文字列として関数に渡し，プログラムオブジェクトを構築する手法である．
反対に，オフラインコンパイルは事前にOpenCLのカーネルコードをコンパイルし，コンパイル結果のバイナリを生成しておく．
そして，生成したバイナリを実行時に読み込んで，プログラムオブジェクトを構築する手法である．
オフラインコンパイルは事前にコンパイルを行う分，実行時のコンパイル時間を削減することができるわけだ．
オフラインコンパイルには，大別して2つの方針がある．
- OpenCLのSDK付属のオフラインコンパイラを利用する
- OpenCLのAPIを用いて，コンパイルプログラムを書く
前者はSDKが必須となるが，後者は不要となる．
この記事では，OpenCLのAPIを用いて，オフラインコンパイラを書くことにする．
オフラインコンパイルを行うプログラムを作成する
clCreateProgramWithSource() ， clBuildProgram() により生成した cl_program から clGetProgramInfo() を用いることで，バイナリの情報およびバイナリそのものを取得する形になる．
具体的には以下のようになる．
このコードは，0番目のプラットフォームID，0番目のデバイスIDを対象に，コマンドライン引数で指定したカーネルのソースコードのオフラインコンパイルを行うものだ．
#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;
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;
}
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;
}
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>());
}
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;
}
static inline std::string
removeSuffix(const std::string& filename) noexcept
{
return filename.substr(0, filename.find_last_of("."));
}
int
main(int argc, char* argv[])
{
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;
}
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 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 が生成されていれば，オフラインコンパイルは成功である．
コンパイルエラーがある場合は，エラーメッセージを表示するようにしてある．
生成したカーネルのバイナリのテストには以下のコードを用いる．
#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;
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 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;
}
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;
}
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(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);
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のオフラインコンパイルには，オンラインコンパイルの結果を保存する手法がある．
