granicus.if.org Git - imagemagick/blob - magick/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-2010 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 "magick/studio.h"
  53 #include "magick/accelerate.h"
  54 #include "magick/artifact.h"
  55 #include "magick/cache-view.h"
  56 #include "magick/color-private.h"
  57 #include "magick/enhance.h"
  58 #include "magick/exception.h"
  59 #include "magick/exception-private.h"
  60 #include "magick/gem.h"
  61 #include "magick/hashmap.h"
  62 #include "magick/image.h"
  63 #include "magick/image-private.h"
  64 #include "magick/list.h"
  65 #include "magick/memory_.h"
  66 #include "magick/monitor-private.h"
  67 #include "magick/accelerate.h"
  68 #include "magick/option.h"
  69 #include "magick/pixel-private.h"
  70 #include "magick/prepress.h"
  71 #include "magick/quantize.h"
  72 #include "magick/registry.h"
  73 #include "magick/semaphore.h"
  74 #include "magick/splay-tree.h"
  75 #include "magick/statistic.h"
  76 #include "magick/string_.h"
  77 #include "magick/string-private.h"
  78 #include "magick/token.h"
  79 \f
  80 /*
  81 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  82 %                                                                             %
  83 %                                                                             %
  84 %                                                                             %
  85 %     A c c e l e r a t e C o n v o l v e I m a g e                           %
  86 %                                                                             %
  87 %                                                                             %
  88 %                                                                             %
  89 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  90 %
  91 %  AccelerateConvolveImage() applies a custom convolution kernel to the image.
  92 %  It is accelerated by taking advantage of speed-ups offered by executing in
  93 %  concert across heterogeneous platforms consisting of CPUs, GPUs, and other
  94 %  processors.
  95 %
  96 %  The format of the AccelerateConvolveImage method is:
  97 %
  98 %      Image *AccelerateConvolveImage(const Image *image,
  99 %        const KernelInfo *kernel,Image *convolve_image,
 100 %        ExceptionInfo *exception)
 101 %
 102 %  A description of each parameter follows:
 103 %
 104 %    o image: the image.
 105 %
 106 %    o kernel: the convolution kernel.
 107 %
 108 %    o convole_image: the convoleed image.
 109 %
 110 %    o exception: return any errors or warnings in this structure.
 111 %
 112 */
 113
 114 #if defined(MAGICKCORE_OPENCL_SUPPORT)
 115
 116 #if defined(MAGICKCORE_HDRI_SUPPORT)
 117 #define CLOptions "-DMAGICKCORE_HDRI_SUPPORT=1 -DCLQuantum=float " \
 118   "-DCLPixelType=float4 -DQuantumRange=%g -DMagickEpsilon=%g"
 119 #define CLPixelPacket  cl_float4
 120 #else
 121 #if (MAGICKCORE_QUANTUM_DEPTH == 8)
 122 #define CLOptions "-DCLQuantum=uchar -DCLPixelType=uchar4 " \
 123   "-DQuantumRange=%g -DMagickEpsilon=%g"
 124 #define CLPixelPacket  cl_uchar4
 125 #elif (MAGICKCORE_QUANTUM_DEPTH == 16)
 126 #define CLOptions "-DCLQuantum=ushort -DCLPixelType=ushort4 " \
 127   "-DQuantumRange=%g -DMagickEpsilon=%g"
 128 #define CLPixelPacket  cl_ushort4
 129 #elif (MAGICKCORE_QUANTUM_DEPTH == 32)
 130 #define CLOptions "-DCLQuantum=uint -DCLPixelType=uint4 " \
 131   "-DQuantumRange=%g -DMagickEpsilon=%g"
 132 #define CLPixelPacket  cl_uint4
 133 #elif (MAGICKCORE_QUANTUM_DEPTH == 32)
 134 #define CLOptions "-DCLQuantum=ussize_t -DCLPixelType=ussize_t4 " \
 135   "-DQuantumRange=%g -DMagickEpsilon=%g"
 136 #define CLPixelPacket  cl_ulong4
 137 #endif
 138 #endif
 139
 140 typedef struct _ConvolveInfo
 141 {
 142   cl_context
 143     context;
 144
 145   cl_device_id
 146     *devices;
 147
 148   cl_command_queue
 149     command_queue;
 150
 151   cl_kernel
 152     kernel;
 153
 154   cl_program
 155     program;
 156
 157   cl_mem
 158     pixels,
 159     convolve_pixels;
 160
 161   cl_ulong
 162     width,
 163     height;
 164
 165   cl_bool
 166     matte;
 167
 168   cl_mem
 169     filter;
 170 } ConvolveInfo;
 171
 172 static char
 173   *ConvolveKernel =
 174     "static inline long ClampToCanvas(const long offset,const unsigned long range)\n"
 175     "{\n"
 176     "  if (offset < 0L)\n"
 177     "    return(0L);\n"
 178     "  if (offset >= range)\n"
 179     "    return((long) (range-1L));\n"
 180     "  return(offset);\n"
 181     "}\n"
 182     "\n"
 183     "static inline CLQuantum ClampToQuantum(const double value)\n"
 184     "{\n"
 185     "#if defined(MAGICKCORE_HDRI_SUPPORT)\n"
 186     "  return((CLQuantum) value)\n"
 187     "#else\n"
 188     "  if (value < 0.0)\n"
 189     "    return((CLQuantum) 0);\n"
 190     "  if (value >= (double) QuantumRange)\n"
 191     "    return((CLQuantum) QuantumRange);\n"
 192     "  return((CLQuantum) (value+0.5));\n"
 193     "#endif\n"
 194     "}\n"
 195     "\n"
 196     "__kernel void Convolve(const __global CLPixelType *input,\n"
 197     "  __constant double *filter,const unsigned long width,const unsigned long height,\n"
 198     "  const bool matte,__global CLPixelType *output)\n"
 199     "{\n"
 200     "  const unsigned long columns = get_global_size(0);\n"
 201     "  const unsigned long rows = get_global_size(1);\n"
 202     "\n"
 203     "  const long x = get_global_id(0);\n"
 204     "  const long y = get_global_id(1);\n"
 205     "\n"
 206     "  const double scale = (1.0/QuantumRange);\n"
 207     "  const long mid_width = (width-1)/2;\n"
 208     "  const long mid_height = (height-1)/2;\n"
 209     "  double4 sum = { 0.0, 0.0, 0.0, 0.0 };\n"
 210     "  double gamma = 0.0;\n"
 211     "  register unsigned long i = 0;\n"
 212     "\n"
 213     "  int method = 0;\n"
 214     "  if (matte != false)\n"
 215     "    method=1;\n"
 216     "  if ((x >= width) && (x < (columns-width-1)) &&\n"
 217     "      (y >= height) && (y < (rows-height-1)))\n"
 218     "    {\n"
 219     "      method=2;\n"
 220     "      if (matte != false)\n"
 221     "        method=3;\n"
 222     "    }\n"
 223     "  switch (method)\n"
 224     "  {\n"
 225     "    case 0:\n"
 226     "    {\n"
 227     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 228     "      {\n"
 229     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 230     "        {\n"
 231     "          const long index=ClampToCanvas(y+v,rows)*columns+\n"
 232     "            ClampToCanvas(x+u,columns);\n"
 233     "          sum.x+=filter[i]*input[index].x;\n"
 234     "          sum.y+=filter[i]*input[index].y;\n"
 235     "          sum.z+=filter[i]*input[index].z;\n"
 236     "          gamma+=filter[i];\n"
 237     "          i++;\n"
 238     "        }\n"
 239     "      }\n"
 240     "      break;\n"
 241     "    }\n"
 242     "    case 1:\n"
 243     "    {\n"
 244     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 245     "      {\n"
 246     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 247     "        {\n"
 248     "          const unsigned long index=ClampToCanvas(y+v,rows)*columns+\n"
 249     "            ClampToCanvas(x+u,columns);\n"
 250     "          const double alpha=scale*(QuantumRange-input[index].w);\n"
 251     "          sum.x+=alpha*filter[i]*input[index].x;\n"
 252     "          sum.y+=alpha*filter[i]*input[index].y;\n"
 253     "          sum.z+=alpha*filter[i]*input[index].z;\n"
 254     "          sum.w+=filter[i]*input[index].w;\n"
 255     "          gamma+=alpha*filter[i];\n"
 256     "          i++;\n"
 257     "        }\n"
 258     "      }\n"
 259     "      break;\n"
 260     "    }\n"
 261     "    case 2:\n"
 262     "    {\n"
 263     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 264     "      {\n"
 265     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 266     "        {\n"
 267     "          const unsigned long index=(y+v)*columns+(x+u);\n"
 268     "          sum.x+=filter[i]*input[index].x;\n"
 269     "          sum.y+=filter[i]*input[index].y;\n"
 270     "          sum.z+=filter[i]*input[index].z;\n"
 271     "          gamma+=filter[i];\n"
 272     "          i++;\n"
 273     "        }\n"
 274     "      }\n"
 275     "      break;\n"
 276     "    }\n"
 277     "    case 3:\n"
 278     "    {\n"
 279     "      for (long v=(-mid_height); v <= mid_height; v++)\n"
 280     "      {\n"
 281     "        for (long u=(-mid_width); u <= mid_width; u++)\n"
 282     "        {\n"
 283     "          const unsigned long index=(y+v)*columns+(x+u);\n"
 284     "          const double alpha=scale*(QuantumRange-input[index].w);\n"
 285     "          sum.x+=alpha*filter[i]*input[index].x;\n"
 286     "          sum.y+=alpha*filter[i]*input[index].y;\n"
 287     "          sum.z+=alpha*filter[i]*input[index].z;\n"
 288     "          sum.w+=filter[i]*input[index].w;\n"
 289     "          gamma+=alpha*filter[i];\n"
 290     "          i++;\n"
 291     "        }\n"
 292     "      }\n"
 293     "      break;\n"
 294     "    }\n"
 295     "  }\n"
 296     "  gamma=1.0/(fabs(gamma) <= MagickEpsilon ? 1.0 : gamma);\n"
 297     "  const unsigned long index = y*columns+x;\n"
 298     "  output[index].x=ClampToQuantum(gamma*sum.x);\n"
 299     "  output[index].y=ClampToQuantum(gamma*sum.y);\n"
 300     "  output[index].z=ClampToQuantum(gamma*sum.z);\n"
 301     "  if (matte == false)\n"
 302     "    output[index].w=input[index].w;\n"
 303     "  else\n"
 304     "    output[index].w=ClampToQuantum(sum.w);\n"
 305     "}\n";
 306
 307 static void ConvolveNotify(const char *message,const void *data,size_t length,
 308   void *user_context)
 309 {
 310   ExceptionInfo
 311     *exception;
 312
 313   (void) data;
 314   (void) length;
 315   exception=(ExceptionInfo *) user_context;
 316   (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 317     "DelegateFailed","`%s'",message);
 318 }
 319
 320 static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
 321   const Image *image,const void *pixels,double *filter,
 322   const size_t width,const size_t height,void *convolve_pixels)
 323 {
 324   cl_int
 325     status;
 326
 327   register cl_uint
 328     i;
 329
 330   size_t
 331     length;
 332
 333   /*
 334     Allocate OpenCL buffers.
 335   */
 336   length=image->columns*image->rows;
 337   convolve_info->pixels=clCreateBuffer(convolve_info->context,(cl_mem_flags)
 338     (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(CLPixelPacket),
 339     (void *) pixels,&status);
 340   if ((convolve_info->pixels == (cl_mem) NULL) || (status != CL_SUCCESS))
 341     return(MagickFalse);
 342   length=width*height;
 343   convolve_info->filter=clCreateBuffer(convolve_info->context,(cl_mem_flags)
 344     (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(cl_double),filter,
 345     &status);
 346   if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS))
 347     return(MagickFalse);
 348   length=image->columns*image->rows;
 349   convolve_info->convolve_pixels=clCreateBuffer(convolve_info->context,
 350     (cl_mem_flags) (CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR),length*
 351     sizeof(CLPixelPacket),convolve_pixels,&status);
 352   if ((convolve_info->convolve_pixels == (cl_mem) NULL) ||
 353       (status != CL_SUCCESS))
 354     return(MagickFalse);
 355   /*
 356     Bind OpenCL buffers.
 357   */
 358   i=0;
 359   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
 360     &convolve_info->pixels);
 361   if (status != CL_SUCCESS)
 362     return(MagickFalse);
 363   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
 364     &convolve_info->filter);
 365   if (status != CL_SUCCESS)
 366     return(MagickFalse);
 367   convolve_info->width=(cl_ulong) width;
 368   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
 369     &convolve_info->width);
 370   if (status != CL_SUCCESS)
 371     return(MagickFalse);
 372   convolve_info->height=(cl_ulong) height;
 373   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
 374     &convolve_info->height);
 375   if (status != CL_SUCCESS)
 376     return(MagickFalse);
 377   convolve_info->matte=(cl_bool) image->matte;
 378   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_bool),(void *)
 379     &convolve_info->matte);
 380   if (status != CL_SUCCESS)
 381     return(MagickFalse);
 382   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
 383     &convolve_info->convolve_pixels);
 384   if (status != CL_SUCCESS)
 385     return(MagickFalse);
 386   status=clFinish(convolve_info->command_queue);
 387   if (status != CL_SUCCESS)
 388     return(MagickFalse);
 389   return(MagickTrue);
 390 }
 391
 392 static void DestroyConvolveBuffers(ConvolveInfo *convolve_info)
 393 {
 394   cl_int
 395     status;
 396
 397   if (convolve_info->convolve_pixels != (cl_mem) NULL)
 398     status=clReleaseMemObject(convolve_info->convolve_pixels);
 399   if (convolve_info->pixels != (cl_mem) NULL)
 400     status=clReleaseMemObject(convolve_info->pixels);
 401   if (convolve_info->filter != (cl_mem) NULL)
 402     status=clReleaseMemObject(convolve_info->filter);
 403 }
 404
 405 static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info)
 406 {
 407   cl_int
 408     status;
 409
 410   if (convolve_info->kernel != (cl_kernel) NULL)
 411     status=clReleaseKernel(convolve_info->kernel);
 412   if (convolve_info->program != (cl_program) NULL)
 413     status=clReleaseProgram(convolve_info->program);
 414   if (convolve_info->command_queue != (cl_command_queue) NULL)
 415     status=clReleaseCommandQueue(convolve_info->command_queue);
 416   if (convolve_info->context != (cl_context) NULL)
 417     status=clReleaseContext(convolve_info->context);
 418   convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info);
 419   return(convolve_info);
 420 }
 421
 422 static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info,
 423   const Image *image,const void *pixels,double *filter,
 424   const size_t width,const size_t height,void *convolve_pixels)
 425 {
 426   cl_int
 427     status;
 428
 429   size_t
 430     global_work_size[2],
 431     length;
 432
 433   length=image->columns*image->rows;
 434   status=clEnqueueWriteBuffer(convolve_info->command_queue,
 435     convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),pixels,0,NULL,
 436     NULL);
 437   length=width*height;
 438   status=clEnqueueWriteBuffer(convolve_info->command_queue,
 439     convolve_info->filter,CL_TRUE,0,length*sizeof(cl_double),filter,0,NULL,
 440     NULL);
 441   if (status != CL_SUCCESS)
 442     return(MagickFalse);
 443   global_work_size[0]=image->columns;
 444   global_work_size[1]=image->rows;
 445   status=clEnqueueNDRangeKernel(convolve_info->command_queue,
 446     convolve_info->kernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
 447   if (status != CL_SUCCESS)
 448     return(MagickFalse);
 449   length=image->columns*image->rows;
 450   status=clEnqueueReadBuffer(convolve_info->command_queue,
 451     convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),
 452     convolve_pixels,0,NULL,NULL);
 453   if (status != CL_SUCCESS)
 454     return(MagickFalse);
 455   status=clFinish(convolve_info->command_queue);
 456   if (status != CL_SUCCESS)
 457     return(MagickFalse);
 458   return(MagickTrue);
 459 }
 460
 461 static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
 462   const char *source,ExceptionInfo *exception)
 463 {
 464   char
 465     options[MaxTextExtent];
 466
 467   cl_int
 468     status;
 469
 470   ConvolveInfo
 471     *convolve_info;
 472
 473   size_t
 474     length,
 475     lengths[] = { strlen(source) };
 476
 477   /*
 478     Create OpenCL info.
 479   */
 480   convolve_info=(ConvolveInfo *) AcquireAlignedMemory(1,sizeof(*convolve_info));
 481   if (convolve_info == (ConvolveInfo *) NULL)
 482     {
 483       (void) ThrowMagickException(exception,GetMagickModule(),
 484         ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
 485       return((ConvolveInfo *) NULL);
 486     }
 487   (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info));
 488   /*
 489     Create OpenCL context.
 490   */
 491   convolve_info->context=clCreateContextFromType((cl_context_properties *)
 492     NULL,(cl_device_type) CL_DEVICE_TYPE_GPU,ConvolveNotify,exception,&status);
 493   if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
 494     convolve_info->context=clCreateContextFromType((cl_context_properties *)
 495       NULL,(cl_device_type) CL_DEVICE_TYPE_CPU,ConvolveNotify,exception,
 496       &status);
 497   if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
 498     convolve_info->context=clCreateContextFromType((cl_context_properties *)
 499       NULL,(cl_device_type) CL_DEVICE_TYPE_DEFAULT,ConvolveNotify,exception,
 500       &status);
 501   if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
 502     {
 503       (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 504         "failed to create OpenCL context","`%s' (%d)",image->filename,status);
 505       convolve_info=DestroyConvolveInfo(convolve_info);
 506       return((ConvolveInfo *) NULL);
 507     }
 508   /*
 509     Detect OpenCL devices.
 510   */
 511   status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL,
 512     &length);
 513   if ((status != CL_SUCCESS) || (length == 0))
 514     {
 515       convolve_info=DestroyConvolveInfo(convolve_info);
 516       return((ConvolveInfo *) NULL);
 517     }
 518   convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length);
 519   if (convolve_info->devices == (cl_device_id *) NULL)
 520     {
 521       (void) ThrowMagickException(exception,GetMagickModule(),
 522         ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
 523       convolve_info=DestroyConvolveInfo(convolve_info);
 524       return((ConvolveInfo *) NULL);
 525     }
 526   status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length,
 527     convolve_info->devices,NULL);
 528   if (status != CL_SUCCESS)
 529     {
 530       convolve_info=DestroyConvolveInfo(convolve_info);
 531       return((ConvolveInfo *) NULL);
 532     }
 533   /*
 534     Create OpenCL command queue.
 535   */
 536   convolve_info->command_queue=clCreateCommandQueue(convolve_info->context,
 537     convolve_info->devices[0],0,&status);
 538   if ((convolve_info->command_queue == (cl_command_queue) NULL) ||
 539       (status != CL_SUCCESS))
 540     {
 541       convolve_info=DestroyConvolveInfo(convolve_info);
 542       return((ConvolveInfo *) NULL);
 543     }
 544   /*
 545     Build OpenCL program.
 546   */
 547   convolve_info->program=clCreateProgramWithSource(convolve_info->context,1,
 548     &source,lengths,&status);
 549   if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
 550     {
 551       convolve_info=DestroyConvolveInfo(convolve_info);
 552       return((ConvolveInfo *) NULL);
 553     }
 554   (void) FormatMagickString(options,MaxTextExtent,CLOptions,(double)
 555     QuantumRange,MagickEpsilon);
 556   status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options,
 557     NULL,NULL);
 558   if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
 559     {
 560       char
 561         *log;
 562
 563       status=clGetProgramBuildInfo(convolve_info->program,
 564         convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,0,NULL,&length);
 565       log=(char *) AcquireMagickMemory(length);
 566       if (log == (char *) NULL)
 567         {
 568           convolve_info=DestroyConvolveInfo(convolve_info);
 569           return((ConvolveInfo *) NULL);
 570         }
 571       status=clGetProgramBuildInfo(convolve_info->program,
 572         convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length);
 573       (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
 574         "failed to build OpenCL program","`%s' (%s)",image->filename,log);
 575       log=DestroyString(log);
 576       convolve_info=DestroyConvolveInfo(convolve_info);
 577       return((ConvolveInfo *) NULL);
 578     }
 579   /*
 580     Get a kernel object.
 581   */
 582   convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status);
 583   if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS))
 584     {
 585       convolve_info=DestroyConvolveInfo(convolve_info);
 586       return((ConvolveInfo *) NULL);
 587     }
 588   return(convolve_info);
 589 }
 590
 591 #endif
 592
 593 MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
 594   const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception)
 595 {
 596   assert(image != (Image *) NULL);
 597   assert(image->signature == MagickSignature);
 598   if (image->debug != MagickFalse)
 599     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
 600   assert(kernel != (KernelInfo *) NULL);
 601   assert(kernel->signature == MagickSignature);
 602   assert(convolve_image != (Image *) NULL);
 603   assert(convolve_image->signature == MagickSignature);
 604   assert(exception != (ExceptionInfo *) NULL);
 605   assert(exception->signature == MagickSignature);
 606   if ((image->storage_class != DirectClass) ||
 607       (image->colorspace == CMYKColorspace))
 608   if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
 609       (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
 610     return(MagickFalse);
 611 #if !defined(MAGICKCORE_OPENCL_SUPPORT)
 612   return(MagickFalse);
 613 #else
 614   {
 615     const void
 616       *pixels;
 617
 618     ConvolveInfo
 619       *convolve_info;
 620
 621     MagickBooleanType
 622       status;
 623
 624     MagickSizeType
 625       length;
 626
 627     void
 628       *convolve_pixels;
 629
 630     convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception);
 631     if (convolve_info == (ConvolveInfo *) NULL)
 632       return(MagickFalse);
 633     pixels=AcquirePixelCachePixels(image,&length,exception);
 634     if (pixels == (const void *) NULL)
 635       {
 636         (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
 637           "UnableToReadPixelCache","`%s'",image->filename);
 638         convolve_info=DestroyConvolveInfo(convolve_info);
 639         return(MagickFalse);
 640       }
 641     convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception);
 642     if (convolve_pixels == (void *) NULL)
 643       {
 644         (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
 645           "UnableToReadPixelCache","`%s'",image->filename);
 646         convolve_info=DestroyConvolveInfo(convolve_info);
 647         return(MagickFalse);
 648       }
 649     status=BindConvolveParameters(convolve_info,image,pixels,kernel->values,
 650       kernel->width,kernel->height,convolve_pixels);
 651     if (status == MagickFalse)
 652       {
 653         DestroyConvolveBuffers(convolve_info);
 654         convolve_info=DestroyConvolveInfo(convolve_info);
 655         return(MagickFalse);
 656       }
 657     status=EnqueueConvolveKernel(convolve_info,image,pixels,kernel->values,
 658       kernel->width,kernel->height,convolve_pixels);
 659     if (status == MagickFalse)
 660       {
 661         DestroyConvolveBuffers(convolve_info);
 662         convolve_info=DestroyConvolveInfo(convolve_info);
 663         return(MagickFalse);
 664       }
 665     DestroyConvolveBuffers(convolve_info);
 666     convolve_info=DestroyConvolveInfo(convolve_info);
 667     return(MagickTrue);
 668   }
 669 #endif
 670 }