granicus.if.org Git - imagemagick/blob - MagickCore/accelerate.c

   1 /*
   2 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
   3 %                                                                             %
   4 %                                                                             %
   5 %                                                                             %
   6 %     AAA     CCCC    CCCC  EEEEE  L      EEEEE  RRRR    AAA   TTTTT  EEEEE   %
   7 %    A   A   C       C      E      L      E      R   R  A   A    T    E       %
   8 %    AAAAA   C       C      EEE    L      EEE    RRRR   AAAAA    T    EEE     %
   9 %    A   A   C       C      E      L      E      R R    A   A    T    E       %
  10 %    A   A    CCCC    CCCC  EEEEE  LLLLL  EEEEE  R  R   A   A    T    EEEEE   %
  11 %                                                                             %
  12 %                                                                             %
  13 %                       MagickCore Acceleration Methods                       %
  14 %                                                                             %
  15 %                              Software Design                                %
  16 %                               John Cristy                                   %
  17 %                               January 2010                                  %
  18 %                                                                             %
  19 %                                                                             %
  20 %  Copyright 1999-2013 ImageMagick Studio LLC, a non-profit organization      %
  21 %  dedicated to making software imaging solutions freely available.           %
  22 %                                                                             %
  23 %  You may not use this file except in compliance with the License.  You may  %
  24 %  obtain a copy of the License at                                            %
  25 %                                                                             %
  26 %    http://www.imagemagick.org/script/license.php                            %
  27 %                                                                             %
  28 %  Unless required by applicable law or agreed to in writing, software        %
  29 %  distributed under the License is distributed on an "AS IS" BASIS,          %
  30 %  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.   %
  31 %  See the License for the specific language governing permissions and        %
  32 %  limitations under the License.                                             %
  33 %                                                                             %
  34 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  35 %
  36 % Morphology is the the application of various kernals, of any size and even
  37 % shape, to a image in various ways (typically binary, but not always).
  38 %
  39 % Convolution (weighted sum or average) is just one specific type of
  40 % accelerate. Just one that is very common for image bluring and sharpening
  41 % effects.  Not only 2D Gaussian blurring, but also 2-pass 1D Blurring.
  42 %
  43 % This module provides not only a general accelerate function, and the ability
  44 % to apply more advanced or iterative morphologies, but also functions for the
  45 % generation of many different types of kernel arrays from user supplied
  46 % arguments. Prehaps even the generation of a kernel from a small image.
  47 */
  48 \f
  49 /*
  50   Include declarations.
  51 */
  52 #include "MagickCore/studio.h"
  53 #include "MagickCore/accelerate.h"
  54 #include "MagickCore/artifact.h"
  55 #include "MagickCore/cache.h"
  56 #include "MagickCore/cache-private.h"
  57 #include "MagickCore/cache-view.h"
  58 #include "MagickCore/color-private.h"
  59 #include "MagickCore/enhance.h"
  60 #include "MagickCore/exception.h"
  61 #include "MagickCore/exception-private.h"
  62 #include "MagickCore/gem.h"
  63 #include "MagickCore/hashmap.h"
  64 #include "MagickCore/image.h"
  65 #include "MagickCore/image-private.h"
  66 #include "MagickCore/list.h"
  67 #include "MagickCore/memory_.h"
  68 #include "MagickCore/monitor-private.h"
  69 #include "MagickCore/accelerate.h"
  70 #include "MagickCore/option.h"
  71 #include "MagickCore/pixel-accessor.h"
  72 #include "MagickCore/prepress.h"
  73 #include "MagickCore/quantize.h"
  74 #include "MagickCore/registry.h"
  75 #include "MagickCore/semaphore.h"
  76 #include "MagickCore/splay-tree.h"
  77 #include "MagickCore/statistic.h"
  78 #include "MagickCore/string_.h"
  79 #include "MagickCore/string-private.h"
  80 #include "MagickCore/token.h"
  81 \f
  82 /*
  83 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  84 %                                                                             %
  85 %                                                                             %
  86 %                                                                             %
  87 %     A c c e l e r a t e C o n v o l v e I m a g e                           %
  88 %                                                                             %
  89 %                                                                             %
  90 %                                                                             %
  91 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  92 %
  93 %  AccelerateConvolveImage() applies a custom convolution kernel to the image.
  94 %  It is accelerated by taking advantage of speed-ups offered by executing in
  95 %  concert across heterogeneous platforms consisting of CPUs, GPUs, and other
  96 %  processors.
  97 %
  98 %  The format of the AccelerateConvolveImage method is:
  99 %
 100 %      Image *AccelerateConvolveImage(const Image *image,
 101 %        const KernelInfo *kernel,Image *convolve_image,
 102 %        ExceptionInfo *exception)
 103 %
 104 %  A description of each parameter follows:
 105 %
 106 %    o image: the image.
 107 %
 108 %    o kernel: the convolution kernel.
 109 %
 110 %    o convole_image: the convoleed image.
 111 %
 112 %    o exception: return any errors or warnings in this structure.
 113 %
 114 */
 115
 116 #if defined(MAGICKCORE_OPENCL_SUPPORT)
 117
 118 #if defined(MAGICKCORE_HDRI_SUPPORT)
 119 #define CLOptions "-DMAGICKCORE_HDRI_SUPPORT=1 -DCLQuantum=float " \
 120   "-DCLPixelType=float4 -DQuantumRange=%g -DMagickEpsilon=%g"
 121 #define CLPixelInfo  cl_float4
 122 #else
 123 #if (MAGICKCORE_QUANTUM_DEPTH == 8)
 124 #define CLOptions "-DCLQuantum=uchar -DCLPixelType=uchar4 " \
 125   "-DQuantumRange=%g -DMagickEpsilon=%g"
 126 #define CLPixelInfo  cl_uchar4
 127 #elif (MAGICKCORE_QUANTUM_DEPTH == 16)
 128 #define CLOptions "-DCLQuantum=ushort -DCLPixelType=ushort4 " \
 129   "-DQuantumRange=%g -DMagickEpsilon=%g"
 130 #define CLPixelInfo  cl_ushort4
 131 #elif (MAGICKCORE_QUANTUM_DEPTH == 32)
 132 #define CLOptions "-DCLQuantum=uint -DCLPixelType=uint4 " \
 133   "-DQuantumRange=%g -DMagickEpsilon=%g"
 134 #define CLPixelInfo  cl_uint4
 135 #elif (MAGICKCORE_QUANTUM_DEPTH == 64)
 136 #define CLOptions "-DCLQuantum=ussize_t -DCLPixelType=ussize_t4 " \
 137   "-DQuantumRange=%g -DMagickEpsilon=%g"
 138 #define CLPixelInfo  cl_ulong4
 139 #endif
 140 #endif
 141
 142 typedef struct _ConvolveInfo
 143 {
 144   cl_context
 145     context;
 146
 147   cl_device_id
 148     *devices;
 149
 150   cl_command_queue
 151     command_queue;
 152
 153   cl_kernel
 154     kernel;
 155
 156   cl_program
 157     program;
 158
 159   cl_mem
 160     pixels,
 161     convolve_pixels;
 162
 163   cl_ulong
 164     width,
 165     height;
 166
 167   cl_uint
 168     matte;
 169
 170   cl_mem
 171     filter;
 172 } ConvolveInfo;
 173
 174 static const char
 175   *ConvolveKernel =
 176     "static inline long ClampToCanvas(const long offset,const unsigned long range)\n"
 177     "{\n"
 178     "  if (offset < 0L)\n"
 179     "    return(0L);\n"
 180     "  if (offset >= range)\n"
 181     "    return((long) (range-1L));\n"
 182     "  return(offset);\n"
 183     "}\n"
 184     "\n"
 185     "static inline CLQuantum ClampToQuantum(const float value)\n"
 186     "{\n"
 187     "#if defined(MAGICKCORE_HDRI_SUPPORT)\n"
 188     "  return((CLQuantum) value);\n"
 189     "#else\n"
 190     "  if (value < 0.0)\n"
 191     "    return((CLQuantum) 0);\n"
 192     "  if (value >= (float) QuantumRange)\n"
 193     "    return((CLQuantum) QuantumRange);\n"
 194     "  return((CLQuantum) (value+0.5));\n"
 195     "#endif\n"
 196     "}\n"
 197     "\n"
 198     "static inline float PerceptibleReciprocal(const float x)\n"
 199     "{\n"
 200     "  float sign = x < (float) 0.0 ? (float) -1.0 : (float) 1.0;\n"
 201     "  return((sign*x) >= MagickEpsilon ? (float) 1.0/x : sign*((float) 1.0/\n"
 202     "    MagickEpsilon));\n"
 203     "}\n"
 204     "\n"
 205     "__kernel void Convolve(const __global CLPixelType *input,\n"
 206     "  __constant float *filter,const unsigned long width,const unsigned long height,\n"
 207     "  const unsigned int matte,__global CLPixelType *output)\n"
 208     "{\n"
 209     "  const unsigned long columns = get_global_size(0);\n"
 210     "  const unsigned long rows = get_global_size(1);\n"
 211     "\n"
 212     "  const long x = get_global_id(0);\n"
 213     "  const long y = get_global_id(1);\n"
 214     "\n"
 215     "  const float scale = (1.0/QuantumRange);\n"
 216     "  const long mid_width = (width-1)/2;\n"
 217     "  const long mid_height = (height-1)/2;\n"
 218     "  float4 sum = { 0.0, 0.0, 0.0, 0.0 };\n"
 219     "  float gamma = 0.0;\n"
 220     "  register unsigned long i = 0;\n"
 221     "\n"
 222     "  int method = 0;\n"
 223     "  if (matte != false)\n"
 224     "    method=1;\n"
 225     "  if ((x >= width) && (x < (columns-width-1)) &&\n"
 226     "      (y >= height) && (y < (rows-height-1)))\n"
 227     "    {\n"
 228     "      method=2;\n"
 229     "      if (matte != false)\n"
 230     "        method=3;\n"
 231     "    }\n"
 232     "  switch (method)\n"
 233     "  {\n"
 234     "    case 0:\n"
 235     "    {\n"
 236     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 237     "      {\n"
 238     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 239     "        {\n"
 240     "          const long index=ClampToCanvas(y+v,rows)*columns+\n"
 241     "            ClampToCanvas(x+u,columns);\n"
 242     "          sum.x+=filter[i]*input[index].x;\n"
 243     "          sum.y+=filter[i]*input[index].y;\n"
 244     "          sum.z+=filter[i]*input[index].z;\n"
 245     "          gamma+=filter[i];\n"
 246     "          i++;\n"
 247     "        }\n"
 248     "      }\n"
 249     "      break;\n"
 250     "    }\n"
 251     "    case 1:\n"
 252     "    {\n"
 253     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 254     "      {\n"
 255     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 256     "        {\n"
 257     "          const unsigned long index=ClampToCanvas(y+v,rows)*columns+\n"
 258     "            ClampToCanvas(x+u,columns);\n"
 259     "          const float alpha=scale*input[index].w;\n"
 260     "          sum.x+=alpha*filter[i]*input[index].x;\n"
 261     "          sum.y+=alpha*filter[i]*input[index].y;\n"
 262     "          sum.z+=alpha*filter[i]*input[index].z;\n"
 263     "          sum.w+=filter[i]*input[index].w;\n"
 264     "          gamma+=alpha*filter[i];\n"
 265     "          i++;\n"
 266     "        }\n"
 267     "      }\n"
 268     "      break;\n"
 269     "    }\n"
 270     "    case 2:\n"
 271     "    {\n"
 272     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 273     "      {\n"
 274     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 275     "        {\n"
 276     "          const unsigned long index=(y+v)*columns+(x+u);\n"
 277     "          sum.x+=filter[i]*input[index].x;\n"
 278     "          sum.y+=filter[i]*input[index].y;\n"
 279     "          sum.z+=filter[i]*input[index].z;\n"
 280     "          gamma+=filter[i];\n"
 281     "          i++;\n"
 282     "        }\n"
 283     "      }\n"
 284     "      break;\n"
 285     "    }\n"
 286     "    case 3:\n"
 287     "    {\n"
 288     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 289     "      {\n"
 290     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 291     "        {\n"
 292     "          const unsigned long index=(y+v)*columns+(x+u);\n"
 293     "          const float alpha=scale*input[index].w;\n"
 294     "          sum.x+=alpha*filter[i]*input[index].x;\n"
 295     "          sum.y+=alpha*filter[i]*input[index].y;\n"
 296     "          sum.z+=alpha*filter[i]*input[index].z;\n"
 297     "          sum.w+=filter[i]*input[index].w;\n"
 298     "          gamma+=alpha*filter[i];\n"
 299     "          i++;\n"
 300     "        }\n"
 301     "      }\n"
 302     "      break;\n"
 303     "    }\n"
 304     "  }\n"
 305     "  gamma=PerceptibleReciprocal(gamma);\n"
 306     "  const unsigned long index = y*columns+x;\n"
 307     "  output[index].x=ClampToQuantum(gamma*sum.x);\n"
 308     "  output[index].y=ClampToQuantum(gamma*sum.y);\n"
 309     "  output[index].z=ClampToQuantum(gamma*sum.z);\n"
 310     "  if (matte == false)\n"
 311     "    output[index].w=input[index].w;\n"
 312     "  else\n"
 313     "    output[index].w=ClampToQuantum(sum.w);\n"
 314     "}\n";
 315
 316 static void ConvolveNotify(const char *message,const void *data,size_t length,
 317   void *user_context)
 318 {
 319   ExceptionInfo
 320     *exception;
 321
 322   (void) data;
 323   (void) length;
 324   exception=(ExceptionInfo *) user_context;
 325   (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 326     "DelegateFailed","`%s'",message);
 327 }
 328
 329 static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
 330   const Image *image,const void *pixels,float *filter,const size_t width,
 331   const size_t height,void *convolve_pixels)
 332 {
 333   cl_int
 334     status;
 335
 336   register cl_uint
 337     i;
 338
 339   size_t
 340     length;
 341
 342   /*
 343     Allocate OpenCL buffers.
 344   */
 345   length=image->columns*image->rows;
 346   convolve_info->pixels=clCreateBuffer(convolve_info->context,(cl_mem_flags)
 347     (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(CLPixelInfo),
 348     (void *) pixels,&status);
 349   if ((convolve_info->pixels == (cl_mem) NULL) || (status != CL_SUCCESS))
 350     return(MagickFalse);
 351   length=width*height;
 352   convolve_info->filter=clCreateBuffer(convolve_info->context,(cl_mem_flags)
 353     (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(cl_float),filter,
 354     &status);
 355   if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS))
 356     return(MagickFalse);
 357   length=image->columns*image->rows;
 358   convolve_info->convolve_pixels=clCreateBuffer(convolve_info->context,
 359     (cl_mem_flags) (CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR),length*
 360     sizeof(CLPixelInfo),convolve_pixels,&status);
 361   if ((convolve_info->convolve_pixels == (cl_mem) NULL) ||
 362       (status != CL_SUCCESS))
 363     return(MagickFalse);
 364   /*
 365     Bind OpenCL buffers.
 366   */
 367   i=0;
 368   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
 369     &convolve_info->pixels);
 370   if (status != CL_SUCCESS)
 371     return(MagickFalse);
 372   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
 373     &convolve_info->filter);
 374   if (status != CL_SUCCESS)
 375     return(MagickFalse);
 376   convolve_info->width=(cl_ulong) width;
 377   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
 378     &convolve_info->width);
 379   if (status != CL_SUCCESS)
 380     return(MagickFalse);
 381   convolve_info->height=(cl_ulong) height;
 382   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
 383     &convolve_info->height);
 384   if (status != CL_SUCCESS)
 385     return(MagickFalse);
 386   convolve_info->matte=(cl_uint) image->alpha_trait == BlendPixelTrait ?
 387     MagickTrue : MagickFalse;
 388   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_uint),(void *)
 389     &convolve_info->matte);
 390   if (status != CL_SUCCESS)
 391     return(MagickFalse);
 392   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
 393     &convolve_info->convolve_pixels);
 394   if (status != CL_SUCCESS)
 395     return(MagickFalse);
 396   status=clFinish(convolve_info->command_queue);
 397   if (status != CL_SUCCESS)
 398     return(MagickFalse);
 399   return(MagickTrue);
 400 }
 401
 402 static void DestroyConvolveBuffers(ConvolveInfo *convolve_info)
 403 {
 404   cl_int
 405     status;
 406
 407   status=0;
 408   if (convolve_info->convolve_pixels != (cl_mem) NULL)
 409     status=clReleaseMemObject(convolve_info->convolve_pixels);
 410   if (convolve_info->pixels != (cl_mem) NULL)
 411     status=clReleaseMemObject(convolve_info->pixels);
 412   if (convolve_info->filter != (cl_mem) NULL)
 413     status=clReleaseMemObject(convolve_info->filter);
 414   (void) status;
 415 }
 416
 417 static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info)
 418 {
 419   cl_int
 420     status;
 421
 422   status=0;
 423   if (convolve_info->kernel != (cl_kernel) NULL)
 424     status=clReleaseKernel(convolve_info->kernel);
 425   if (convolve_info->program != (cl_program) NULL)
 426     status=clReleaseProgram(convolve_info->program);
 427   if (convolve_info->command_queue != (cl_command_queue) NULL)
 428     status=clReleaseCommandQueue(convolve_info->command_queue);
 429   if (convolve_info->context != (cl_context) NULL)
 430     status=clReleaseContext(convolve_info->context);
 431   (void) status;
 432   convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info);
 433   return(convolve_info);
 434 }
 435
 436 static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info,
 437   const Image *image,const void *pixels,float *filter,const size_t width,
 438   const size_t height,void *convolve_pixels)
 439 {
 440   cl_int
 441     status;
 442
 443   size_t
 444     global_work_size[2],
 445     length;
 446
 447   length=image->columns*image->rows;
 448   status=clEnqueueWriteBuffer(convolve_info->command_queue,
 449     convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelInfo),pixels,0,NULL,
 450     NULL);
 451   length=width*height;
 452   status=clEnqueueWriteBuffer(convolve_info->command_queue,
 453     convolve_info->filter,CL_TRUE,0,length*sizeof(cl_float),filter,0,NULL,
 454     NULL);
 455   if (status != CL_SUCCESS)
 456     return(MagickFalse);
 457   global_work_size[0]=image->columns;
 458   global_work_size[1]=image->rows;
 459   status=clEnqueueNDRangeKernel(convolve_info->command_queue,
 460     convolve_info->kernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
 461   if (status != CL_SUCCESS)
 462     return(MagickFalse);
 463   length=image->columns*image->rows;
 464   status=clEnqueueReadBuffer(convolve_info->command_queue,
 465     convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelInfo),
 466     convolve_pixels,0,NULL,NULL);
 467   if (status != CL_SUCCESS)
 468     return(MagickFalse);
 469   status=clFinish(convolve_info->command_queue);
 470   if (status != CL_SUCCESS)
 471     return(MagickFalse);
 472   return(MagickTrue);
 473 }
 474
 475 static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
 476   const char *source,ExceptionInfo *exception)
 477 {
 478   char
 479     options[MaxTextExtent];
 480
 481   cl_context_properties
 482     context_properties[3];
 483
 484   cl_int
 485     status;
 486
 487   cl_platform_id
 488     platforms[1];
 489
 490   cl_uint
 491     number_platforms;
 492
 493   ConvolveInfo
 494     *convolve_info;
 495
 496   size_t
 497     length,
 498     lengths[] = { strlen(source) };
 499
 500   /*
 501     Create OpenCL info.
 502   */
 503   convolve_info=(ConvolveInfo *) AcquireMagickMemory(sizeof(*convolve_info));
 504   if (convolve_info == (ConvolveInfo *) NULL)
 505     {
 506       (void) ThrowMagickException(exception,GetMagickModule(),
 507         ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
 508       return((ConvolveInfo *) NULL);
 509     }
 510   (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info));
 511   /*
 512     Create OpenCL context.
 513   */
 514   status=clGetPlatformIDs(0,(cl_platform_id *) NULL,&number_platforms);
 515   if ((status == CL_SUCCESS) && (number_platforms > 0))
 516     status=clGetPlatformIDs(1,platforms,NULL);
 517   if (status != CL_SUCCESS)
 518     {
 519       (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 520         "failed to create OpenCL context","'%s' (%d)",image->filename,status);
 521       convolve_info=DestroyConvolveInfo(convolve_info);
 522       return((ConvolveInfo *) NULL);
 523     }
 524   context_properties[0]=CL_CONTEXT_PLATFORM;
 525   context_properties[1]=(cl_context_properties) platforms[0];
 526   context_properties[2]=0;
 527   convolve_info->context=clCreateContextFromType(context_properties,
 528     (cl_device_type) CL_DEVICE_TYPE_GPU,ConvolveNotify,exception,&status);
 529   if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
 530     convolve_info->context=clCreateContextFromType(context_properties,
 531       (cl_device_type) CL_DEVICE_TYPE_CPU,ConvolveNotify,exception,&status);
 532   if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
 533     convolve_info->context=clCreateContextFromType(context_properties,
 534       (cl_device_type) CL_DEVICE_TYPE_DEFAULT,ConvolveNotify,exception,&status);
 535   if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
 536     {
 537       (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 538         "failed to create OpenCL context","'%s' (%d)",image->filename,status);
 539       convolve_info=DestroyConvolveInfo(convolve_info);
 540       return((ConvolveInfo *) NULL);
 541     }
 542   /*
 543     Detect OpenCL devices.
 544   */
 545   status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL,
 546     &length);
 547   if ((status != CL_SUCCESS) || (length == 0))
 548     {
 549       convolve_info=DestroyConvolveInfo(convolve_info);
 550       return((ConvolveInfo *) NULL);
 551     }
 552   convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length);
 553   if (convolve_info->devices == (cl_device_id *) NULL)
 554     {
 555       (void) ThrowMagickException(exception,GetMagickModule(),
 556         ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
 557       convolve_info=DestroyConvolveInfo(convolve_info);
 558       return((ConvolveInfo *) NULL);
 559     }
 560   status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length,
 561     convolve_info->devices,NULL);
 562   if (status != CL_SUCCESS)
 563     {
 564       convolve_info=DestroyConvolveInfo(convolve_info);
 565       return((ConvolveInfo *) NULL);
 566     }
 567   if (image->debug != MagickFalse)
 568     {
 569       char
 570         attribute[MaxTextExtent];
 571
 572       size_t
 573         length;
 574
 575       clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_NAME,
 576         sizeof(attribute),attribute,&length);
 577       (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Name: %s",
 578         attribute);
 579       clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VENDOR,
 580         sizeof(attribute),attribute,&length);
 581       (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Vendor: %s",
 582         attribute);
 583       clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VERSION,
 584         sizeof(attribute),attribute,&length);
 585       (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
 586         "Driver Version: %s",attribute);
 587       clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_PROFILE,
 588         sizeof(attribute),attribute,&length);
 589       (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Profile: %s",
 590         attribute);
 591       clGetDeviceInfo(convolve_info->devices[0],CL_DRIVER_VERSION,
 592         sizeof(attribute),attribute,&length);
 593       (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Driver: %s",
 594         attribute);
 595       clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_EXTENSIONS,
 596         sizeof(attribute),attribute,&length);
 597       (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Extensions: %s",
 598         attribute);
 599     }
 600   /*
 601     Create OpenCL command queue.
 602   */
 603   convolve_info->command_queue=clCreateCommandQueue(convolve_info->context,
 604     convolve_info->devices[0],0,&status);
 605   if ((convolve_info->command_queue == (cl_command_queue) NULL) ||
 606       (status != CL_SUCCESS))
 607     {
 608       convolve_info=DestroyConvolveInfo(convolve_info);
 609       return((ConvolveInfo *) NULL);
 610     }
 611   /*
 612     Build OpenCL program.
 613   */
 614   convolve_info->program=clCreateProgramWithSource(convolve_info->context,1,
 615     &source,lengths,&status);
 616   if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
 617     {
 618       convolve_info=DestroyConvolveInfo(convolve_info);
 619       return((ConvolveInfo *) NULL);
 620     }
 621   (void) FormatLocaleString(options,MaxTextExtent,CLOptions,(float)
 622     QuantumRange,MagickEpsilon);
 623   status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options,
 624     NULL,NULL);
 625   if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
 626     {
 627       char
 628         *log;
 629
 630       status=clGetProgramBuildInfo(convolve_info->program,
 631         convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,0,NULL,&length);
 632       log=(char *) AcquireMagickMemory(length);
 633       if (log == (char *) NULL)
 634         {
 635           convolve_info=DestroyConvolveInfo(convolve_info);
 636           return((ConvolveInfo *) NULL);
 637         }
 638       status=clGetProgramBuildInfo(convolve_info->program,
 639         convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length);
 640       (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 641         "failed to build OpenCL program","'%s' (%s)",image->filename,log);
 642       log=DestroyString(log);
 643       convolve_info=DestroyConvolveInfo(convolve_info);
 644       return((ConvolveInfo *) NULL);
 645     }
 646   /*
 647     Get a kernel object.
 648   */
 649   convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status);
 650   if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS))
 651     {
 652       convolve_info=DestroyConvolveInfo(convolve_info);
 653       return((ConvolveInfo *) NULL);
 654     }
 655   return(convolve_info);
 656 }
 657
 658 #endif
 659
 660 MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
 661   const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception)
 662 {
 663   assert(image != (Image *) NULL);
 664   assert(image->signature == MagickSignature);
 665   if (image->debug != MagickFalse)
 666     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
 667   assert(kernel != (KernelInfo *) NULL);
 668   assert(kernel->signature == MagickSignature);
 669   assert(convolve_image != (Image *) NULL);
 670   assert(convolve_image->signature == MagickSignature);
 671   assert(exception != (ExceptionInfo *) NULL);
 672   assert(exception->signature == MagickSignature);
 673   if ((image->storage_class != DirectClass) ||
 674       (image->colorspace == CMYKColorspace))
 675     return(MagickFalse);
 676   if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
 677       (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
 678     return(MagickFalse);
 679   if (GetPixelChannels(image) != 4)
 680     return(MagickFalse);
 681 #if !defined(MAGICKCORE_OPENCL_SUPPORT)
 682   return(MagickFalse);
 683 #else
 684   {
 685     const void
 686       *pixels;
 687
 688     float
 689       *filter;
 690
 691     ConvolveInfo
 692       *convolve_info;
 693
 694     MagickBooleanType
 695       status;
 696
 697     MagickSizeType
 698       length;
 699
 700     register ssize_t
 701       i;
 702
 703     void
 704       *convolve_pixels;
 705
 706     convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception);
 707     if (convolve_info == (ConvolveInfo *) NULL)
 708       return(MagickFalse);
 709     pixels=AcquirePixelCachePixels(image,&length,exception);
 710     if (pixels == (const void *) NULL)
 711       {
 712         convolve_info=DestroyConvolveInfo(convolve_info);
 713         (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
 714           "UnableToReadPixelCache","`%s'",image->filename);
 715         return(MagickFalse);
 716       }
 717     convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception);
 718     if (convolve_pixels == (void *) NULL)
 719       {
 720         convolve_info=DestroyConvolveInfo(convolve_info);
 721         (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
 722           "UnableToReadPixelCache","`%s'",image->filename);
 723         return(MagickFalse);
 724       }
 725     filter=(float *) AcquireQuantumMemory(kernel->width,kernel->height*
 726       sizeof(*filter));
 727     if (filter == (float *) NULL)
 728       {
 729         DestroyConvolveBuffers(convolve_info);
 730         convolve_info=DestroyConvolveInfo(convolve_info);
 731         (void) ThrowMagickException(exception,GetMagickModule(),
 732           ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
 733         return(MagickFalse);
 734       }
 735     for (i=0; i < (ssize_t) (kernel->width*kernel->height); i++)
 736       filter[i]=(float) kernel->values[i];
 737     status=BindConvolveParameters(convolve_info,image,pixels,filter,
 738       kernel->width,kernel->height,convolve_pixels);
 739     if (status == MagickFalse)
 740       {
 741         filter=(float *) RelinquishMagickMemory(filter);
 742         DestroyConvolveBuffers(convolve_info);
 743         convolve_info=DestroyConvolveInfo(convolve_info);
 744         return(MagickFalse);
 745       }
 746     status=EnqueueConvolveKernel(convolve_info,image,pixels,filter,
 747       kernel->width,kernel->height,convolve_pixels);
 748     filter=(float *) RelinquishMagickMemory(filter);
 749     if (status == MagickFalse)
 750       {
 751         DestroyConvolveBuffers(convolve_info);
 752         convolve_info=DestroyConvolveInfo(convolve_info);
 753         return(MagickFalse);
 754       }
 755     DestroyConvolveBuffers(convolve_info);
 756     convolve_info=DestroyConvolveInfo(convolve_info);
 757     return(MagickTrue);
 758   }
 759 #endif
 760 }