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-2010 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 % Morpology 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 "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"
81 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
85 % A c c e l e r a t e C o n v o l v e I m a g e %
89 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
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
96 % The format of the AccelerateConvolveImage method is:
98 % Image *AccelerateConvolveImage(const Image *image,
99 % const KernelInfo *kernel,Image *convolve_image,
100 % ExceptionInfo *exception)
102 % A description of each parameter follows:
104 % o image: the image.
106 % o kernel: the convolution kernel.
108 % o convole_image: the convoleed image.
110 % o exception: return any errors or warnings in this structure.
114 #if defined(MAGICKCORE_OPENCL_SUPPORT)
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
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=ulong -DCLPixelType=ulong4 " \
135 "-DQuantumRange=%g -DMagickEpsilon=%g"
136 #define CLPixelPacket cl_ulong4
140 typedef struct _ConvolveInfo
174 "static inline long ClampToCanvas(const long offset,const ulong range)\n"
176 " if (offset < 0L)\n"
178 " if (offset >= range)\n"
179 " return((long) (range-1L));\n"
183 "static inline CLQuantum ClampToQuantum(const double value)\n"
185 "#if defined(MAGICKCORE_HDRI_SUPPORT)\n"
186 " return((CLQuantum) value)\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"
196 "__kernel void Convolve(const __global CLPixelType *input,\n"
197 " __constant double *filter,const ulong width,const ulong height,\n"
198 " const bool matte,__global CLPixelType *output)\n"
200 " const ulong columns = get_global_size(0);\n"
201 " const ulong rows = get_global_size(1);\n"
203 " const long x = get_global_id(0);\n"
204 " const long y = get_global_id(1);\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 ulong i = 0;\n"
214 " if (matte != false)\n"
216 " if ((x >= width) && (x < (columns-width-1)) &&\n"
217 " (y >= height) && (y < (rows-height-1)))\n"
220 " if (matte != false)\n"
227 " for (long v=(-mid_height); v <= mid_height; v++)\n"
229 " for (long u=(-mid_width); u <= mid_width; u++)\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"
244 " for (long v=(-mid_height); v <= mid_height; v++)\n"
246 " for (long u=(-mid_width); u <= mid_width; u++)\n"
248 " const ulong 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"
263 " for (long v=(-mid_height); v <= mid_height; v++)\n"
265 " for (long u=(-mid_width); u <= mid_width; u++)\n"
267 " const ulong 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"
279 " for (long v=(-mid_height); v <= mid_height; v++)\n"
281 " for (long u=(-mid_width); u <= mid_width; u++)\n"
283 " const ulong 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"
296 " gamma=1.0/(fabs(gamma) <= MagickEpsilon ? 1.0 : gamma);\n"
297 " const ulong 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"
304 " output[index].w=ClampToQuantum(sum.w);\n"
307 static void ConvolveNotify(const char *message,const void *data,size_t length,
315 exception=(ExceptionInfo *) user_context;
316 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
317 "DelegateFailed","`%s'",message);
320 static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
321 const Image *image,const void *pixels,double *filter,
322 const unsigned long width,const unsigned long height,void *convolve_pixels)
334 Allocate OpenCL buffers.
336 length=image->columns*image->rows;
337 convolve_info->pixels=clCreateBuffer(convolve_info->context,CL_MEM_READ_ONLY |
338 CL_MEM_USE_HOST_PTR,length*sizeof(CLPixelPacket),(void *) pixels,&status);
339 if ((convolve_info->pixels == (cl_mem) NULL) || (status != CL_SUCCESS))
342 convolve_info->filter=clCreateBuffer(convolve_info->context,CL_MEM_READ_ONLY |
343 CL_MEM_USE_HOST_PTR,length*sizeof(cl_double),filter,&status);
344 if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS))
346 length=image->columns*image->rows;
347 convolve_info->convolve_pixels=clCreateBuffer(convolve_info->context,
348 CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,length*sizeof(CLPixelPacket),
349 convolve_pixels,&status);
350 if ((convolve_info->convolve_pixels == (cl_mem) NULL) ||
351 (status != CL_SUCCESS))
357 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
358 &convolve_info->pixels);
359 if (status != CL_SUCCESS)
361 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
362 &convolve_info->filter);
363 if (status != CL_SUCCESS)
365 convolve_info->width=(cl_ulong) width;
366 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
367 &convolve_info->width);
368 if (status != CL_SUCCESS)
370 convolve_info->height=(cl_ulong) height;
371 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_ulong),(void *)
372 &convolve_info->height);
373 if (status != CL_SUCCESS)
375 convolve_info->matte=(cl_bool) image->matte;
376 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_bool),(void *)
377 &convolve_info->matte);
378 if (status != CL_SUCCESS)
380 status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_mem),(void *)
381 &convolve_info->convolve_pixels);
382 if (status != CL_SUCCESS)
384 status=clFinish(convolve_info->command_queue);
385 if (status != CL_SUCCESS)
390 static void DestroyConvolveBuffers(ConvolveInfo *convolve_info)
395 if (convolve_info->convolve_pixels != (cl_mem) NULL)
396 status=clReleaseMemObject(convolve_info->convolve_pixels);
397 if (convolve_info->pixels != (cl_mem) NULL)
398 status=clReleaseMemObject(convolve_info->pixels);
399 if (convolve_info->filter != (cl_mem) NULL)
400 status=clReleaseMemObject(convolve_info->filter);
403 static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info)
408 if (convolve_info->kernel != (cl_kernel) NULL)
409 status=clReleaseKernel(convolve_info->kernel);
410 if (convolve_info->program != (cl_program) NULL)
411 status=clReleaseProgram(convolve_info->program);
412 if (convolve_info->command_queue != (cl_command_queue) NULL)
413 status=clReleaseCommandQueue(convolve_info->command_queue);
414 if (convolve_info->context != (cl_context) NULL)
415 status=clReleaseContext(convolve_info->context);
416 convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info);
417 return(convolve_info);
420 static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info,
421 const Image *image,const void *pixels,double *filter,
422 const unsigned long width,const unsigned long height,void *convolve_pixels)
431 length=image->columns*image->rows;
432 status=clEnqueueWriteBuffer(convolve_info->command_queue,
433 convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),pixels,0,NULL,
436 status=clEnqueueWriteBuffer(convolve_info->command_queue,
437 convolve_info->filter,CL_TRUE,0,length*sizeof(cl_double),filter,0,NULL,
439 if (status != CL_SUCCESS)
441 global_work_size[0]=image->columns;
442 global_work_size[1]=image->rows;
443 status=clEnqueueNDRangeKernel(convolve_info->command_queue,
444 convolve_info->kernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
445 if (status != CL_SUCCESS)
447 length=image->columns*image->rows;
448 status=clEnqueueReadBuffer(convolve_info->command_queue,
449 convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),
450 convolve_pixels,0,NULL,NULL);
451 if (status != CL_SUCCESS)
453 status=clFinish(convolve_info->command_queue);
454 if (status != CL_SUCCESS)
459 static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
460 const char *source,ExceptionInfo *exception)
463 options[MaxTextExtent];
473 lengths[] = { strlen(source) };
478 convolve_info=(ConvolveInfo *) AcquireAlignedMemory(1,sizeof(*convolve_info));
479 if (convolve_info == (ConvolveInfo *) NULL)
481 (void) ThrowMagickException(exception,GetMagickModule(),
482 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
483 return((ConvolveInfo *) NULL);
485 (void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info));
487 Create OpenCL context.
489 convolve_info->context=clCreateContextFromType((cl_context_properties *)
490 NULL,CL_DEVICE_TYPE_GPU,ConvolveNotify,exception,&status);
491 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
492 convolve_info->context=clCreateContextFromType((cl_context_properties *)
493 NULL,CL_DEVICE_TYPE_CPU,ConvolveNotify,exception,&status);
494 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
495 convolve_info->context=clCreateContextFromType((cl_context_properties *)
496 NULL,CL_DEVICE_TYPE_DEFAULT,ConvolveNotify,exception,&status);
497 if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
499 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
500 "failed to create OpenCL context","`%s' (%d)",image->filename,status);
501 DestroyConvolveInfo(convolve_info);
502 return((ConvolveInfo *) NULL);
505 Detect OpenCL devices.
507 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,0,NULL,
509 if ((status != CL_SUCCESS) || (length == 0))
511 DestroyConvolveInfo(convolve_info);
512 return((ConvolveInfo *) NULL);
514 convolve_info->devices=(cl_device_id *) AcquireMagickMemory(length);
515 if (convolve_info->devices == (cl_device_id *) NULL)
517 (void) ThrowMagickException(exception,GetMagickModule(),
518 ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
519 DestroyConvolveInfo(convolve_info);
520 return((ConvolveInfo *) NULL);
522 status=clGetContextInfo(convolve_info->context,CL_CONTEXT_DEVICES,length,
523 convolve_info->devices,NULL);
524 if (status != CL_SUCCESS)
526 DestroyConvolveInfo(convolve_info);
527 return((ConvolveInfo *) NULL);
530 Create OpenCL command queue.
532 convolve_info->command_queue=clCreateCommandQueue(convolve_info->context,
533 convolve_info->devices[0],0,&status);
534 if ((convolve_info->command_queue == (cl_command_queue) NULL) ||
535 (status != CL_SUCCESS))
537 DestroyConvolveInfo(convolve_info);
538 return((ConvolveInfo *) NULL);
541 Build OpenCL program.
543 convolve_info->program=clCreateProgramWithSource(convolve_info->context,1,
544 &source,lengths,&status);
545 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
547 DestroyConvolveInfo(convolve_info);
548 return((ConvolveInfo *) NULL);
550 (void) FormatMagickString(options,MaxTextExtent,CLOptions,(double)
551 QuantumRange,MagickEpsilon);
552 status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options,
554 if ((convolve_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
559 status=clGetProgramBuildInfo(convolve_info->program,
560 convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,0,NULL,&length);
561 log=(char *) AcquireMagickMemory(length);
562 if (log == (char *) NULL)
564 DestroyConvolveInfo(convolve_info);
565 return((ConvolveInfo *) NULL);
567 status=clGetProgramBuildInfo(convolve_info->program,
568 convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length);
569 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
570 "failed to build OpenCL program","`%s' (%s)",image->filename,log);
571 log=DestroyString(log);
572 DestroyConvolveInfo(convolve_info);
573 return((ConvolveInfo *) NULL);
578 convolve_info->kernel=clCreateKernel(convolve_info->program,name,&status);
579 if ((convolve_info->kernel == (cl_kernel) NULL) || (status != CL_SUCCESS))
581 DestroyConvolveInfo(convolve_info);
582 return((ConvolveInfo *) NULL);
584 return(convolve_info);
589 MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
590 const KernelInfo *kernel,Image *convolve_image,ExceptionInfo *exception)
592 assert(image != (Image *) NULL);
593 assert(image->signature == MagickSignature);
594 if (image->debug != MagickFalse)
595 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
596 assert(kernel != (KernelInfo *) NULL);
597 assert(kernel->signature == MagickSignature);
598 assert(convolve_image != (Image *) NULL);
599 assert(convolve_image->signature == MagickSignature);
600 assert(exception != (ExceptionInfo *) NULL);
601 assert(exception->signature == MagickSignature);
602 if ((image->storage_class != DirectClass) ||
603 (image->colorspace == CMYKColorspace))
604 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
605 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
607 #if !defined(MAGICKCORE_OPENCL_SUPPORT)
626 convolve_info=GetConvolveInfo(image,"Convolve",ConvolveKernel,exception);
627 if (convolve_info == (ConvolveInfo *) NULL)
629 pixels=AcquirePixelCachePixels(image,&length,exception);
630 if (pixels == (const void *) NULL)
632 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
633 "UnableToReadPixelCache","`%s'",image->filename);
634 convolve_info=DestroyConvolveInfo(convolve_info);
637 convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception);
638 if (convolve_pixels == (void *) NULL)
640 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
641 "UnableToReadPixelCache","`%s'",image->filename);
642 convolve_info=DestroyConvolveInfo(convolve_info);
645 status=BindConvolveParameters(convolve_info,image,pixels,kernel->values,
646 kernel->width,kernel->height,convolve_pixels);
647 if (status == MagickFalse)
649 DestroyConvolveBuffers(convolve_info);
650 convolve_info=DestroyConvolveInfo(convolve_info);
653 status=EnqueueConvolveKernel(convolve_info,image,pixels,kernel->values,
654 kernel->width,kernel->height,convolve_pixels);
655 if (status == MagickFalse)
657 DestroyConvolveBuffers(convolve_info);
658 convolve_info=DestroyConvolveInfo(convolve_info);
661 DestroyConvolveBuffers(convolve_info);
662 convolve_info=DestroyConvolveInfo(convolve_info);