この記事について

OpenCLを使ったプログラムの基本を確認するために、OpenCLのHello Worldプログラムを作成し、その内容を説明していきます。

OpenCLについて

OpenCLは並列コンピューティングを行うために標準化されているAPIの一つです。Appleにより提案され、現在は標準化団体であるKhronos Group(公式サイト)によりメンテナンスされています。OpenCLを使うことで、マルチコアCPU、GPU、DSPなど異種混在(ヘテロジニアス)な計算資源を使ったアプリケーションをクロスプラットフォームで開発できます。類似する仕様の代表的なもののひとつにnVidiaのCUDAがあります。

OpenCLプログラムの構成

OpenCLプログラムはホストコードとカーネルコードで構成されます。

ホストコード

CPUで動作するソフトウェアです
デバイスの初期化や並列処理プログラムの呼び出しタイミングを制御します

カーネルコード

並列処理を行うデバイス(GPU、DSPなど)上で動作するソフトウェアです
並列処理を実行し、結果をホスト側に返します
C言語に似たOpenCL Cという形式で記述します

Hello World

OpenCLを使ってHello Worldを表示するプログラムを作成します。Hello Worldの文字列をデバイス側で作成し、その情報をホスト側に戻して画面に表示させます。

ホストコードの作成

ホストコードの大まかな流れは以下のようになります。

デバイスの情報を取得します
OpenCLコンテキストと呼ばれるオブジェクトを作成します
ホスト↔デバイス間でメッセージをやり取りするためのOpenCLメッセージキューを作成します
カーネルコードをソースからビルドします
カーネルコードが扱うことができるデバイス側のメモリを確保します
カーネルコードを実行します
カーネルコードの実行結果をメッセージキューを使って取得します
確保した各種リソースを解放します

OpenCL API

まずは、OpenCLのAPIを読み込むための準備です。以下のヘッダを読み込みます。

#include <CL/cl.h>

OpenCL APIを利用するために、さまざまな管理オブジェクトを作成していくことになります。構造体（クラス）にまとめていきながら作成していきます。

デバイス情報の取得

// OpenCLのデバイス情報を管理するクラスを作成します
typedef struct ___opencl_device
{
    cl_platform_id    platform_id;       // PlatformのID
    cl_device_id      device_id;         // DeviceのID
    cl_int            num_of_platforms;  // システムが保持しているPlatform数
    cl_int            num_of_devices;    // システムが保持しているDevice数
}OpenCLDevice;

void OpenCLDevice_Init(OpenCLDevice* p_this){
    p_this->platform_id      = NULL;
    p_this->device_id        = NULL;
    p_this->num_of_platforms = 0;
    p_this->num_of_devices   = 0;
}

Error OpenCLDevice_GetDevices(OpenCLDevice* p_this){
    cl_int ret;
    // プラットフォームIDを取得します
    ret = clGetPlatformIDs(1, &(p_this->platform_id), &(p_this->num_of_platforms) );

    // プラットフォームが持っているデバイスIDとデバイス数を取得します
    ret = clGetDeviceIDs(p_this->platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &(p_this->device_id), &(p_this->num_of_devices) );

    // retにはOpenCLが提供するエラー番号が格納されます
    // 成功した場合は、CL_SUCCESSが格納されます

    // OpenCLのエラー番号を返却してもよいですが、
    // ここではシステムやプロジェクトで定義されているエラー定義を返しています
    return Success;
}

OpenCLコンテキストの取得

// OpenCLコンテキストを管理するクラスです
typedef struct ___opencl_context
{
    cl_context        context;  // OpenCLコンテキスト
    OpenCLDevice*     p_device; // 関連するデバイス
}OpenCLContext;

void OpenCLContext_Init(OpenCLContext* p_this, OpenCLDevice* p_dev){
    p_this->context       = NULL;
    p_this->p_device      = p_dev; // 関連するデバイスへの参照を保存
}

Error OpenCLContext_Create(OpenCLContext* p_this){
    cl_int ret;
    // コンテキストを取得します
    p_this->context = clCreateContext(NULL, 1, &(p_this->p_device->device_id), NULL, NULL, &ret );

    return Success;
}

Error OpenCLContext_Finalize(OpenCLContext* p_this){
    cl_int ret;
    // コンテキストをリリースします
    ret = clReleaseContext(p_this->context);

    return Success;
}

OpenCLのメッセージキューを作成

// OpenCLコマンドキュー情報を管理するクラス
typedef struct ___opencl_queue
{
    cl_command_queue  command_queue;  // OpenCLコマンドキュー
    OpenCLDevice*     p_device;       // 関連するデバイス
    OpenCLContext*    p_context;      // 関連するコンテキスト
}OpenCLQueue;

void  OpenCLQueue_Init( OpenCLQueue* p_this, OpenCLDevice* p_dev, OpenCLContext* p_ctx ){
    p_this->command_queue = NULL;
    p_this->p_device      = p_dev;
    p_this->p_context     = p_ctx;
}

Error OpenCLQueue_Create(OpenCLQueue* p_this){
    cl_int ret;

    // コマンドキューを作成します
    p_this->command_queue = clCreateCommandQueue(p_this->p_context->context, p_this->p_device->device_id, 0, &ret);

    return Success;
}

Error OpenCLQueue_Read(OpenCLQueue* p_this, const OpenCLBuffer* p_rbuff){
    cl_int ret;

    // コマンドキューからデータを読み込みます
    // 読み込むために、OpenCLのバッファオブジェクトを利用します
    clEnqueueReadBuffer(p_this->command_queue, p_rbuff->memobj, CL_TRUE, 0, p_rbuff->size, p_rbuff->p_buffer, 0, NULL, NULL);
    
    return Success;
}

Error OpenCLQueue_Finalize(OpenCLQueue* p_this){
    cl_int ret;

    // コマンドキューをリリースします
    ret = clFlush(p_this->command_queue);
    ret = clFinish(p_this->command_queue);
    ret = clReleaseCommandQueue(p_this->command_queue);
    
    return Success;
}

OpenCLバッファの作成

typedef struct ___opencl_buffer
{
    cl_mem            memobj;     // OpenCLメモリオブジェクト
    OpenCLContext*    p_context;  // 関連するコンテキスト

    char*             p_buffer;   // メモリオブジェクトを経由して情報をやりとりするメモリ
    size_t            size;       // p_bufferのサイズ
}OpenCLBuffer;


void OpenCLBuffer_Init(OpenCLBuffer* p_this, OpenCLContext* p_ctx){
    p_this->memobj    = NULL;
    p_this->p_context = p_ctx;

    p_this->p_buffer  = NULL;
    p_this->size      = 0;
}

Error OpenCLBuffer_Create(OpenCLBuffer* p_this, size_t size){
    cl_int ret;

    // メモリオブジェクトを作成します
    p_this->memobj   = clCreateBuffer(p_this->p_context->context, CL_MEM_READ_WRITE, size, NULL, &ret);
    
    // 後にメモリオブジェクトに渡すホスト側のメモリ領域を確保します
    p_this->p_buffer = malloc(size);
    p_this->size     = size;

    return Success;
}

Error OpenCLBuffer_Finalize(OpenCLBuffer* p_this){
    cl_int ret;

    // メモリオブジェクトをリリースします
    ret = clReleaseMemObject(p_this->memobj);

    // mallocで格納したホスト側のメモリも解放します
    free(p_this->p_buffer);

    return Success;
}

カーネル

typedef struct ___opencl_kernel
{
    cl_program          program;       // OpenCLプログラムオブジェクト
    cl_kernel           kernel;        // OpenCLカーネルオブジェクト

    OpenCLDevice*       p_device;      // 関連するデバイス

    OpenCLContext       context;       // カーネル用のコンテキスト
    OpenCLQueue         queue;         // カーネル用のコマンドキュー

    char*               p_kernel_name; // カーネル名
}OpenCLKernel;


void OpenCLKernel_Init(OpenCLKernel* p_this, OpenCLDevice* p_dev){
    p_this->program   = NULL;
    p_this->kernel    = NULL;
    p_this->p_device  = p_dev;

    p_this->p_kernel_name = NULL;
}


Error OpenCLKernel_Setup(OpenCLKernel* p_this){
    cl_int ret;

    // 変数を初期化します
    OpenCLContext_Init( &(p_this->context), p_this->p_device );
    OpenCLQueue_Init( &(p_this->queue), p_this->p_device, &(p_this->context) );

    // コンテキストとコマンドキューを作成
    OpenCLContext_Create( &(p_this->context) );
    OpenCLQueue_Create( &(p_this->queue) );

    return Success;
}



Error OpenCLKernel_CreateWithSource(OpenCLKernel* p_this, const char* p_src, size_t src_size ){
    cl_int ret;

    // カーネルソースコードからプログラムオブジェクトを作成します
    p_this->program = clCreateProgramWithSource(p_this->context.context, 1, (const char **)&p_src, (const size_t *)&src_size, &ret);

    // プログラムをビルドします
    ret = clBuildProgram(p_this->program, 1, &(p_this->p_device->device_id), NULL, NULL, NULL);

    // ビルド結果を表示します
    if (ret != CL_SUCCESS)
    {
        printf("Build Failure:%s\n", p_this->p_kernel_name);
    }
    else
    {
        printf("Build Success:%s\n", p_this->p_kernel_name);
    }

    // カーネルオブジェクトを作成します
    p_this->kernel = clCreateKernel(p_this->program, p_this->p_kernel_name, &ret);

    return Success;
}

Error OpenCLKernel_Execute(OpenCLKernel* p_this, OpenCLBuffer* p_args, int len){
    cl_int ret;

    // 並列数を指定するための変数です
    // この指定は最もシンプルな単一ワークアイテムのみのカーネルで実行する設定です
    cl_uint work_dim = 1;
    size_t  global_work_size[] = { 1 };
    size_t  local_work_size[]  = { 1 };

    // カーネル関数の左側から順番に引数を設定していきます
    int i;
    for(i=0;i<len;i++)
    {
       ret = clSetKernelArg( p_this->kernel, i, sizeof(cl_mem), &(p_args[i].memobj) );
    }

    // カーネルを実行します
    ret = clEnqueueNDRangeKernel(p_this->queue.command_queue, p_this->kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, NULL);

    // カーネルの実行結果を取得します
    // カーネル関数の第一引数が戻り値であることを前提にしています
    OpenCLQueue_Read( &(p_this->queue), &(p_args[0]) );

    return Success;
}


Error OpenCLKernel_Finalize(OpenCLKernel* p_this){
    cl_int ret;
    
    // コンテキストとコマンドキューオブジェクトをリリース
    OpenCLContext_Finalize( &(p_this->context) );
    OpenCLQueue_Finalize( &(p_this->queue) );

    // カーネルオブジェクトをリリース
    ret = clReleaseKernel(p_this->kernel);
    
    // プログラムオブジェクトをリリース
    ret = clReleaseProgram(p_this->program);
    
    return Success;
}

カーネルコードの作成

OpenCL Cの形式で記述します。外部ファイルに記述して、本体プログラムから読み込むこともできますが、今回は、カーネルソースをC言語の配列内に直接書き込みます。

static char g_src[] =
    "__kernel void hello(__global char* string){\n"
    "string[ 0] = 'H';\n"
    "string[ 1] = 'e';\n"
    "string[ 2] = 'l';\n"
    "string[ 3] = 'l';\n"
    "string[ 4] = 'o';\n"
    "string[ 5] = ',';\n"
    "string[ 6] = 'w';\n"
    "string[ 7] = 'o';\n"
    "string[ 8] = 'r';\n"
    "string[ 9] = 'l';\n"
    "string[10] = 'd';\n"
    "string[11] = '.';\n"
    "string[12] = '\\n';\n"
    "string[13] = '\\0';\n"
    "}\n";

サンプルプログラム

int main(){
    OpenCLDevice  device;  // デバイス情報を管理します
    OpenCLKernel  kernel;  // カーネル情報を管理します
    OpenCLBuffer  arg;     // カーネルに渡す引数の情報を格納するメモリオブジェクトです

    // デバイス情報の取得
    OpenCLDevice_Init( &device );
    OpenCLDevice_GetDevices( &device );

    // カーネルコードのビルド
    OpenCLKernel_Init( &kernel, &device );
    OpenCLKernel_Setup( &kernel );
    kernel.p_kernel_name = "hello";
    OpenCLKernel_CreateWithSource( &kernel, g_src, sizeof(g_src) );

    // 引数に渡すバッファを作成
    OpenCLBuffer_Create( &arg, 32 );

    // カーネルを実行
    OpenCLKernel_Execute( &kernel, &arg, 1 );

    // カーネルからの戻り値を表示
    printf("from device=%s",arg.p_buffer);
    
    // 作成したオブジェクトをリリース
    OpenCLKernel_Finalize(&kernel);
    OpenCLBuffer_Finalize(&arg);

    return 0;
}

成功すると、以下のようにデバイスで作成したメッセージがホスト側に渡されて表示されます。

Build Success:hello
from device=Hello,world.