A probléma megoldásához szükségünk lesz két képobjektumra és egy kernelt tartalmazó memória területre, amelyetz egy egy dimenziós tömbben adunk át és a gazdagép main függvényében állítunk elő.

Ezenkívűl szükségünk lesz még egy mintavételező objektumra, amely segítségével tudjuk lekérdezni a betöltött kép pixel intenzitás értékeit.

Természetesen itt is érdemes először ellenőrizni, hogy a kiválasztott eszköz tudja-e kezelni a képobjektumokat.

Mindenek előtt be kell töltenünk az input képet. Ehhez a következő függvényt használjuk

cl_mem LoadImage(cl_context context, char *fileName, int &width, int &height)
{
  FREE_IMAGE_FORMAT format = FreeImage_GetFileType(fileName, 0);
  FIBITMAP* image = FreeImage_Load(format, fileName);

  FIBITMAP* temp = image;
  image = FreeImage_ConvertTo32Bits(image);
  FreeImage_Unload(temp);

  width = FreeImage_GetWidth(image);
  height = FreeImage_GetHeight(image);

  char *buffer = new char[width * height * 4];
  memcpy(buffer, FreeImage_GetBits(image), width * height * 4);

  FreeImage_Unload(image);

  cl_image_format clImageFormat;
  clImageFormat.image_channel_order = CL_BGRA;
  clImageFormat.image_channel_data_type = CL_UNORM_INT8;

  cl_int errNum;
  cl_mem clImage;
  clImage = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &clImageFormat, width, height, 0, buffer, &errNum);

  if (errNum != CL_SUCCESS)
  {
    std::cerr << "Error creating CL image object" << std::endl;
    return 0;
  }

  return clImage;
}

Az input kép betöltéséhez a függvényt a következő képpen hívhatjuk meg a main függvényen belül.

cl_mem imgObjects[2] = { 0, 0 };
int width, height;
imgObjects[0] = LoadImage(context, "d:/MyProjects/OpenCL/OpenCL_conv/input.png", width, height);
if (imgObjects[0] == 0)
{
  std::cerr << "Error loading: " << std::string(argv[1]) << std::endl;
  Cleanup(context, commandQueue, program, kernel, imgObjects, sampler, kernelBuffer);
  return 1;
}

A kép mintavételezéséhez szükségünk lesz egy mintavételező objektumra.

cl_sampler sampler = 0;
sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, 	CL_FILTER_NEAREST, &errNum);

Hozzuk létre az eredményt tároló képobjektumot a következőképpen!

cl_image_format clImageFormat;
clImageFormat.image_channel_order = CL_BGRA;
clImageFormat.image_channel_data_type = CL_UNORM_INT8;
imgObjects[1] = clCreateImage2D(context, CL_MEM_WRITE_ONLY, &clImageFormat, N, N, 0, NULL, &errNum);

Definiáljuk a mask függvényt!

const int kernelWidth = 5;
const int sum = 25;

cl_float  kernelMask[kernelWidth * kernelWidth] =
{
  1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum,
  1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum,
  1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum,
  1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum,
  1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum, 1.0f / sum
};

cl_mem kernelBuffer = 0;
kernelBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY |  CL_MEM_COPY_HOST_PTR, 
                               sizeof(cl_uint) * kernelWidth * kernelWidth, kernelMask, &errNum);

Ezekután már csak meg kell adnunk a kernel argumentumokat.

errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &imgObjects[0]);	
errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &imgObjects[1]);
errNum |= clSetKernelArg(kernel, 2, sizeof(cl_sampler), &sampler);
errNum |= clSetKernelArg(kernel, 3, sizeof(cl_int), &width);
errNum |= clSetKernelArg(kernel, 4, sizeof(cl_int), &height);
errNum |= clSetKernelArg(kernel, 5, sizeof(cl_mem), &kernelBuffer);
errNum |= clSetKernelArg(kernel, 6, sizeof(cl_int), &kernelWidth);

Végezetül a kernel végrehajtása után az eredmény képobjektum tartalmát a gazdagép memóriájába másoljuk és elmentjük egy png típusú állományba.

unsigned char* result = new  unsigned char[width * height * 4];
size_t origin[3] = { 0, 0, 0 };
size_t region[3] = { width, height, 1 };

errNum = clEnqueueReadImage(commandQueue, imgObjects[1], CL_TRUE, origin, region, 0, 0, result, 0, NULL, NULL);

SaveImage("Output.bmp", result, width, height);

A teljes forrás letölthető itt.

Az eszköz kernel kódja a következő.

// Add you device OpenCL code

__kernel void smoothing_filter(__read_only image2d_t srcImg,
                               __write_only image2d_t dstImg,
                               sampler_t sampler,
                               const int width, const int height, 
                               __constant float *const kernelWeights, 
                               const int kernelWidth)
{
  int2 outImageCoord = (int2) (get_global_id(0), get_global_id(1));
// A kezdő és vég képkoordináták az input képen a kernel méretét figyelembe véve
  int2 startImageCoord = (int2) (outImageCoord.x - (kernelWidth-1)/2, outImageCoord.y - (kernelWidth-1)/2);
  int2 endImageCoord   = (int2) (outImageCoord.x + (kernelWidth-1)/2, outImageCoord.y + (kernelWidth+1)/2);

// A kimeneteti koordináta ellenőrzése    
  if (outImageCoord.x < width && outImageCoord.y < height)
  {
// A kernel aktuális súlyának a pozíciója az egy-dimenziós tömbben
    int weight = 0;
    float4 outColor = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
// for ciklus, ami bejárja a kernel területét
// A mintavételezés beállítása miatt a külső koordináták is le vannak kezelve
    for( int y = startImageCoord.y; y <= endImageCoord.y; y++)
    {
       for( int x = startImageCoord.x; x <= endImageCoord.x; x++)
       {
          outColor += (read_imagef(srcImg, sampler, (int2)(x, y)) * (kernelWeights[weight]));
          weight += 1;
       }
   }      
   write_imagef(dstImg, outImageCoord, outColor);
  }
}

A forrás letölthető itt.