C
OpenCL
C言語

OpenCLを嗜む ~ 入門編 ~

More than 1 year has passed since last update.

諸用でOpenCLを始めることとなったのですが、OpenCL系の(日本語)記事が少ないようなと思ったのでメモ書き代りに書いてみます。

環境はmacを想定しています。

構造

OpenCLは制御用のホストコードと並列処理用のカーネルコードの2つに分かれています。役割を見てみましょう。

用語

  • デバイス: GPUやマルチコアCPUといった並列演算用のLSI

ホストコード

  • CPU側で動くプログラムです。
  • 並列処理を行わない部分の処理を書きます。
  • あるいは、デバイス(GPUなど)のセットアップ(メモリにデータ置いたり、結果読み出したりとか)やデバイスコードのコンパイル(場合による)、デバイスの制御(カーネルコードをデバイスに送る)を担当します。
  • CにOpenCLの関数を追加しただけです。

カーネルコード

  • デバイスで動く処理を書きます。
  • OpenCL Cと呼ばれる独自に拡張したCで書きます。

書いてみる

とりあえず書いてみようということで、デバイスに整数を渡して10倍するコードを書きます。突っ込まれる前に言っておくと実用性は皆無です。

環境構築

macの人はxcodeだけ入れれば大丈夫です。他の人は各種OpenCLのSDKを入れましょう。

ホストコード

main.c
#include <stdio.h>
#include <stdlib.h>
#include <OpenCL/opencl.h>

#define MEM_SIZE 128
#define MAX_SOURCE_SIZE 0x100000

int main(void) {
  cl_device_id device_id = NULL;
  cl_context context = NULL;
  cl_command_queue command_queue = NULL;
  cl_mem memobj = NULL;
  cl_program program = NULL;
  cl_kernel kernel = NULL;
  cl_platform_id platform_id = NULL;
  cl_uint ret_num_devices, ret_num_platforms;
  cl_int ret;

  float p;
  float f[] = {10.0};

  FILE *fp;
  char fileName[] = "./test.cl";
  char *source_str;
  size_t source_size;

  fp = fopen(fileName, "r");
  if(!fp) {
    exit(1);
  }
  source_str = (char*)malloc(MAX_SOURCE_SIZE);
  source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
  fclose(fp);

  ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
  ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

  context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

  memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float), NULL, &ret);
  ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, sizeof(float), f, 0, NULL, NULL);

  program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

  ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

  kernel = clCreateKernel(program, "test", &ret);

  ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

  ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

  ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, sizeof(float), &p, 0, NULL, NULL);

  printf("%f\n", p);

  ret = clFlush(command_queue);
  ret = clFinish(command_queue);
  ret = clReleaseKernel(kernel);
  ret = clReleaseProgram(program);
  ret = clReleaseMemObject(memobj);
  ret = clReleaseCommandQueue(command_queue);
  ret = clReleaseContext(context);

  free(source_str);

  return 0;
}

mac以外の人は#include <OpenCL/opencl.h>#include <CL/cl.h>にすると動くようです。

カーネルコード

test.cl
__kernel void test(__global float *a) {
  *a *= 10;
}

コンパイル

$ gcc main.c -framework OpenCL

mac以外の人はまちまちなので省略します。

解説

ホストコード

カーネルコードを読み込んでいるコードです。

  fp = fopen(fileName, "r");
  if(!fp) {
    exit(1);
  }
  source_str = (char*)malloc(MAX_SOURCE_SIZE);
  source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
  fclose(fp);

デバイスの情報を取得しています。

  ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
  ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

コンテキストというOpenCLを動作させる環境を作ります。

  context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

コマンドキューと呼ばれるデバイスに対して指示を送るためのキューを作成します。

  command_queue = clCreateCommandQueue(context, device_id, 0, &ret);

読み書き可能なバッファー(メモリ領域)を作成して初期化しています。メモリに書き込むやりとりもキューを通じて行います。

  memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float), NULL, &ret);
  ret = clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, sizeof(float), f, 0, NULL, NULL);

カーネルコードをオブジェクト化してコンパイルします。

  program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

  ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

コンパイルしたカーネルコードのうち呼び出すカーネル関数を選択します。この場合は「test」という関数が選ばれています。

  kernel = clCreateKernel(program, "test", &ret);

選択したカーネルの引数の代入をします。この場合test関数の0番目の引数にmemobjのアドレスを代入しています。

ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);

カーネルをデバイスで実行します。

ret = clEnqueueTask(command_queue, kernel, 0, NULL, NULL);

バッファから値を取得します。

ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, sizeof(float), &p, 0, NULL, NULL);

後始末を行うコードです。一応書いておきましょう。

  ret = clFlush(command_queue);
  ret = clFinish(command_queue);
  ret = clReleaseKernel(kernel);
  ret = clReleaseProgram(program);
  ret = clReleaseMemObject(memobj);
  ret = clReleaseCommandQueue(command_queue);
  ret = clReleaseContext(context);

カーネルコード

それと言って解説することがないのですが、カーネル関数の先頭には修飾子__kernelをつけましょう。あと変数の宣言の前についている__globalはCPUからもデバイスからも読み書き可能なバッファという意味があるのですが、このあたりの話は別のところでまとめます。

今後も別記事での解説や加筆を進めていくつもりです。