OpenCLでlanczos3を使った画像の拡大処理(4)

呼出側を書いておく。
実行する場合は、カレントディレクトリに前回の lanczos3.cl を置いておく。
カーネルの文字コードには注意した方が良いかもしれない。

//////// begin: main.cpp
#include <Windows.h>
#include <comdef.h>
#include <wincodec.h>
#include <Shlwapi.h>
#include <iostream>
#include <fstream>
#include <string>
#include <vector>
#include <CL/cl.hpp>

#pragma comment(lib, "Windowscodecs")
#pragma comment(lib, "Shlwapi")
#pragma comment(lib, "OpenCL")

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("lanczos3.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_setupTable(program, "lanczos3_setupTable");
    cl::Kernel kernel_passFirst(program, "lanczos3_passFirst");
    cl::Kernel kernel_passSecond(program, "lanczos3_passSecond");


    const cl_int widthDst = 1920;
    const cl_int heightDst = 1080;
    cl_int elementsHor;
    cl_int elementsVert;
    void *pData;
    UINT width, height;

   
    read(&pData, &width, &height, L"src.bmp");

    if (widthDst < width)
        elementsHor = 6 * width / widthDst;
    else
        elementsHor = 6;

    if (heightDst < height)
        elementsVert = 6 * height / heightDst;
    else
        elementsVert = 6;

    Buffer *pBufferTable = new Buffer(context, CL_MEM_READ_WRITE, (sizeof (float))*((elementsHor*widthDst)+(elementsVert*heightDst)));
    ImageFormat format(CL_BGRA, CL_UNORM_INT8);
    Image2D *pImgSrc = new Image2D(context, CL_MEM_READ_ONLY, format, width, height);
    Image2D *pImgDst = new Image2D(context, CL_MEM_WRITE_ONLY, format, widthDst, heightDst);
    Image2D *pImgTmp = new Image2D(context, CL_MEM_READ_WRITE, format, widthDst, height);

    //src
    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);
   
    //table
    kernel_setupTable.setArg(0, *pBufferTable);
    kernel_setupTable.setArg(1, cl_int(width));
    kernel_setupTable.setArg(2, cl_int(height));
    kernel_setupTable.setArg(3, cl_int(widthDst));
    kernel_setupTable.setArg(4, cl_int(heightDst));

    NDRange wsGlobalTable(widthDst+heightDst);
    queue.enqueueNDRangeKernel(kernel_setupTable, NullRange, wsGlobalTable);
    queue.finish();


    //pass 1
    kernel_passFirst.setArg(0, *pImgTmp);
    kernel_passFirst.setArg(1, *pImgSrc);
    kernel_passFirst.setArg(2, *pBufferTable);
    NDRange wsGlobal(widthDst, height);
    queue.enqueueNDRangeKernel(kernel_passFirst, NullRange, wsGlobal);
    queue.finish();

    //pass 2
    kernel_passSecond.setArg(0, *pImgDst);
    kernel_passSecond.setArg(1, *pImgTmp);
    kernel_passSecond.setArg(2, *pBufferTable);
    kernel_passSecond.setArg(3, width);
    NDRange wsGlobal2(widthDst, heightDst);
    queue.enqueueNDRangeKernel(kernel_passSecond, NullRange, wsGlobal2);
    queue.finish();
   
    BYTE *pData2 = new BYTE[4*widthDst*heightDst];
    region[0] = widthDst;
    region[1] = heightDst;
    queue.enqueueReadImage(*pImgDst, CL_TRUE, origin, region, 0, 0, pData2);
    writeBmp(L"dst.bmp", pData2, widthDst, heightDst);

    CoUninitialize();

    return 0;
}


_COM_SMARTPTR_TYPEDEF(IWICImagingFactory, __uuidof(IWICImagingFactory));
_COM_SMARTPTR_TYPEDEF(IWICBitmapDecoder, __uuidof(IWICBitmapDecoder));
_COM_SMARTPTR_TYPEDEF(IWICBitmapFrameDecode, __uuidof(IWICBitmapFrameDecode));
_COM_SMARTPTR_TYPEDEF(IWICBitmapSource, __uuidof(IWICBitmapSource));
_COM_SMARTPTR_TYPEDEF(IStream, __uuidof(IStream));
_COM_SMARTPTR_TYPEDEF(IWICBitmapEncoder, __uuidof(IWICBitmapEncoder));
_COM_SMARTPTR_TYPEDEF(IWICBitmapFrameEncode, __uuidof(IWICBitmapFrameEncode));

HRESULT read(void **ppData, UINT *pWidth, UINT *pHeight, LPCWSTR pName)
{
    IWICImagingFactoryPtr pFactory;
    IWICBitmapDecoderPtr pDecoder;
    IWICBitmapFrameDecodePtr pFrame;
    IWICBitmapSourcePtr pBitmap;
    UINT stride;
    UINT size;


    pFactory.CreateInstance(CLSID_WICImagingFactory, NULL, CLSCTX_INPROC_SERVER);
    pFactory->CreateDecoderFromFilename(pName, NULL, GENERIC_READ, WICDecodeMetadataCacheOnDemand, &pDecoder);
    pDecoder->GetFrame(0, &pFrame);
    WICConvertBitmapSource(GUID_WICPixelFormat32bppBGR, pFrame, &pBitmap); //DIB section compatible
    pBitmap->GetSize(pWidth, pHeight);
    stride = 4 * *pWidth;
    size = stride * *pHeight;
    *ppData = new BYTE[size];
    pBitmap->CopyPixels(NULL, stride, size, (BYTE *)*ppData);

    return S_OK;
}

HRESULT writeBmp(LPCWSTR pName, void *pData, UINT width, UINT height)
{
    IWICImagingFactoryPtr pFactory;
    IWICBitmapEncoderPtr pEncoder;
    IStreamPtr pStream;
    IWICBitmapFrameEncodePtr pFrame;
    WICPixelFormatGUID guidFormat;


    pFactory.CreateInstance(CLSID_WICImagingFactory, NULL, CLSCTX_INPROC_SERVER);
    pFactory->CreateEncoder(GUID_ContainerFormatBmp, NULL, &pEncoder);
    SHCreateStreamOnFileEx(pName, STGM_READWRITE| STGM_CREATE, FILE_ATTRIBUTE_NORMAL, TRUE, NULL, &pStream);
    pEncoder->Initialize(pStream, WICBitmapEncoderNoCache);
    pEncoder->CreateNewFrame(&pFrame, NULL);
    pFrame->Initialize(NULL);
    guidFormat = GUID_WICPixelFormat32bppBGR;
    pFrame->SetPixelFormat(&guidFormat);
    pFrame->SetSize(width, height);
    pFrame->WritePixels(height, 4*width, 4*width*height, (BYTE *)pData);
    pFrame->Commit();
    pEncoder->Commit();

    return S_OK;
}
//////// end: main.cpp

OpenCLでlanczos3を使った画像の拡大処理(3)

重みのテーブルを作成する。
構造は、
float [dstの幅][水平方向タップ数]
float [dstの高さ][垂直方向タップ数]
を隙間なく詰めた感じだ。

//////// begin: lanczos3.cl
float lanczos3(float distance)
{
    float weight;


    if (FLT_EPSILON > fabs(distance))
        weight = 1.0f;
    else if (3.0f > fabs(distance))
        weight = 3.0f * sin(M_PI_F*distance) * sin(M_PI_F*distance/3.0f) / (M_PI_F*M_PI_F*distance*distance);
    else
        weight = 0.0f;

    return weight;
}

void setupTablePartial(__global float *pWeight, int index, int lengthSrc, int lengthDst, int elements)
{
    float center = (index+0.5f) * lengthSrc / lengthDst;
    float beginSrc = trunc(center-elements/2);
    float distance;
    float sumWeight = 0.0f;


    //拡大もしくは等倍の場合
    if (lengthDst >= lengthSrc)
    {
        distance = (beginSrc+0.5f)-center;

        for (int i=0; i<elements; i++)
        {
            sumWeight += pWeight[i] = lanczos3(distance);
            distance += 1.0f;
        }
    }
    //縮小の場合
    else
    {
        float positionSrc = beginSrc + 0.5f;

        for (int i=0; i<elements; i++)
        {
            distance = (index+0.5f) - (positionSrc * lengthDst / lengthSrc);
            sumWeight += pWeight[i] = lanczos3(distance);
            positionSrc += 1.0f;
        }
    }

    //ここで割っておくことで、各ドットの計算時に割る手間をなくす
    for (int i=0; i<elements; i++)
        pWeight[i] /= sumWeight;
}


//テーブルの要素数: elementsHor*widthDst + elementsVert*heightDst;
//elementsHor: 拡大もしくは等倍の場合 6, 縮小の場合 6 * widthSrc / widthDst;
//elementsVert: 拡大もしくは等倍の場合 6, 縮小の場合 6 * heightSrc / heightDst;
__kernel void lanczos3_setupTable(__global float *pTable, int widthSrc, int heightSrc, int widthDst, int heightDst)
{
    int index = get_global_id(0);
    int elementsHor;
    int elementsVert;


    if (widthDst+heightDst <= index)
        return;

    if (widthDst < widthSrc)
        elementsHor = 6 * widthSrc / widthDst;
    else
        elementsHor = 6;

    if (heightDst < heightSrc)
        elementsVert = 6 * heightSrc / heightDst;
    else
        elementsVert = 6;

    //水平方向の場合
    if (widthDst > index)
        setupTablePartial(pTable+(elementsHor*index), index, widthSrc, widthDst, elementsHor);
    //垂直方向の場合
    else
        setupTablePartial(pTable+(elementsHor*widthDst)+(elementsVert*(index-widthDst)), index-widthDst, heightSrc, heightDst, elementsVert);
}
//////// end: lanczos3.cl
手元の環境では丸めモードがOpenCL(GPU)とC++(CPU)で異なり、OpenCL(GPU)では0方向への丸めで、C++(CPU)では直近値への丸め(同距離なら偶数寄り)だった。
他にもいくつかの相違があり同一の計算結果を得るのは難しいため、テーブルをCPUで作るのは好ましくない。

OpenCLでlanczos3を使った画像の拡大処理(2)

1: imgSrcの画像を水平方向に補完してimgTmpに書込む。
2: imgTmpの画像を垂直方向に補完してimgDstに書込む

imgTmp.width = imgDst.width
imgTmp.height = imgSrc.height
とする。

//////// begin:
#define GAMMA_CORRECT 0
__kernel void lanczos3_hmag_vmag_passFirst(__write_only image2d_t imgTmp, __read_only image2d_t imgSrc)
{
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_MIRRORED_REPEAT| CLK_FILTER_NEAREST;

    int2 coordDst = (int2)(get_global_id(0), get_global_id(1));
    float widthSrc = get_image_width(imgSrc);
    float widthDst = get_image_width(imgTmp);
    float heightDst = get_image_height(imgTmp);


    if (widthDst<=coordDst.x || heightDst<=coordDst.y)
        return;

    float center = (coordDst.x+0.5f) * (widthSrc/widthDst);
    float beginSrc = trunc(center-3.0f);
    float distance = (beginSrc+0.5f)-center;
    float2 coordSample = (float2)(beginSrc, coordDst.y);
    float sumWeight = 0.0f;
    float4 sumColor = 0.0f;

    for (int i=0; i<6; i++)
    {
        float weight = lanczos3(distance);
        float4 color = read_imagef(imgSrc, sampler, coordSample);
#if GAMMA_CORRECT
        color = pow(color, 2.2);
#endif
        sumColor += color * weight;
        sumWeight += weight;
        coordSample.x += 1.0f;
        distance += 1.0f;
    }

    write_imagef(imgTmp, coordDst, sumColor / sumWeight);
}

__kernel void lanczos3_hmag_vmag_passSecond(__write_only image2d_t imgDst, __read_only image2d_t imgTmp)
{
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_MIRRORED_REPEAT| CLK_FILTER_NEAREST;

    int2 coordDst = (int2)(get_global_id(0), get_global_id(1));
    float heightSrc = get_image_height(imgTmp);
    float heightDst = get_image_height(imgDst);
    float widthDst = get_image_width(imgTmp);


    if (widthDst<=coordDst.x || heightDst<=coordDst.y)
        return;

    float center = (coordDst.y+0.5f) * (heightSrc/heightDst);
    float beginSrc = trunc(center-3.0f);
    float distance = (beginSrc+0.5f)-center;
    float2 coordSample = (float2)(coordDst.x, beginSrc);
    float sumWeight = 0.0f;
    float4 sumColor = 0.0f;

    for (int i=0; i<6; i++)
    {
        float weight = lanczos3(distance);
        sumColor += weight * read_imagef(imgTmp, sampler, coordSample);
        sumWeight += weight;
        coordSample.y += 1.0f;
        distance += 1.0f;
    }

#if GAMMA_CORRECT
    write_imagef(imgDst, coordDst, pow(sumColor / sumWeight, 1.0f/2.2f));
#else
    write_imagef(imgDst, coordDst, sumColor / sumWeight);
#endif
}
//////// end:
遅いカーネルでは平均38.7[ms]、今回のカーネルでは平均8.69[ms]といった具合。
処理点数は36+6から6+6への減少なのにスループットは4倍以上になっているが、原因は分からない。

今回のカーネルでは重みを毎回計算しているが、事前にテーブルを作成しておき、そこから引けば速くなりそうだ。
次回はテーブルの作成のみ書く。

OpenCLでlanczos3を使った画像の拡大処理(1)

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

手順は次の通りである。
1: 水平方向に補完する。
2: 1で補完したデータを垂直方向に補完する。

水平方向のみの補間は、次のようになる。
1: dstの書込み位置がsrcのどの位置に対応するか求め、これをcとする
2: srcにあるcから距離3以内の6点の輝度に対して、lanczos窓を掛けて総和を求め、重みの総和で割り、これをdstの点の輝度とする

距離の考え方は少々ややこしい。
dst座標での書込み位置をsrc座標で表したもの = (dst座標での書込み位置 + 0.5) * (srcの幅 / dstの幅);
srcのサンプリング開始座標 = trunc(書込み位置をsrc座標で表したもの - 3);
最初のサンプリング点の距離 = (srcのサンプリング開始座標 + 0.5) - 書込み位置をsrc座標で表したもの;
．．．としておこう。

画像端を容易に処理するために、アドレッシングモードを CLK_ADDRESS_MIRRORED_REPEAT にする。
そのため read_imagef の座標指定をfloat2で行うが、環境によってはCPU処理の場合に期待通り読み取れない場合がある。

今回は遅く汚いカーネルを示す。
より正しく計算する場合はガンマ補正も考慮すべきだが、ここでは割愛する。
//////// begin:
float lanczos3(float d)
{
    float weight;


    if (0.0f == d)
        weight = 1.0f;
    else if (3.0f > fabs(d))
        weight = 3.0f * sin(M_PI_F*d) * sin(M_PI_F*d/3.0f) / (M_PI_F*M_PI_F*d*d);
    else
        weight = 0.0f;

    return weight;
}

__kernel void lanczos3_hmag_vmag_slow_fuck(__write_only image2d_t imgDst, __read_only image2d_t imgSrc)
{
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_MIRRORED_REPEAT| CLK_FILTER_NEAREST;

    float2 coordDst = (float2)(get_global_id(0), get_global_id(1));
    float2 sizeSrc = (float2)(get_image_width(imgSrc), get_image_height(imgSrc));
    float2 sizeDst = (float2)(get_image_width(imgDst), get_image_height(imgDst));
    float2 center;
    float2 beginSrc;
    float2 distance;
    float2 coordSrc;
    float4 colorVert[6];


    if (sizeDst.x<=coordDst.x || sizeDst.y<=coordDst.y)
        return;

    center = (coordDst+0.5f) * (sizeSrc / sizeDst);
    beginSrc = trunc(center - 3.0f);
    distance = (beginSrc+0.5f) - center;

    //水平方向に補完してcolorVert[]に入れる
    //垂直方向ループ
    coordSrc.y = beginSrc.y;
    for (int i=0; i<6; i++)
    {
        float distanceTmp = distance.x;
        float sumWeight = 0.0f;
        float4 sumColor = 0.0f;
        coordSrc.x = beginSrc.x;

        //水平方向ループ
        for (int j=0; j<6; j++)
        {
            float weight = lanczos3(distanceTmp);
            sumColor += weight * read_imagef(imgSrc, sampler, coordSrc);
            sumWeight += weight;
            coordSrc.x += 1.0f;;
            distanceTmp += 1.0f;
        }

        colorVert[i] = sumColor / sumWeight;
        coordSrc.y += 1.0f;
    }

    //colorVert[]を補完して書込む
    {
        float distanceTmp = distance.y;
        float sumWeight = 0.0f;
        float4 sumColor = 0.0f;

        for (int i=0; i<6; i++)
        {
            float weight = lanczos3(distanceTmp);
            sumColor += weight * colorVert[i];
            sumWeight += weight;
            distanceTmp += 1.0f;
        }

        write_imagef(imgDst, (int2)(coordDst.x, coordDst.y), sumColor / sumWeight);
    }
}
//////// end:

Radeon HD 6870 使用時に1280*720の画像を1920*1080に拡大する処理でローカルワークサイズを自動設定とした場合、平均38[ms]と中々に遅い。
ローカルワークサイズを(16, 12)とした場合は平均36[ms]程度であったが、誤差に近い。
次回は作業領域を使い2パスで処理して計算量を減らすカーネルを挙げる。

CodeXL

CodeXL – Powerful Debugging, Profiling & Analysis
http://developer.amd.com/tools-and-sdks/opencl-zone/opencl-tools-sdks/codexl/

カーネルのコンパイルを行うだけでも十分に有用だが、詰めたい場合に GPU: Performance Counter, Application Timeline Trace を見ると良さそうだ。

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++

Windowsでの画像ファイルの読書き

細かい事を気にしてはいけない。
Windowscodecs.lib, Shlwapi.lib をリンクすること。


#include <Windows.h>
#include <comdef.h>
#include <wincodec.h>
#include <Shlwapi.h>

_COM_SMARTPTR_TYPEDEF(IWICImagingFactory, __uuidof(IWICImagingFactory));
_COM_SMARTPTR_TYPEDEF(IWICBitmapDecoder, __uuidof(IWICBitmapDecoder));
_COM_SMARTPTR_TYPEDEF(IWICBitmapFrameDecode, __uuidof(IWICBitmapFrameDecode));
_COM_SMARTPTR_TYPEDEF(IWICBitmapSource, __uuidof(IWICBitmapSource));
_COM_SMARTPTR_TYPEDEF(IStream, __uuidof(IStream));
_COM_SMARTPTR_TYPEDEF(IWICBitmapEncoder, __uuidof(IWICBitmapEncoder));
_COM_SMARTPTR_TYPEDEF(IWICBitmapFrameEncode, __uuidof(IWICBitmapFrameEncode));


HRESULT read(void **ppData, UINT *pWidth, UINT *pHeight, LPCWSTR pName)
{
    IWICImagingFactoryPtr pFactory;
    IWICBitmapDecoderPtr pDecoder;
    IWICBitmapFrameDecodePtr pFrame;
    IWICBitmapSourcePtr pBitmap;
    UINT stride;
    UINT size;


    pFactory.CreateInstance(CLSID_WICImagingFactory, NULL, CLSCTX_INPROC_SERVER);
    pFactory->CreateDecoderFromFilename(pName, NULL, GENERIC_READ, WICDecodeMetadataCacheOnDemand, &pDecoder);
    pDecoder->GetFrame(0, &pFrame);
    WICConvertBitmapSource(GUID_WICPixelFormat32bppBGR, pFrame, &pBitmap); //DIB section compatible
    pBitmap->GetSize(pWidth, pHeight);
    stride = 4 * *pWidth;
    size = stride * *pHeight;
    *ppData = new BYTE[size];
    pBitmap->CopyPixels(NULL, stride, size, (BYTE *)*ppData);

    return S_OK;
}

HRESULT writeBmp(LPCWSTR pName, void *pData, UINT width, UINT height)
{
    IWICImagingFactoryPtr pFactory;
    IWICBitmapEncoderPtr pEncoder;
    IStreamPtr pStream;
    IWICBitmapFrameEncodePtr pFrame;
    WICPixelFormatGUID guidFormat;


    pFactory.CreateInstance(CLSID_WICImagingFactory, NULL, CLSCTX_INPROC_SERVER);
    pFactory->CreateEncoder(GUID_ContainerFormatBmp, NULL, &pEncoder);
    SHCreateStreamOnFileEx(pName, STGM_READWRITE| STGM_CREATE, FILE_ATTRIBUTE_NORMAL, TRUE, NULL, &pStream);
    pEncoder->Initialize(pStream, WICBitmapEncoderNoCache);
    pEncoder->CreateNewFrame(&pFrame, NULL);
    pFrame->Initialize(NULL);
    guidFormat = GUID_WICPixelFormat32bppBGR;
    pFrame->SetPixelFormat(&guidFormat);
    pFrame->SetSize(width, height);
    pFrame->WritePixels(height, 4*width, 4*width*height, (BYTE *)pData);
    pFrame->Commit();
    pEncoder->Commit();

    return S_OK;
}

OpenCLを使う準備

手元の環境が、
Windows 7 X64 SP1
Radeon HD 6870
なので、それに応じた環境を整える。

APP SDK
http://developer.amd.com/tools-and-sdks/opencl-zone/opencl-tools-sdks/amd-accelerated-parallel-processing-app-sdk/
AMD APP SDK の Windows-64 を拾ってくる。

Khronos OpenCL Registry
http://www.khronos.org/registry/cl/
OpenCL 1.2 API and C Language Specification
OpenCL 1.2 C++ Bindings Specification
OpenCL 1.2 Reference Card
を拾ってくる。

Microsoft Visual Studio Express 2013 for Windows Desktop
http://www.microsoft.com/ja-jp/download/details.aspx?id=40787

Microsoft Visual Studio 2013 Update 2
http://www.microsoft.com/ja-jp/download/details.aspx?id=42666

インストール及び設定については割愛する。