読者です 読者をやめる 読者になる 読者になる

OpenCL をさわってみる ver.2

主旨

OpenCL をさわってみる - kawa0810の日記 のオフラインコンパイル版です.
OpenCL の環境は Apple OpenCL (Mac 環境) を想定しています.

参考:Mac Developer Library

カーネルコード

/*--------------------
  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];
  }
}

カーネルコードのコンパイル方法

$ /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
$ /System/Library/Frameworks/OpenCL.framework/Libraries/openclc -x cl \
   -triple gpu_32-applecl-darwin -emit-llvm-bc kernel.cl -o kernel.gpu32.bc

基本的には -triple でカーネルコードの実行環境を決定するようです.

ソースコード

/*--------------------
  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 で実行環境を指定します.確認した実行環境だと以下のものがあるようです.

所感

オンラインコンパイル版とオフラインコンパイル版の違いは以下の通り.

  1. カーネルコードをコンパイルしておく
  2. カーネルをバイナリとして読み込む
  3. clCreateProgramWithSource と clCreateProgramWithBinary の違い
  4. オフラインコンパイル版は(本来)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 以外は未確認のため不明)