OpenCL をさわってみる ver.2
カーネルコード
/*-------------------- kernel.cl --------------------*/ __kernel void vec_add(size_t n, __global float* z, __global float* x, __global float* y){ int gid = get_global_id(0); if(gid < n){ z[gid] = x[gid] + y[gid]; } }
カーネルコードのコンパイル方法
- i386 環境向け
$ /System/Library/Frameworks/OpenCL.framework/Libraries/openclc -x cl \ -triple i386-applecl-darwin -emit-llvm-bc kernel.cl -o kernel.cpu32.bc
$ /System/Library/Frameworks/OpenCL.framework/Libraries/openclc -x cl \
-triple x86_64-applecl-darwin -emit-llvm-bc kernel.cl -o kernel.cpu64.bc
- GPU 向け
$ /System/Library/Frameworks/OpenCL.framework/Libraries/openclc -x cl \ -triple gpu_32-applecl-darwin -emit-llvm-bc kernel.cl -o kernel.gpu32.bc
ソースコード
/*-------------------- opencl_test2.cpp --------------------*/ #include <iostream> #include <cstdio> #include <cstdlib> #include <cstring> #include <numeric> #include <sys/time.h> #ifdef __APPLE__ #include <OpenCL/opencl.h> #else #include <CL/cl.h> #endif #define MAX_BINARY_SIZE (0x100000) int main(void){ const char filename[] = "./kernel.cpu64.bc"; std::size_t const n = 10; cl_platform_id platform_id = NULL; cl_device_id device_id = NULL; cl_context context = NULL; cl_command_queue command_queue = NULL; cl_program program = NULL; cl_kernel kernel = NULL; cl_uint ret_num_devices; cl_uint ret_num_platforms; cl_int ret; cl_int binary_status; cl_mem x_dev, y_dev, z_dev;//for device memory float* x = (float*)malloc(sizeof(float) * n); float* y = (float*)malloc(sizeof(float) * n); float* z = (float*)malloc(sizeof(float) * n); srand( static_cast<unsigned>(time(NULL)) ); for(std::size_t i=0; i<n; ++i) x[i] = static_cast<float>( rand() ) / RAND_MAX; for(std::size_t i=0; i<n; ++i) y[i] = static_cast<float>( rand() ) / RAND_MAX; memset(z, 0x00, sizeof(float) * n); //Read a Kernel code FILE* fp = fopen(filename, "r"); if(!fp){ std::cerr << "Failed to load kernel" << std::endl; return -1; } char* binary_buf = (char*)malloc(MAX_BINARY_SIZE); std::size_t binary_size = fread(binary_buf, 1, MAX_BINARY_SIZE, fp); fclose(fp); //Get the platforms ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms); //Get the device ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices); //Create context context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret); //Create Command Queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret); //------------------------------------------------------------// //Create Program Object //program = clCreateProgramWithSource(context, 1, (const char**)&source_str, // (const size_t*)&source_size, &ret); program = clCreateProgramWithBinary(context, 1, &device_id, (const size_t*)&binary_size, (const unsigned char**)&binary_buf, &binary_status, &ret); //Compile of kernel's source code #ifdef __APPLE__ ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); #endif //------------------------------------------------------------// //Create Kernel object kernel = clCreateKernel(program, "vec_add", &ret); if(ret != CL_SUCCESS){ std::cerr << "Error - clCreateKernel" << std::endl; return -1; } //Create memory object x_dev = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret); y_dev = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret); z_dev = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(float), NULL, &ret); //memory copy host to device ret = clEnqueueWriteBuffer(command_queue, x_dev, CL_TRUE, 0, n * sizeof(float), x, 0, NULL, NULL); ret = clEnqueueWriteBuffer(command_queue, y_dev, CL_TRUE, 0, n * sizeof(float), y, 0, NULL, NULL); ret = clEnqueueWriteBuffer(command_queue, z_dev, CL_TRUE, 0, n * sizeof(float), z, 0, NULL, NULL); //Set args for kernel object ret = clSetKernelArg(kernel, 0, sizeof(std::size_t), &n); ret += clSetKernelArg(kernel, 1, sizeof(cl_mem), &z_dev); ret += clSetKernelArg(kernel, 2, sizeof(cl_mem), &x_dev); ret += clSetKernelArg(kernel, 3, sizeof(cl_mem), &y_dev); size_t global_work_size[3] = {n, 0, 0}; size_t local_work_size[3] = {n, 0, 0}; //Execute Kernel Program ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); //memory copy device to host ret = clEnqueueReadBuffer(command_queue, z_dev, CL_TRUE, 0, n * sizeof(float), z, 0, NULL, NULL); for(std::size_t i=0; i<n; ++i){ std::cout << x[i] << " " << y[i] << " = " << z[i]<< std::endl; } //memory free ret = clFlush(command_queue); ret = clFinish(command_queue); ret = clReleaseKernel(kernel); ret = clReleaseProgram(program); ret = clReleaseCommandQueue(command_queue); ret = clReleaseContext(context); ret = clReleaseMemObject(x_dev); ret = clReleaseMemObject(y_dev); ret = clReleaseMemObject(z_dev); free(x); free(y); free(z); free(binary_buf); return 0; }
ソースコードのコンパイル方法
$ g++ opencl_test2.cpp -framework opencl -arch x86_64
arch で実行環境を指定します．確認した実行環境だと以下のものがあるようです．
所感
オンラインコンパイル版とオフラインコンパイル版の違いは以下の通り．
- カーネルコードをコンパイルしておく
- カーネルをバイナリとして読み込む
- clCreateProgramWithSource と clCreateProgramWithBinary の違い
- オフラインコンパイル版は（本来）clBuildProgram が不要
「OpenCL 入門」によるとオフラインコンパイルの際は 「clBuildProgram()」は不要となっていますが，本環境では clBuildProgram() がないと実行できませんでした（clCreateKernel() がエラーコードを返す）．Web を探してみると似たような（？）情報を発見
知らんうちにOpenCLのオフラインコンパイルのが出てた。：毎回サブタイトル考えるの面倒いなぁ：So-net blog．
また，上記参考のサンプルコードでは clBuildProgram を使用しており，以下のコメントが記載されてます（OpenCLOfflineCompile の main.c より一部抜粋）．
... // The above tells OpenCL how to locate the intermediate bitcode, but we // still must build the program to produce executable bits for our // *specific* device. This transforms gpu32 bitcode into actual executable // bits for an AMD or Intel compute device (for example). err = clBuildProgram(program, 1, &device, NULL, NULL, NULL); ...
・・・現状は”オフラインコンパイル”も clBuildProgram でビルドする必要があるぽいですね（ただし，Apple OpenCL 以外は未確認のため不明）
