以前書いたCUDAの記事のOpenCL版。
こちらも、いちばん簡単なんてタイトル付けてるけど実は他のやり方をあまり調べてないというか、OpenCLはじめて数時間の初心者なので、もっといい方法があれば教えてください。

0. 準備

今回使ったのはXcode 4.6.2。
CUDAのときはCUDAドライバやtoolkitのインストールが必要だったけど、最近のMacはOpenCLに標準対応してるのでXcodeさえあればすぐに開発をはじめられるはず。

OpenCLは特にGPU限定のAPIってワケじゃないけど、今回はGPUの利用を前提とした説明となる。

1. Xcodeで新規プロジェクトを作成

OS X → Application → Command Line Tool を選択。
今回は言語をC++とした。

ちなみにCUDAのときはホストコードの言語をCにしたけど、CUDAカーネルはC++で書ける。
OpenCLを使う今回はホストコードの言語をC++としたけど、OpenCLカーネルはCで書く必要がある。
つまり言語の選択があべこべなんだけど、ホストコードの言語を自由に選べるという事を示すために敢えてそうしている。

プロジェクト名はOpenCLTest。
TERGETS → Build Phases → Link Binary With Libraries にて OpenCL.framework を追加。

2. カーネルを書く

kernel.cl というファイルを新規追加しコードを書く。
今回はテストということで以下のようにデータをインクリメントするだけの簡単なカーネルを書いた。

__kernel
void addone(__global float* data, const int n)
{
    int index = get_global_id(0);
    
    if (index < n) {
        data[index] += 1.0f;
    }
}

CUDAのときはインデクス計算部分をあえて汎用的に書いたというのもあるけど、OpenCLの場合はget_global_idだけで済むのでその点楽。
get_global_id(0)だけで済むかどうかはホストコードからのカーネルの呼び出し方に依存する。

3. kernel.clをコンパイルする設定

XcodeはOpenCLに対応しているので、CUDAの時のように独自のビルドルールを追加したりする必要はない。
普通のCのソースのように、単に TERGETS → Build Phases → Compile Souces に kernel.cl を追加するのみ。

4. ホストコードを書く。

ここまで来たらあとは main.cpp で #include <OpenCL/opencl.h> するだけでOpenCLを使ったコードが書ける。
とりあえず今回は以下のようなコードを書いた。
処理内容としては、1000要素のfloatの配列に+1.0するのみ。
初期値が0.0〜999.0なのでGPUを通した後は1.0〜1000.0になるはず。

#include <iostream>
#include <vector>
#include <OpenCL/opencl.h>

#define PLATFORM_MAX 4
#define DEVICE_MAX 4


void EC(cl_int result, const char *title)
{
    if (result != CL_SUCCESS) {
        std::cout << "Error: " << title << "(" << result << ")\n";
    }
}


cl_int err = CL_SUCCESS;
void EC2(const char *title)
{
    if (err != CL_SUCCESS) {
        std::cout << "Error: " << title << "(" << err << ")\n";
    }
    err = CL_SUCCESS;
}


int main(int argc, const char * argv[])
{
    // プラットフォーム一覧を取得
    cl_platform_id platforms[PLATFORM_MAX];
    cl_uint platformCount;
    EC(clGetPlatformIDs(PLATFORM_MAX, platforms, &platformCount), "clGetPlatformIDs");
    if (platformCount == 0) {
        std::cerr << "No platform.\n";
        return EXIT_FAILURE;
    }
    
    // 見つかったプラットフォームの情報を印字
    for (int i = 0; i < platformCount; i++) {
        char vendor[100] = {0};
        char version[100] = {0};
        EC(clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, nullptr), "clGetPlatformInfo");
        EC(clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(version), version, nullptr), "clGetPlatformInfo");
        std::cout << "Platform id: " << platforms[i] << ", Vendor: " << vendor << ", Version: " << version << "\n";
    }
    
    // デバイス一覧を取得
    cl_device_id devices[DEVICE_MAX];
    cl_uint deviceCount;
    EC(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, DEVICE_MAX, devices, &deviceCount), "clGetDeviceIDs");
    if (deviceCount == 0) {
        std::cerr << "No device.\n";
        return EXIT_FAILURE;
    }
    
    // 見つかったデバイスの情報を印字
    std::cout << deviceCount << " device(s) found.\n";
    for (int i = 0; i < deviceCount; i++) {
        char name[100] = {0};
        size_t len;
        EC(clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(name), name, &len), "clGetDeviceInfo");
        std::cout << "Device id: " << i << ", Name: " << name << "\n";
    }
    
    // コンテキストの作成
    cl_context ctx = clCreateContext(nullptr, 1, devices, nullptr, nullptr, &err);
    EC2("clCreateContext");
    
    // コンパイル済みclプログラムの読み込み
    const char* bitcode_path = "OpenCL/kernel.cl.gpu_32.bc";
    size_t len = strlen(bitcode_path);
    cl_program program = clCreateProgramWithBinary(ctx, 1, devices, &len, (const unsigned char**)&bitcode_path, nullptr, &err);
    EC2("clCreateProgramWithBinary");
    
    // プログラムのビルド
    EC(clBuildProgram(program, 1, devices, nullptr, nullptr, nullptr), "clBuildProgram");
    
    // カーネルの作成
    cl_kernel kernel = clCreateKernel(program, "addone", &err);
    EC2("clCreateKernel");
    
    // データを用意
    int n = 1000;
    std::vector<float> data(n);
    for (int i = 0; i < n; i++) {
        data[i] = float(i);
    }
    
    // デバイスメモリを確保しつつデータをコピー
    cl_mem device_mem = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float) * n, data.data(), &err);
    EC2("clCreateBuffer");
    
    // カーネルの引数をセット
    EC(clSetKernelArg(kernel, 0, sizeof(cl_mem), &device_mem), "clSetKernelArg");
    EC(clSetKernelArg(kernel, 1, sizeof(int), &n), "clSetKernelArg");
    
    // コマンドキューの作成
    cl_command_queue q = clCreateCommandQueue(ctx, devices[0], 0, &err);
    EC2("clCreateCommandQueue");
    
    // カーネルの実行
    size_t global = n;
    EC(clEnqueueNDRangeKernel(q, kernel, 1, nullptr, &global, nullptr, 0, nullptr, nullptr), "clEnqueueNDRangeKernel");
    
    // 結果を読み込み
    EC(clEnqueueReadBuffer(q, device_mem, CL_TRUE, 0, sizeof(float) * n, data.data(), 0, nullptr, nullptr), "clEnqueueReadBuffer");
    
    // 結果の印字
    for (int i = 0; i < n; i++) {
        std::cout << data[i] << ", ";
    }
    std::cout << "\n";
    
    // コマンドキューの解放
    EC(clReleaseCommandQueue(q), "clReleaseCommandQueue");
    
    // デバイスメモリを解放
    EC(clReleaseMemObject(device_mem), "clReleaseMemObject");
    
    // カーネルの解放
    EC(clReleaseKernel(kernel), "clReleaseKernel");
    
    // プログラムの解放
    EC(clReleaseProgram(program), "clReleaseProgram");
    
    // コンテキストの解放
    EC(clReleaseContext(ctx), "clReleaseContext");
    
    std::cout << "Done.\n";
    return EXIT_SUCCESS;
}

処理の流れはCUDA Driver APIの場合とほとんど同じという印象。

5. 実行結果

Platform id: 0x7fff0000, Vendor: Apple, Version: OpenCL 1.2 (Dec  4 2012 18:26:30)
1 device(s) found.
Device id: 0, Name: GeForce 9400M
1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 
ーーー　大胆に省略　ーーー
971, 972, 973, 974, 975, 976, 977, 978, 979, 980, 981, 982, 983, 984, 985, 986, 987, 988, 989, 990, 991, 992, 993, 994, 995, 996, 997, 998, 999, 1000, 
Done.

6. 感想

CUDA Driver APIを使った経験があればOpenCLはそう難しくないのではないかという印象を受けた。
OpenCLはCUDAより面倒という話はよく聞くが、Driver APIばっかり使ってる身としては、Xcodeのサポートのおかげもあって、OpenCLの方が簡単に感じる。
Macは少し古くなると新しいGPUドライバがアップデートで配信されなくなるみたいなので、CUDAを使うよりは標準対応のOpenCLを使ったほうがいいかもしれない。

ただ、OpenGLのシェーダもそうだけど、GPUメーカーのドライバの出来不出来で、コンパイルしたシェーダやカーネルの動きに違いが出るらしいという話もあるので注意が必要。
CUDAのほうはNVIDIA限定なのでそういう問題はあまりないと思われる。
（まぁこれは当たり前で、メリットというよりデメリットなんだけど。）

個人的には、カーネルをCで書かなきゃいけないのがつらい。
やっぱクラス使いたいし。
AMDのツールを使ったり、LLVM Clangで工夫すればC++でカーネルを書くのも不可能ではないらしい。
nvccと比べてどのくらいのC++対応具合か後で調べてみたい。
ただ、環境が限られるなら普通にCUDA使ったほうがマシかも。

あと、OpenCLの関数はexternで引数名無しで宣言されてるので、Xcodeのコード補完時に引数の型しか出てこなくて、具体的に何を指定すればいいのかが推測しづらい。
細かい点だけどこれは記憶力が弱い私には地味につらい。
後から知った事だが、OpenCLにはopencl.hのかわりにcl.hppをincludeする事で使える、C++ Bindingsというものがあって、C++でホストコードを書く場合はそれを使うといいらしい。

もうひとつ、OpenCLの関数はエラーコードを返り値で返すものと、ポインタ引数で返すものが混在してるので、対応するのが少し面倒かもしれない。

個人でNVIDIAのGPUを使ってGPGPUする場合はCUDAを、配布を想定する場合や、Macで本格的にGPGPUする場合はOpenCLを選択するのが良いんじゃないかと思った。