OS XのOpenCL - その11

さて、久しぶりのOpenCLの抄訳の続き。もうただの惰性。あれですな、「無精の代参」

「背中をぽーんと突いてくれ。突かれた弾みで歩き出す」
「けったいな人やな。ほな、ぽーんと」
「うわ、そんな強う突いたらあかんがな。丼池通り越してしまうがな」

い、いかん、これを始めるとキリがない．．．

7.3  カーネルからイメージオブジェクトをアクセスする

gcl_copy云々関数を使ってイメージとホストメモリとの間でコピーできる。デバイス上で実際にこのイメージを処理するには、デバイス上で実行するワークアイテムに対して利用可能にしなければならない。つぎのセクションはカーネルのさらなる処理のためにどうやってデータを渡すかを示す。

7.4  イメージオブジェクトのマップ

イメージオブジェクトのある領域をホストのアドレス空間にマップするには

void *gcl_map_image(cl_image image,
                    cl_map_flags map_flags,
                    const size_t origin[3],
                    const size_t region[3]);

マップされた領域へのポインタが返る。

image

マップされるイメージ

map_flags

戻り値のポインタをどう使うかによってCL_MAP_READとCL_MAP_WRITEのビットフィールドを指定する

origin

マップする先頭のピクセル座標(x,y,z)

region

ピクセル単位のマップ領域

注意：バイトではなくピクセルで指定するので、実際のバイト数はイメージフォーマットに依存する。この関数はOpenCL規格のclEnqueueMapImage関数と同じような動作をする。

7.5  マップの解除

マップの解除は

void gcl_unmap(void *ptr);

を呼ぶ。ptrはデバイスメモリ、またはイメージへのポインタである。

注意：gcl_map_ptrやgcl_map_imageが返したポインタを渡すのではないことに注意。

7.6  イメージオブジェクトの参照カウントの操作

void gcl_retain_image(cl_image image)
void gcl_release_image(cl_image image)

7.7  例

つぎの例では、ホストは入力用と出力用にイメージオブジェクトを作って、カーネルで赤と緑のピクセルを入れ替え、そのあとで結果をチェックしている。

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

// いつもと同じようにXcodeが自動生成したファイルをインクルード
#include "kernels.cl.h"

#define COUNT 2048

static void display_device(cl_device_id device)
{
    char name_buf[128];
    char vendor_buf[128];
    
    clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(char)*128, name_buf, NULL);
    clGetDeviceInfo(device, CL_DEVICE_VENDOR, sizeof(char)*128, vendor_buf, NULL);
    
    fprintf(stdout, "Using OpenCL device: %s %s\n", vendor_buf, name_buf);
}

static void image_test(const dispatch_queue_t dq)
{
    // ホストアプリとOpenCLデバイスの作業終了の同期を取るために
    // ディスパッチセマフォを使っている
    dispatch_semaphore_t dsema = dispatch_semaphore_create(0);

    // この例では8ビットRGBAイメージのフェイクを使っている。赤に値を詰める
    // 実際のプログラムではモノホンのラスタデータを使え。
    // 多くのOpenCLデバイスは広範囲のイメージフォーマットをサポートしている。
    
    unsigned int i;
    size_t height = 2048, width = 2048;
    
    unsigned int *pixels =
    (unsigned int*)malloc( sizeof(unsigned int) * width * height );
    
    for (i = 0; i < width*height; i++)
        pixels[i] = 0xFF0000FF; // 0xAABBGGRR: 8bits per channel, all red.
    
    // このイメージデータはホスト側にある
    // 操作のためには入力と出力のふたつのOpenCLイメージを必要がある
    
    // これはイメージデータのフォーマットを記述している
    cl_image_format format;
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_UNSIGNED_INT8;
    
    cl_mem input_image = gcl_create_image(&format, width, height, 1, NULL);
    cl_mem output_image = gcl_create_image(&format, width, height, 1, NULL);
    
    dispatch_async(dq, ^{

        // このカーネルはそれぞれのアークアイテムがひとつのピクセルを操作するように書いてある
        // 従って、2次元の幅と高さ全部にわたって実行することになる
        
        cl_ndrange range = {
            2,                  // 2次元である
            {0},                // 最初からである
            {width, height},    // 幅×高さ
            {0}                 // ワークグループへの分解はOpenCLに任せる
        };
        
        // ホスト側の初期ピクセルデータをOpenCLデバイスのイメージオブジェクトにコピーする
        // ここでは画像全体をコピーしているが、画像の一部をoriginとregionとして
        // 領域とオフセットを指定してもいい
        const size_t origin[3] = { 0, 0, 0 };
        const size_t region[3] = { width, height, 1 };
        gcl_copy_ptr_to_image(input_image, pixels, origin, region);
        
        // やらんかい！
        red_to_green_kernel(&range, input_image, output_image);
        
        // 結果を読み出す。ホスト側のバッファを再利用している
        gcl_copy_image_to_ptr(pixels, output_image, origin, region);

        // 終わったかどうかをホストに知らせる
        dispatch_semaphore_signal(dsema);
    });
    
    // 他にすることがあれば．．．．

    // ．．．しかし最後にはOpenCLが完了するのを待つ
    dispatch_semaphore_wait(dsema, DISPATCH_TIME_FOREVER);

    // それぞれのピクセルが0xFF00FF00になっているはず
    // ずーっと緑が埋まっている
    int results_ok = 1;
    for (i = 0; i < width*height; i++) {
        if (pixels[i] != 0xFF00FF00) {
            fprintf(stdout,
                    "Oh dear. Pixel %d was not correct.
                    Expected 0xFF00FF00, saw %x\n",
                    i, pixels[i]);
            results_ok = 0;
            break;
        }
    }
    
    if (results_ok)
        fprintf(stdout, "Image results OK!\n");
    
    // デバイス側のアロケーションを解放する
    // これはOpenCL規格のAPIであることに注意
    clReleaseMemObject(input_image);
    clReleaseMemObject(output_image);

    // ホスト側を解放する
    free(pixels);
}

int main (int argc, const char * argv[])
{
    // CPUでのディスパッチキューをとってくる
    dispatch_queue_t dq = gcl_create_dispatch_queue(CL_DEVICE_TYPE_CPU, NULL);
    if (!dq)
    {
        fprintf(stdout, "Unable to create a CPU-based dispatch queue.\n");
        exit(1);
    }
    
    // このディスパッチキューに関連付けられたOpenCLデバイスを表示する
    display_device(gcl_get_device_id_with_dispatch_queue(dq));
    
    image_test(dq);
    
    fprintf(stdout, "\nDone.\n\n");
    
    dispatch_release(dq);
}

つぎのリスト8-2は赤と緑のチャンネルを入れ替えるカーネルのコードである。

// 赤と緑のチャンネルを入れ替える簡単なカーネル

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST;

kernel void red_to_green(read_only image2d_t input, write_only image2d_t output)
{
    size_t x = get_global_id(0);
    size_t y = get_global_id(1);
    
    uint4 tap = read_imageui(input, sampler, (int2)(x,y));
    write_imageui(output, (int2)(x,y), tap.yxzw);
}