2 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
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 %
13 % MagickCore Acceleration Methods %
20 % Copyright 1999-2012 ImageMagick Studio LLC, a non-profit organization %
21 % dedicated to making software imaging solutions freely available. %
23 % You may not use this file except in compliance with the License. You may %
24 % obtain a copy of the License at %
26 % http://www.imagemagick.org/script/license.php %
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. %
34 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
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).
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.
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.
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"
83 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
87 % A c c e l e r a t e C o n v o l v e I m a g e %
91 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
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
98 % The format of the AccelerateConvolveImage method is:
100 % Image *AccelerateConvolveImage(const Image *image,
101 % const KernelInfo *kernel,Image *convolve_image,
102 % ExceptionInfo *exception)
104 % A description of each parameter follows:
106 % o image: the image.
108 % o kernel: the convolution kernel.
110 % o convole_image: the convoleed image.
112 % o exception: return any errors or warnings in this structure.
116 #if defined(MAGICKCORE_OPENCL_SUPPORT)
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
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
142 typedef struct _ConvolveInfo
176 "static inline long ClampToCanvas(const long offset,const unsigned long range)\n"
178 " if (offset < 0L)\n"
180 " if (offset >= range)\n"
181 " return((long) (range-1L));\n"
185 "static inline CLQuantum ClampToQuantum(const float value)\n"
187 "#if defined(MAGICKCORE_HDRI_SUPPORT)\n"
188 " return((CLQuantum) value);\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"
198 "static inline float PerceptibleReciprocal(const float x)\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"
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"
209 " const unsigned long columns = get_global_size(0);\n"
210 " const unsigned long rows = get_global_size(1);\n"
212 " const long x = get_global_id(0);\n"
213 " const long y = get_global_id(1);\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"
223 " if (matte != false)\n"
225 " if ((x >= width) && (x < (columns-width-1)) &&\n"
226 " (y >= height) && (y < (rows-height-1)))\n"
229 " if (matte != false)\n"
236 " for (long v=(-mid_height); v <= mid_height; v++)\n"
238 " for (long u=(-mid_width); u <= mid_width; u++)\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"
253 " for (long v=(-mid_height); v <= mid_height; v++)\n"
255 " for (long u=(-mid_width); u <= mid_width; u++)\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"
272 " for (long v=(-mid_height); v <= mid_height; v++)\n"
274 " for (long u=(-mid_width); u <= mid_width; u++)\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"
288 " for (long v=(-mid_height); v <= mid_height; v++)\n"
290 " for (long u=(-mid_width); u <= mid_width; u++)\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"
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"
313 " output[index].w=ClampToQuantum(sum.w);\n"
316 static void ConvolveNotify(const char *message,const void *data,size_t length,
324 exception=(ExceptionInfo *) user_context;
325 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
326 "DelegateFailed","'%s'",message);
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)
343 Allocate OpenCL buffers.
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))
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,
355 if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS))
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))
368 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
369 &convolve_info->pixels);
370 if (status != CL_SUCCESS)
372 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
373 &convolve_info->filter);
374 if (status != CL_SUCCESS)
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)
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)
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)
392 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
393 &convolve_info->convolve_pixels);
394 if (status != CL_SUCCESS)
396 status=clFinish(convolve_info->command_queue);
397 if (status != CL_SUCCESS)
402 static void DestroyConvolveBuffers(ConvolveInfo *convolve_info)
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);
417 static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info)
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);
432 convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info);
433 return(convolve_info);
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)
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,
452 status=clEnqueueWriteBuffer(convolve_info->command_queue,
453 convolve_info->filter,CL_TRUE,0,length*sizeof(cl_float),filter,0,NULL,
455 if (status != CL_SUCCESS)
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)
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)
469 status=clFinish(convolve_info->command_queue);
470 if (status != CL_SUCCESS)
475 static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
476 const char *source,ExceptionInfo *exception)
479 options[MaxTextExtent];
481 cl_context_properties
482 context_properties[3];
498 lengths[] = { strlen(source) };
503 convolve_info=(ConvolveInfo *) AcquireMagickMemory(sizeof(*convolve_info));
504 if (convolve_info == (ConvolveInfo *) NULL)
506 (void) ThrowMagickException(exception,GetMagickModule(),
507 ResourceLimitError,"MemoryAllocationFailed","'%s'",image->filename);
508 return((ConvolveInfo *) NULL);
510 (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info));
512 Create OpenCL context.
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)
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);
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))
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);
543 Detect OpenCL devices.
545 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL,
547 if ((status != CL_SUCCESS) || (length == 0))
549 convolve_info=DestroyConvolveInfo(convolve_info);
550 return((ConvolveInfo *) NULL);
552 convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length);
553 if (convolve_info->devices == (cl_device_id *) NULL)
555 (void) ThrowMagickException(exception,GetMagickModule(),
556 ResourceLimitError,"MemoryAllocationFailed","'%s'",image->filename);
557 convolve_info=DestroyConvolveInfo(convolve_info);
558 return((ConvolveInfo *) NULL);
560 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length,
561 convolve_info->devices,NULL);
562 if (status != CL_SUCCESS)
564 convolve_info=DestroyConvolveInfo(convolve_info);
565 return((ConvolveInfo *) NULL);
567 if (image->debug != MagickFalse)
570 attribute[MaxTextExtent];
575 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_NAME,
576 sizeof(attribute),attribute,&length);
577 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Name: %s",
579 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VENDOR,
580 sizeof(attribute),attribute,&length);
581 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Vendor: %s",
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",
591 clGetDeviceInfo(convolve_info->devices[0],CL_DRIVER_VERSION,
592 sizeof(attribute),attribute,&length);
593 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Driver: %s",
595 clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_EXTENSIONS,
596 sizeof(attribute),attribute,&length);
597 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Extensions: %s",
601 Create OpenCL command queue.
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))
608 convolve_info=DestroyConvolveInfo(convolve_info);
609 return((ConvolveInfo *) NULL);
612 Build OpenCL program.
614 convolve_info->program=clCreateProgramWithSource(convolve_info->context,1,
615 &source,lengths,&status);
616 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
618 convolve_info=DestroyConvolveInfo(convolve_info);
619 return((ConvolveInfo *) NULL);
621 (void) FormatLocaleString(options,MaxTextExtent,CLOptions,(float)
622 QuantumRange,MagickEpsilon);
623 status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options,
625 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
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)
635 convolve_info=DestroyConvolveInfo(convolve_info);
636 return((ConvolveInfo *) NULL);
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);
649 convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status);
650 if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS))
652 convolve_info=DestroyConvolveInfo(convolve_info);
653 return((ConvolveInfo *) NULL);
655 return(convolve_info);
660 MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
661 const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception)
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))
676 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
677 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
679 if (GetPixelChannels(image) != 4)
681 #if !defined(MAGICKCORE_OPENCL_SUPPORT)
706 convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception);
707 if (convolve_info == (ConvolveInfo *) NULL)
709 pixels=AcquirePixelCachePixels(image,&length,exception);
710 if (pixels == (const void *) NULL)
712 convolve_info=DestroyConvolveInfo(convolve_info);
713 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
714 "UnableToReadPixelCache","'%s'",image->filename);
717 convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception);
718 if (convolve_pixels == (void *) NULL)
720 convolve_info=DestroyConvolveInfo(convolve_info);
721 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
722 "UnableToReadPixelCache","'%s'",image->filename);
725 filter=(float *) AcquireQuantumMemory(kernel->width,kernel->height*
727 if (filter == (float *) NULL)
729 DestroyConvolveBuffers(convolve_info);
730 convolve_info=DestroyConvolveInfo(convolve_info);
731 (void) ThrowMagickException(exception,GetMagickModule(),
732 ResourceLimitError,"MemoryAllocationFailed","'%s'",image->filename);
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)
741 filter=(float *) RelinquishMagickMemory(filter);
742 DestroyConvolveBuffers(convolve_info);
743 convolve_info=DestroyConvolveInfo(convolve_info);
746 status=EnqueueConvolveKernel(convolve_info,image,pixels,filter,
747 kernel->width,kernel->height,convolve_pixels);
748 filter=(float *) RelinquishMagickMemory(filter);
749 if (status == MagickFalse)
751 DestroyConvolveBuffers(convolve_info);
752 convolve_info=DestroyConvolveInfo(convolve_info);
755 DestroyConvolveBuffers(convolve_info);
756 convolve_info=DestroyConvolveInfo(convolve_info);