OpenCLのImageObjectsでの画像の扱いについて

細かい事を気にしてはいけない。

ホスト側の関数。
clCreateImage: イメージオブジェクトの作成。
clReleaseMemObject: バッファ或いはイメージオブジェクトの解放。
clEnqueueWriteImage: イメージオブジェクトへの書込み。
clEnqueueReadImage: イメージオブジェクトからの読出し。

cl_channel_order について。
メモリ上でのバイト順を指定する。
CL_RGBA: 下位から順にR, G, B, A。
CL_BGRA: 下位から順にB, G, R, A、これはWindowsのDIBSectionと同じである。

cl_channel_type について。
CL_UNORM_INT8: 符号無し, カーネルでの読書きは0.0~1.0, メモリ上のデータは符号なし8ビット整数、一般的にはこれを扱う。


カーネル側の関数など。
read_imagef: イメージオブジェクトから読出す。座標を整数で指定する場合はサンプラに制限がある。AMD製ソフトウェアでCPUデバイスを使用する場合、実数座標だと不具合が出るようだ。
write_imagef: イメージオブジェクトへ書込む。

仮引数image2d_t は、__write_only か __read_only の一方のみで修飾する必要がある。(デフォで __read_only 修飾だが。)


例として、画像をグレイスケールに変換するプログラムを示す。
画像の読書きについては、前回の投稿を参照すること。
Windowscodecs.lib, Shlwapi.lib, OpenCL.lib をリンクすること。

//////// begin: grayscale.cl
__kernel void grayscale(__write_only image2d_t imgDst, __read_only image2d_t imgSrc)
{
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_NONE| CLK_FILTER_NEAREST;
    int2 coord;
    float4 color;

    coord.x = get_global_id(0);
    coord.y = get_global_id(1);
    if (get_image_width(imgDst)<=coord.x || get_image_height(imgDst)<=coord.y)
        return;
    color = read_imagef(imgSrc, sampler, coord);
    color.xyz = dot(color.xyz, (float3)(0.299f, 0.587f, 0.114f)); //r=g=b = r*0.299 + g*0.587 + b*0.114
    color.w = 0.0f; //a = 0
    write_imagef(imgDst, coord, color);
}
//////// end: grayscale.cl

//////// begin: C++
#include <Windows.h>
#include <comdef.h>
#include <wincodec.h>
#include <Shlwapi.h>

#include <iostream>
#include <fstream>
#include <string>
#include <vector>
#include <CL/cl.hpp>

using namespace std;
using namespace cl;

HRESULT read(void **ppData, UINT *pWidth, UINT *pHeight, LPCWSTR pName);
HRESULT writeBmp(LPCWSTR pName, void *pData, UINT width, UINT height);

int main()
{
    CoInitializeEx(NULL, COINIT_MULTITHREADED| COINIT_DISABLE_OLE1DDE);

    vector<Platform> platforms;
    Platform::get(&platforms);
    cl_context_properties cps[3] = {CL_CONTEXT_PLATFORM, (cl_context_properties)(platforms[0])(), 0};
    Context context(CL_DEVICE_TYPE_ALL, cps);
    vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
    CommandQueue queue = CommandQueue(context, devices[0], 0);
   
    cout << platforms.size() << " platforms" << endl;
    cout << devices.size() << " devices" << endl;

    ifstream sourceFileName("grayscale.cl", ios::in);
    string sourceFile(istreambuf_iterator<char>(sourceFileName), (istreambuf_iterator<char>()));
    Program::Sources source(1, make_pair(sourceFile.c_str(), sourceFile.length()+1));
    Program program(context, source);
    program.build(devices);
    cl::Kernel kernel_grayscale(program, "grayscale");

    void *pData;
    UINT width, height;
    read(&pData, &width, &height, L"src.jpg");

    ImageFormat format(CL_BGRA, CL_UNORM_INT8);
    Image2D *pImgDst = new Image2D(context, CL_MEM_WRITE_ONLY, format, width, height);
    Image2D *pImgSrc = new Image2D(context, CL_MEM_READ_ONLY, format, width, height);
   
    cl::size_t<3> origin, region;
    region[0] = width;
    region[1] = height;
    region[2] = 1;
    queue.enqueueWriteImage(*pImgSrc, CL_TRUE, origin, region, 0, 0, pData);
   
    kernel_grayscale.setArg(0, *pImgDst);
    kernel_grayscale.setArg(1, *pImgSrc);

    NDRange wsGlobal(width, height);
    queue.enqueueNDRangeKernel(kernel_grayscale, NullRange, wsGlobal);
    queue.enqueueReadImage(*pImgDst, CL_TRUE, origin, region, 0, 0, pData);

    writeBmp(L"dst.bmp", pData, width, height);

    CoUninitialize();

    return 0;
}
//////// end: C++