今更なEICAR

EICARテストファイル


0100 58            POP AX   ; AX=0
0101 354F21        XOR AX,214F  ; AX=214F
0104 50            PUSH AX   ; 214F
0105 254041        AND AX,4140  ; AX=0140
0108 50            PUSH AX   ; 0140 214F
0109 5B            POP BX   ; BX=0140, 214F
010A 345C          XOR AL,5C  ; AX=011C
010C 50            PUSH AX   ; 011C 214F
010D 5A            POP DX   ; DX=011C, 214F
010E 58            POP AX   ; AX=214F
010F 353428        XOR AX,2834  ; AX=097B
0112 50            PUSH AX   ; 097B
0113 5E            POP SI   ; SI=097B
0114 2937          SUB [BX],SI  ; [0140]=2B48-097B=21CD
0116 43            INC BX   ; BX=0141
0117 43            INC BX   ; BX=0142
0118 2937          SUB [BX],SI  ; [0142]=2A48-097B=20CD
011A 7D24          JGE 0140  ; JMP 0140: CD 21 CD 20, AH=09 DX=011C INT21, INT 20
011C 45            INC BP   ; '$'(24H) terminated output text.
011D 49            DEC CX
011E 43            INC BX
011F 41            INC CX
0120 52            PUSH DX
0121 2D5354        SUB AX,5453
0124 41            INC CX
0125 4E            DEC SI
0126 44            INC SP
0127 41            INC CX
0128 52            PUSH DX
0129 44            INC SP
012A 2D414E        SUB AX,4E41
012D 54            PUSH SP
012E 49            DEC CX
012F 56            PUSH SI
0130 49            DEC CX
0131 52            PUSH DX
0132 55            PUSH BP
0133 53            PUSH BX
0134 2D5445        SUB AX,4554
0137 53            PUSH BX
0138 54            PUSH SP
0139 2D4649        SUB AX,4946
013C 4C            DEC SP
013D 45            INC BP
013E 2124          AND [SI],SP
0140 48            DEC AX
0141 2B482A        SUB CX,[BX+SI+2A]

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