OpenCL clCreateImage2D Sample

//================================
// OpenCL Kernel Function for element by element pixel access
//================================
__kernel void VidCapPixelAccess(__read_only image2d_t imageIn,__write_only image2d_t imageOut)
{
  const sampler_t sampler=CLK_NORMALIZED_COORDS_FALSE
  CLK_ADDRESS_CLAMP
  CLK_FILTER_NEAREST;

  int gid0 = get_global_id(0); // x
  int gid1 = get_global_id(1); // y

  uint4 pixel;

  pixel=read_imageui(imageIn,sampler,(int2)(gid0,gid1));
  write_imageui (imageOut,(int2)(gid0,gid1),pixel);
}



//================================
// OpenCL Main Program Segment
//================================
void *imageIn;
//void *imageOut;
cl_mem myClImageInBuffer;
cl_mem myClImageOutBuffer;
//imageOut = (void *)malloc(8 * (640*480*4));

//====> Allocate the OpenCL buffer memory object IMAGE
cl_image_format format;
format.image_channel_order = CL_RGBA;
format.image_channel_data_type = CL_UNSIGNED_INT8;

cl_mem_flags flags1 = CL_MEM_READ_ONLY;
myClImageInBuffer = clCreateImage2D (
  cxGPUContext,
  flags1,
  &format,
  640,
  480,
  0,
  0, //imageIn
  &ciErr1);

cl_mem_flags flags2 = CL_MEM_READ_WRITE;
myClImageOutBuffer = clCreateImage2D (
  cxGPUContext,
  flags2,
  &format,
  640,
  480,
  0,
  0, //imageOut
  &ciErr2);
ciErr1 = clSetKernelArg( ckKernel, 0, sizeof(cl_mem), (void*)&myClImageInBuffer);
ciErr1 = clSetKernelArg( ckKernel, 1, sizeof(cl_mem), (void*)&myClImageOutBuffer);

// 將 BMP imagedata 放入 image buffer 中
Rect rect(0, 0, 640, 480);
BitmapData* pBitmapData = new BitmapData;
pBitmap->LockBits(
  &rect,
  ImageLockModeRead,
  PixelFormat32bppARGB,
  pBitmapData
);
imageIn = pBitmapData->Scan0;
size_t origin[3] = {0, 0, 0};
size_t region[3] = {640, 480, 1}; //width, height depth

// Asynchronous write of data to GPU device (將來源:imageBuffer 放入 GPU: myClImageInBuffer)
// The row pitch is the number of bytes in each row
// Thus: width * 4 bytes each * sizeof(byte)
ciErr1 = clEnqueueWriteImage(
    cqCommandQueue, // command queue
    myClImageInBuffer, // to which buffer
    CL_TRUE, // whether is a blocking or non-blocking read
    origin, // const size_t origin[3]
    region, // const size_t region[3]
    640*sizeof(unsigned char) * 4, //width is bytes
    0,
    imageIn,
    0, // number events in wait list
    NULL, // event wait list pointer
    NULL); // event pointer

size_t szGlobalWorkSize[] = {640, 480};
size_t szLocalWorkSize[] = {16, 16};
// Launch kernel
// clEnqueueNDRangeKernel 會把執行一個 kernel 的動作加到 command queue 裡面
ciErr1 = clEnqueueNDRangeKernel(
    cqCommandQueue, // command queue
    ckKernel, // Which Kernel
    2, // Choose if we are using 1D, 2D or 3D work-items and work-group
    NULL, // Global work offset
    szGlobalWorkSize, // The total number of work-items (must have work_dim dimensions)
    szLocalWorkSize, // The number of work-items per work-group (must have work_dim dimensions)
    0, // number events in wait list
    NULL, // event wait list pointer
    NULL); // event pointer

// Synchronous/blocking read of results, and check accumulated errors
// 將 image buffer 放入 BMP imagedata 中
if (pOutBitmap == 0)
  pOutBitmap = ::new Gdiplus::Bitmap(640, 480, PixelFormat32bppRGB); //PixelFormat32bppARGB
BitmapData* pOutBitmapData = new BitmapData;
Rect outRect(0, 0, 640, 480);
pOutBitmap->LockBits(
  &outRect,
  ImageLockModeWrite,
  PixelFormat32bppARGB,
  pOutBitmapData
);

// clEnqueueReadBuffer 函式會把「將記憶體資料從 OpenCL 裝置複製到主記憶體」的動作加到 command queue 中
// The row pitch is the number of bytes in each row
// Thus: width * 4 bytes each * sizeof(byte)
ciErr1 = clEnqueueReadImage( cqCommandQueue, // command queue
  myClImageOutBuffer, // to which buffer
  CL_TRUE, // whether is a blocking or non-blocking read
  origin, // const size_t origin[3]
  region, // const size_t region[3]
  640*sizeof(unsigned char) * 4, //width is bytes
  0,
  pOutBitmapData->Scan0, //imageOut
  0, // number events in wait list
  NULL, // event wait list pointer
  NULL); // event pointer
pOutBitmap->UnlockBits(pOutBitmapData);

g.DrawImage(pOutBitmap, Gdiplus::Rect(0, 0, r.right, r.bottom));

//==== DONE ====