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 %
22 % Copyright 1999-2015 ImageMagick Studio LLC, a non-profit organization %
23 % dedicated to making software imaging solutions freely available. %
25 % You may not use this file except in compliance with the License. You may %
26 % obtain a copy of the License at %
28 % http://www.imagemagick.org/script/license.php %
30 % Unless required by applicable law or agreed to in writing, software %
31 % distributed under the License is distributed on an "AS IS" BASIS, %
32 % WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. %
33 % See the License for the specific language governing permissions and %
34 % limitations under the License. %
36 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
42 #include "MagickCore/studio.h"
43 #include "MagickCore/accelerate.h"
44 #include "MagickCore/accelerate-private.h"
45 #include "MagickCore/artifact.h"
46 #include "MagickCore/cache.h"
47 #include "MagickCore/cache-private.h"
48 #include "MagickCore/cache-view.h"
49 #include "MagickCore/color-private.h"
50 #include "MagickCore/delegate-private.h"
51 #include "MagickCore/enhance.h"
52 #include "MagickCore/exception.h"
53 #include "MagickCore/exception-private.h"
54 #include "MagickCore/gem.h"
55 #include "MagickCore/hashmap.h"
56 #include "MagickCore/image.h"
57 #include "MagickCore/image-private.h"
58 #include "MagickCore/list.h"
59 #include "MagickCore/memory_.h"
60 #include "MagickCore/monitor-private.h"
61 #include "MagickCore/accelerate.h"
62 #include "MagickCore/opencl.h"
63 #include "MagickCore/opencl-private.h"
64 #include "MagickCore/option.h"
65 #include "MagickCore/pixel-accessor.h"
66 #include "MagickCore/pixel-private.h"
67 #include "MagickCore/prepress.h"
68 #include "MagickCore/quantize.h"
69 #include "MagickCore/quantum-private.h"
70 #include "MagickCore/random_.h"
71 #include "MagickCore/random-private.h"
72 #include "MagickCore/registry.h"
73 #include "MagickCore/resize.h"
74 #include "MagickCore/resize-private.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"
82 #ifdef MAGICKCORE_CLPERFMARKER
83 #include "CLPerfMarker.h"
86 #define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
87 #define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
89 #if defined(MAGICKCORE_OPENCL_SUPPORT)
91 #define ALIGNED(pointer,type) ((((size_t)(pointer)) & (sizeof(type)-1)) == 0)
93 /* pad the global workgroup size to the next multiple of
94 the local workgroup size */
95 inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
96 const unsigned int orgGlobalSize,const unsigned int localGroupSize)
98 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
101 static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception)
109 clEnv=GetDefaultOpenCLEnv();
111 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,
112 sizeof(MagickBooleanType),&flag,exception);
113 if (flag != MagickFalse)
116 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED,
117 sizeof(MagickBooleanType),&flag,exception);
118 if (flag == MagickFalse)
120 if (InitOpenCLEnv(clEnv,exception) == MagickFalse)
123 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,
124 sizeof(MagickBooleanType),&flag,exception);
125 if (flag != MagickFalse)
132 static MagickBooleanType checkAccelerateCondition(const Image* image,
133 const ChannelType channel)
135 /* check if the image's colorspace is supported */
136 if (image->colorspace != RGBColorspace &&
137 image->colorspace != sRGBColorspace &&
138 image->colorspace != GRAYColorspace)
141 /* check if the channel is supported */
142 if (((channel & RedChannel) == 0) ||
143 ((channel & GreenChannel) == 0) ||
144 ((channel & BlueChannel) == 0))
147 /* check if the virtual pixel method is compatible with the OpenCL implementation */
148 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
149 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
152 /* check if the image has read / write mask */
153 if (image->read_mask != MagickFalse || image->write_mask != MagickFalse)
156 /* check if pixel order is RGBA */
157 if (GetPixelChannelOffset(image,RedPixelChannel) != 0 ||
158 GetPixelChannelOffset(image,GreenPixelChannel) != 1 ||
159 GetPixelChannelOffset(image,BluePixelChannel) != 2 ||
160 GetPixelChannelOffset(image,AlphaPixelChannel) != 3)
163 /* check if all channels are available */
164 if (((GetPixelRedTraits(image) & UpdatePixelTrait) == 0) ||
165 ((GetPixelGreenTraits(image) & UpdatePixelTrait) == 0) ||
166 ((GetPixelBlueTraits(image) & UpdatePixelTrait) == 0) ||
167 ((GetPixelAlphaTraits(image) & UpdatePixelTrait) == 0))
173 static MagickBooleanType checkHistogramCondition(Image *image,
174 const ChannelType channel)
176 /* ensure this is the only pass get in for now. */
177 if ((channel & SyncChannels) == 0)
180 if (image->intensity == Rec601LuminancePixelIntensityMethod ||
181 image->intensity == Rec709LuminancePixelIntensityMethod)
184 if (image->colorspace != sRGBColorspace)
190 static MagickBooleanType splitImage(const Image* image)
202 clEnv=GetDefaultOpenCLEnv();
204 allocSize=GetOpenCLDeviceMaxMemAllocSize(clEnv);
205 tempSize=(unsigned long) (image->columns * image->rows * 4 * 4);
207 split = ((tempSize > allocSize) ? MagickTrue : MagickFalse);
212 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
216 % C o n v o l v e I m a g e w i t h O p e n C L %
220 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
222 % ConvolveImage() applies a custom convolution kernel to the image.
224 % The format of the ConvolveImage method is:
226 % Image *ConvolveImage(const Image *image,const size_t order,
227 % const double *kernel,ExceptionInfo *exception)
228 % Image *ConvolveImageChannel(const Image *image,const ChannelType channel,
229 % const size_t order,const double *kernel,ExceptionInfo *exception)
231 % A description of each parameter follows:
233 % o image: the image.
235 % o channel: the channel type.
237 % o kernel: kernel info.
239 % o exception: return any errors or warnings in this structure.
243 static Image *ComputeConvolveImage(const Image* image,
244 const ChannelType channel,const KernelInfo *kernel,ExceptionInfo *exception)
271 deviceLocalMemorySize;
294 localMemoryRequirement;
311 /* intialize all CL objects to NULL */
314 filteredImageBuffer = NULL;
315 convolutionKernel = NULL;
319 filteredImage = NULL;
320 filteredImage_view = NULL;
321 outputReady = MagickFalse;
323 clEnv = GetDefaultOpenCLEnv();
324 context = GetOpenCLContext(clEnv);
326 image_view=AcquireVirtualCacheView(image,exception);
327 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
328 if (inputPixels == (const void *) NULL)
330 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
334 /* Create and initialize OpenCL buffers. */
336 /* If the host pointer is aligned to the size of CLPixelPacket,
337 then use the host buffer directly from the GPU; otherwise,
338 create a buffer on the GPU and copy the data over */
339 if (ALIGNED(inputPixels,CLPixelPacket))
341 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
345 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
347 /* create a CL buffer from image pixel buffer */
348 length = image->columns * image->rows;
349 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
350 if (clStatus != CL_SUCCESS)
352 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
356 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
357 assert(filteredImage != NULL);
358 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
360 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
363 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
364 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
365 if (filteredPixels == (void *) NULL)
367 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
371 if (ALIGNED(filteredPixels,CLPixelPacket))
373 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
374 hostPtr = filteredPixels;
378 mem_flags = CL_MEM_WRITE_ONLY;
381 /* create a CL buffer from image pixel buffer */
382 length = image->columns * image->rows;
383 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
384 if (clStatus != CL_SUCCESS)
386 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
390 kernelSize = (unsigned int) (kernel->width * kernel->height);
391 convolutionKernel = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus);
392 if (clStatus != CL_SUCCESS)
394 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
398 queue = AcquireOpenCLCommandQueue(clEnv);
400 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize * sizeof(float)
401 , 0, NULL, NULL, &clStatus);
402 if (clStatus != CL_SUCCESS)
404 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
407 for (i = 0; i < kernelSize; i++)
409 kernelBufferPtr[i] = (float) kernel->values[i];
411 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
412 if (clStatus != CL_SUCCESS)
414 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
417 clEnv->library->clFlush(queue);
419 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
421 /* Compute the local memory requirement for a 16x16 workgroup.
422 If it's larger than 16k, reduce the workgroup size to 8x8 */
423 localGroupSize[0] = 16;
424 localGroupSize[1] = 16;
425 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
426 + kernel->width*kernel->height*sizeof(float);
428 if (localMemoryRequirement > deviceLocalMemorySize)
430 localGroupSize[0] = 8;
431 localGroupSize[1] = 8;
432 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
433 + kernel->width*kernel->height*sizeof(float);
435 if (localMemoryRequirement <= deviceLocalMemorySize)
437 /* get the OpenCL kernel */
438 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized");
439 if (clkernel == NULL)
441 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
445 /* set the kernel arguments */
447 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
448 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
449 imageWidth = (unsigned int) image->columns;
450 imageHeight = (unsigned int) image->rows;
451 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
452 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
453 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
454 filterWidth = (unsigned int) kernel->width;
455 filterHeight = (unsigned int) kernel->height;
456 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
457 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
458 matte = (image->alpha_trait != BlendPixelTrait)?1:0;
459 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
460 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
461 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*sizeof(CLPixelPacket),NULL);
462 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL);
463 if (clStatus != CL_SUCCESS)
465 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
469 /* pad the global size to a multiple of the local work size dimension */
470 global_work_size[0] = ((image->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ;
471 global_work_size[1] = ((image->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1];
473 /* launch the kernel */
474 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, NULL);
475 if (clStatus != CL_SUCCESS)
477 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
483 /* get the OpenCL kernel */
484 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve");
485 if (clkernel == NULL)
487 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
491 /* set the kernel arguments */
493 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
494 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
495 imageWidth = (unsigned int) image->columns;
496 imageHeight = (unsigned int) image->rows;
497 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
498 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
499 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
500 filterWidth = (unsigned int) kernel->width;
501 filterHeight = (unsigned int) kernel->height;
502 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
503 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
504 matte = (image->alpha_trait != BlendPixelTrait)?1:0;
505 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
506 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
507 if (clStatus != CL_SUCCESS)
509 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
513 localGroupSize[0] = 8;
514 localGroupSize[1] = 8;
515 global_work_size[0] = (image->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0];
516 global_work_size[1] = (image->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1];
517 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, NULL);
519 if (clStatus != CL_SUCCESS)
521 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
525 clEnv->library->clFlush(queue);
527 if (ALIGNED(filteredPixels,CLPixelPacket))
529 length = image->columns * image->rows;
530 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
534 length = image->columns * image->rows;
535 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
537 if (clStatus != CL_SUCCESS)
539 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
543 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
546 OpenCLLogException(__FUNCTION__,__LINE__,exception);
548 image_view=DestroyCacheView(image_view);
549 if (filteredImage_view != NULL)
550 filteredImage_view=DestroyCacheView(filteredImage_view);
552 if (imageBuffer != NULL)
553 clEnv->library->clReleaseMemObject(imageBuffer);
555 if (filteredImageBuffer != NULL)
556 clEnv->library->clReleaseMemObject(filteredImageBuffer);
558 if (convolutionKernel != NULL)
559 clEnv->library->clReleaseMemObject(convolutionKernel);
561 if (clkernel != NULL)
562 RelinquishOpenCLKernel(clEnv, clkernel);
565 RelinquishOpenCLCommandQueue(clEnv, queue);
567 if (outputReady == MagickFalse)
569 if (filteredImage != NULL)
571 DestroyImage(filteredImage);
572 filteredImage = NULL;
576 return(filteredImage);
579 MagickExport Image *AccelerateConvolveImageChannel(const Image *image,
580 const ChannelType channel,const KernelInfo *kernel,ExceptionInfo *exception)
585 assert(image != NULL);
586 assert(kernel != (KernelInfo *) NULL);
587 assert(exception != (ExceptionInfo *) NULL);
589 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
590 (checkAccelerateCondition(image, channel) == MagickFalse))
593 filteredImage=ComputeConvolveImage(image, channel, kernel, exception);
594 return(filteredImage);
597 static MagickBooleanType ComputeFunctionImage(Image *image,
598 const ChannelType channel,const MagickFunction function,
599 const size_t number_parameters,const double *parameters,
600 ExceptionInfo *exception)
625 *parametersBufferPtr;
645 status = MagickFalse;
651 parametersBuffer = NULL;
653 clEnv = GetDefaultOpenCLEnv();
654 context = GetOpenCLContext(clEnv);
656 image_view=AcquireAuthenticCacheView(image,exception);
657 pixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
658 if (pixels == (void *) NULL)
660 (void) OpenCLThrowMagickException(exception, GetMagickModule(), CacheWarning,
661 "GetPixelCachePixels failed.",
662 "'%s'", image->filename);
667 if (ALIGNED(pixels,CLPixelPacket))
669 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
673 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
675 /* create a CL buffer from image pixel buffer */
676 length = image->columns * image->rows;
677 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)pixels, &clStatus);
678 if (clStatus != CL_SUCCESS)
680 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
684 parametersBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus);
685 if (clStatus != CL_SUCCESS)
687 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
691 queue = AcquireOpenCLCommandQueue(clEnv);
693 parametersBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, parametersBuffer, CL_TRUE, CL_MAP_WRITE, 0, number_parameters * sizeof(float)
694 , 0, NULL, NULL, &clStatus);
695 if (clStatus != CL_SUCCESS)
697 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
700 for (i = 0; i < number_parameters; i++)
702 parametersBufferPtr[i] = (float)parameters[i];
704 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL);
705 if (clStatus != CL_SUCCESS)
707 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
710 clEnv->library->clFlush(queue);
712 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "FunctionImage");
713 if (clkernel == NULL)
715 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
719 /* set the kernel arguments */
721 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
722 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
723 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(MagickFunction),(void *)&function);
724 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&number_parameters);
725 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)¶metersBuffer);
726 if (clStatus != CL_SUCCESS)
728 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
732 globalWorkSize[0] = image->columns;
733 globalWorkSize[1] = image->rows;
734 /* launch the kernel */
735 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, NULL);
736 if (clStatus != CL_SUCCESS)
738 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
741 clEnv->library->clFlush(queue);
744 if (ALIGNED(pixels,CLPixelPacket))
746 length = image->columns * image->rows;
747 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
751 length = image->columns * image->rows;
752 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), pixels, 0, NULL, NULL);
754 if (clStatus != CL_SUCCESS)
756 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
759 status=SyncCacheViewAuthenticPixels(image_view,exception);
762 OpenCLLogException(__FUNCTION__,__LINE__,exception);
764 image_view=DestroyCacheView(image_view);
766 if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel);
767 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
768 if (imageBuffer != NULL) clEnv->library->clReleaseMemObject(imageBuffer);
769 if (parametersBuffer != NULL) clEnv->library->clReleaseMemObject(parametersBuffer);
774 MagickExport MagickBooleanType AccelerateFunctionImage(Image *image,
775 const ChannelType channel,const MagickFunction function,
776 const size_t number_parameters,const double *parameters,
777 ExceptionInfo *exception)
782 assert(image != NULL);
783 assert(exception != (ExceptionInfo *) NULL);
785 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
786 (checkAccelerateCondition(image, channel) == MagickFalse))
789 status=ComputeFunctionImage(image, channel, function, number_parameters, parameters, exception);
794 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
798 % B l u r I m a g e w i t h O p e n C L %
802 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
804 % BlurImage() blurs an image. We convolve the image with a Gaussian operator
805 % of the given radius and standard deviation (sigma). For reasonable results,
806 % the radius should be larger than sigma. Use a radius of 0 and BlurImage()
807 % selects a suitable radius for you.
809 % The format of the BlurImage method is:
811 % Image *BlurImage(const Image *image,const double radius,
812 % const double sigma,ExceptionInfo *exception)
813 % Image *BlurImageChannel(const Image *image,const ChannelType channel,
814 % const double radius,const double sigma,ExceptionInfo *exception)
816 % A description of each parameter follows:
818 % o image: the image.
820 % o channel: the channel type.
822 % o radius: the radius of the Gaussian, in pixels, not counting the center
825 % o sigma: the standard deviation of the Gaussian, in pixels.
827 % o exception: return any errors or warnings in this structure.
831 static Image *ComputeBlurImage(const Image* image,const ChannelType channel,
832 const double radius,const double sigma,ExceptionInfo *exception)
839 geometry[MaxTextExtent];
895 filteredImage = NULL;
896 filteredImage_view = NULL;
898 tempImageBuffer = NULL;
899 filteredImageBuffer = NULL;
900 imageKernelBuffer = NULL;
901 blurRowKernel = NULL;
902 blurColumnKernel = NULL;
906 outputReady = MagickFalse;
908 clEnv = GetDefaultOpenCLEnv();
909 context = GetOpenCLContext(clEnv);
910 queue = AcquireOpenCLCommandQueue(clEnv);
912 /* Create and initialize OpenCL buffers. */
914 image_view=AcquireVirtualCacheView(image,exception);
915 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
916 if (inputPixels == (const void *) NULL)
918 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
921 /* If the host pointer is aligned to the size of CLPixelPacket,
922 then use the host buffer directly from the GPU; otherwise,
923 create a buffer on the GPU and copy the data over */
924 if (ALIGNED(inputPixels,CLPixelPacket))
926 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
930 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
932 /* create a CL buffer from image pixel buffer */
933 length = image->columns * image->rows;
934 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
935 if (clStatus != CL_SUCCESS)
937 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
944 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
945 assert(filteredImage != NULL);
946 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
948 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
951 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
952 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
953 if (filteredPixels == (void *) NULL)
955 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
959 if (ALIGNED(filteredPixels,CLPixelPacket))
961 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
962 hostPtr = filteredPixels;
966 mem_flags = CL_MEM_WRITE_ONLY;
969 /* create a CL buffer from image pixel buffer */
970 length = image->columns * image->rows;
971 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
972 if (clStatus != CL_SUCCESS)
974 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
979 /* create processing kernel */
981 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
982 kernel=AcquireKernelInfo(geometry,exception);
983 if (kernel == (KernelInfo *) NULL)
985 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
989 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus);
990 if (clStatus != CL_SUCCESS)
992 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
995 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
996 if (clStatus != CL_SUCCESS)
998 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
1002 for (i = 0; i < kernel->width; i++)
1004 kernelBufferPtr[i] = (float) kernel->values[i];
1007 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
1008 if (clStatus != CL_SUCCESS)
1010 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1017 /* create temp buffer */
1019 length = image->columns * image->rows;
1020 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
1021 if (clStatus != CL_SUCCESS)
1023 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1028 /* get the OpenCL kernels */
1030 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
1031 if (blurRowKernel == NULL)
1033 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1037 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn");
1038 if (blurColumnKernel == NULL)
1040 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1046 /* need logic to decide this value */
1047 int chunkSize = 256;
1050 imageColumns = (unsigned int) image->columns;
1051 imageRows = (unsigned int) image->rows;
1053 /* set the kernel arguments */
1055 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1056 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1057 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
1058 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1059 kernelWidth = (unsigned int) kernel->width;
1060 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1061 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1062 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1063 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
1064 if (clStatus != CL_SUCCESS)
1066 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1071 /* launch the kernel */
1076 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
1077 gsize[1] = image->rows;
1078 wsize[0] = chunkSize;
1081 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1082 if (clStatus != CL_SUCCESS)
1084 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1087 clEnv->library->clFlush(queue);
1092 /* need logic to decide this value */
1093 int chunkSize = 256;
1096 imageColumns = (unsigned int) image->columns;
1097 imageRows = (unsigned int) image->rows;
1099 /* set the kernel arguments */
1101 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1102 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1103 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
1104 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1105 kernelWidth = (unsigned int) kernel->width;
1106 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1107 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1108 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1109 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *) NULL);
1110 if (clStatus != CL_SUCCESS)
1112 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1117 /* launch the kernel */
1122 gsize[0] = image->columns;
1123 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
1125 wsize[1] = chunkSize;
1127 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1128 if (clStatus != CL_SUCCESS)
1130 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1133 clEnv->library->clFlush(queue);
1140 if (ALIGNED(filteredPixels,CLPixelPacket))
1142 length = image->columns * image->rows;
1143 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1147 length = image->columns * image->rows;
1148 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1150 if (clStatus != CL_SUCCESS)
1152 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1156 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
1159 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1161 image_view=DestroyCacheView(image_view);
1162 if (filteredImage_view != NULL)
1163 filteredImage_view=DestroyCacheView(filteredImage_view);
1165 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
1166 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
1167 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
1168 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
1169 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
1170 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
1171 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1172 if (kernel!=NULL) DestroyKernelInfo(kernel);
1173 if (outputReady == MagickFalse && filteredImage != NULL)
1174 filteredImage=DestroyImage(filteredImage);
1175 return(filteredImage);
1178 static Image* ComputeBlurImageSection(const Image* image,
1179 const ChannelType channel,const double radius,const double sigma,
1180 ExceptionInfo *exception)
1183 *filteredImage_view,
1187 geometry[MaxTextExtent];
1202 filteredImageBuffer,
1243 filteredImage = NULL;
1244 filteredImage_view = NULL;
1246 tempImageBuffer = NULL;
1247 filteredImageBuffer = NULL;
1248 imageKernelBuffer = NULL;
1249 blurRowKernel = NULL;
1250 blurColumnKernel = NULL;
1254 outputReady = MagickFalse;
1256 clEnv = GetDefaultOpenCLEnv();
1257 context = GetOpenCLContext(clEnv);
1258 queue = AcquireOpenCLCommandQueue(clEnv);
1260 /* Create and initialize OpenCL buffers. */
1262 image_view=AcquireVirtualCacheView(image,exception);
1263 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
1264 if (inputPixels == (const void *) NULL)
1266 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
1269 /* If the host pointer is aligned to the size of CLPixelPacket,
1270 then use the host buffer directly from the GPU; otherwise,
1271 create a buffer on the GPU and copy the data over */
1272 if (ALIGNED(inputPixels,CLPixelPacket))
1274 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1278 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1280 /* create a CL buffer from image pixel buffer */
1281 length = image->columns * image->rows;
1282 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1283 if (clStatus != CL_SUCCESS)
1285 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1292 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
1293 assert(filteredImage != NULL);
1294 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1296 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1299 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
1300 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
1301 if (filteredPixels == (void *) NULL)
1303 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1307 if (ALIGNED(filteredPixels,CLPixelPacket))
1309 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1310 hostPtr = filteredPixels;
1314 mem_flags = CL_MEM_WRITE_ONLY;
1317 /* create a CL buffer from image pixel buffer */
1318 length = image->columns * image->rows;
1319 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1320 if (clStatus != CL_SUCCESS)
1322 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1327 /* create processing kernel */
1329 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
1330 kernel=AcquireKernelInfo(geometry,exception);
1331 if (kernel == (KernelInfo *) NULL)
1333 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
1337 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus);
1338 if (clStatus != CL_SUCCESS)
1340 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1343 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
1344 if (clStatus != CL_SUCCESS)
1346 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
1350 for (i = 0; i < kernel->width; i++)
1352 kernelBufferPtr[i] = (float) kernel->values[i];
1355 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
1356 if (clStatus != CL_SUCCESS)
1358 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1364 unsigned int offsetRows;
1367 /* create temp buffer */
1369 length = image->columns * (image->rows / 2 + 1 + (kernel->width-1) / 2);
1370 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
1371 if (clStatus != CL_SUCCESS)
1373 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1378 /* get the OpenCL kernels */
1380 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection");
1381 if (blurRowKernel == NULL)
1383 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1387 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumnSection");
1388 if (blurColumnKernel == NULL)
1390 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1395 for (sec = 0; sec < 2; sec++)
1398 /* need logic to decide this value */
1399 int chunkSize = 256;
1402 imageColumns = (unsigned int) image->columns;
1404 imageRows = (unsigned int) (image->rows / 2 + (kernel->width-1) / 2);
1406 imageRows = (unsigned int) ((image->rows - image->rows / 2) + (kernel->width-1) / 2);
1408 offsetRows = (unsigned int) (sec * image->rows / 2);
1410 kernelWidth = (unsigned int) kernel->width;
1412 /* set the kernel arguments */
1414 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1415 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1416 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
1417 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1418 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1419 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1420 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1421 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
1422 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
1423 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec);
1424 if (clStatus != CL_SUCCESS)
1426 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1431 /* launch the kernel */
1436 gsize[0] = chunkSize*((imageColumns+chunkSize-1)/chunkSize);
1437 gsize[1] = imageRows;
1438 wsize[0] = chunkSize;
1441 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1442 if (clStatus != CL_SUCCESS)
1444 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1447 clEnv->library->clFlush(queue);
1452 /* need logic to decide this value */
1453 int chunkSize = 256;
1456 imageColumns = (unsigned int) image->columns;
1458 imageRows = (unsigned int) (image->rows / 2);
1460 imageRows = (unsigned int) ((image->rows - image->rows / 2));
1462 offsetRows = (unsigned int) (sec * image->rows / 2);
1464 kernelWidth = (unsigned int) kernel->width;
1466 /* set the kernel arguments */
1468 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1469 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1470 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
1471 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1472 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1473 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1474 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1475 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *) NULL);
1476 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
1477 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&sec);
1478 if (clStatus != CL_SUCCESS)
1480 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1485 /* launch the kernel */
1490 gsize[0] = imageColumns;
1491 gsize[1] = chunkSize*((imageRows+chunkSize-1)/chunkSize);
1493 wsize[1] = chunkSize;
1495 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1496 if (clStatus != CL_SUCCESS)
1498 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1501 clEnv->library->clFlush(queue);
1509 if (ALIGNED(filteredPixels,CLPixelPacket))
1511 length = image->columns * image->rows;
1512 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1516 length = image->columns * image->rows;
1517 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1519 if (clStatus != CL_SUCCESS)
1521 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1525 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
1528 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1530 image_view=DestroyCacheView(image_view);
1531 if (filteredImage_view != NULL)
1532 filteredImage_view=DestroyCacheView(filteredImage_view);
1534 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
1535 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
1536 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
1537 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
1538 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
1539 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
1540 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1541 if (kernel!=NULL) DestroyKernelInfo(kernel);
1542 if (outputReady == MagickFalse)
1544 if (filteredImage != NULL)
1546 DestroyImage(filteredImage);
1547 filteredImage = NULL;
1550 return filteredImage;
1553 static Image *ComputeUnsharpMaskImageSingle(const Image *image,
1554 const ChannelType channel,const double radius,const double sigma,
1555 const double gain,const double threshold,int blurOnly, ExceptionInfo *exception);
1557 static Image* ComputeBlurImageSingle(const Image* image,
1558 const ChannelType channel,const double radius,const double sigma,
1559 ExceptionInfo *exception)
1561 return ComputeUnsharpMaskImageSingle(image, channel, radius, sigma, 0.0, 0.0, 1, exception);
1564 MagickExport Image* AccelerateBlurImage(const Image *image,
1565 const ChannelType channel,const double radius,const double sigma,
1566 ExceptionInfo *exception)
1571 assert(image != NULL);
1572 assert(exception != (ExceptionInfo *) NULL);
1574 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1575 (checkAccelerateCondition(image, channel) == MagickFalse))
1579 filteredImage=ComputeBlurImageSingle(image, channel, radius, sigma, exception);
1580 else if (splitImage(image) && (image->rows / 2 > radius))
1581 filteredImage=ComputeBlurImageSection(image, channel, radius, sigma, exception);
1583 filteredImage=ComputeBlurImage(image, channel, radius, sigma, exception);
1585 return(filteredImage);
1589 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1593 % R o t a t i o n a l B l u r I m a g e w i t h O p e n C L %
1597 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1599 % RotationalBlurImage() applies a rotational blur to the image.
1601 % Andrew Protano contributed this effect.
1603 % The format of the RotationalBlurImage method is:
1605 % Image *RotationalBlurImage(const Image *image,const double angle,
1606 % ExceptionInfo *exception)
1607 % Image *RotationalBlurImageChannel(const Image *image,const ChannelType channel,
1608 % const double angle,ExceptionInfo *exception)
1610 % A description of each parameter follows:
1612 % o image: the image.
1614 % o channel: the channel type.
1616 % o angle: the angle of the rotational blur.
1618 % o exception: return any errors or warnings in this structure.
1622 static Image* ComputeRotationalBlurImage(const Image *image,
1623 const ChannelType channel,const double angle,ExceptionInfo *exception)
1627 *filteredImage_view;
1646 filteredImageBuffer,
1654 rotationalBlurKernel;
1682 global_work_size[2];
1693 outputReady = MagickFalse;
1695 filteredImage = NULL;
1696 filteredImage_view = NULL;
1698 filteredImageBuffer = NULL;
1699 sinThetaBuffer = NULL;
1700 cosThetaBuffer = NULL;
1702 rotationalBlurKernel = NULL;
1705 clEnv = GetDefaultOpenCLEnv();
1706 context = GetOpenCLContext(clEnv);
1709 /* Create and initialize OpenCL buffers. */
1711 image_view=AcquireVirtualCacheView(image,exception);
1712 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
1713 if (inputPixels == (const void *) NULL)
1715 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
1719 /* If the host pointer is aligned to the size of CLPixelPacket,
1720 then use the host buffer directly from the GPU; otherwise,
1721 create a buffer on the GPU and copy the data over */
1722 if (ALIGNED(inputPixels,CLPixelPacket))
1724 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1728 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1730 /* create a CL buffer from image pixel buffer */
1731 length = image->columns * image->rows;
1732 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1733 if (clStatus != CL_SUCCESS)
1735 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1740 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
1741 assert(filteredImage != NULL);
1742 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1744 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1747 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
1748 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
1749 if (filteredPixels == (void *) NULL)
1751 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1755 if (ALIGNED(filteredPixels,CLPixelPacket))
1757 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1758 hostPtr = filteredPixels;
1762 mem_flags = CL_MEM_WRITE_ONLY;
1765 /* create a CL buffer from image pixel buffer */
1766 length = image->columns * image->rows;
1767 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1768 if (clStatus != CL_SUCCESS)
1770 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1774 blurCenter.s[0] = (float) (image->columns-1)/2.0;
1775 blurCenter.s[1] = (float) (image->rows-1)/2.0;
1776 blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]);
1777 cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL);
1779 /* create a buffer for sin_theta and cos_theta */
1780 sinThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
1781 if (clStatus != CL_SUCCESS)
1783 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1786 cosThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
1787 if (clStatus != CL_SUCCESS)
1789 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1794 queue = AcquireOpenCLCommandQueue(clEnv);
1795 sinThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
1796 if (clStatus != CL_SUCCESS)
1798 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
1802 cosThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
1803 if (clStatus != CL_SUCCESS)
1805 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
1809 theta=DegreesToRadians(angle)/(MagickRealType) (cossin_theta_size-1);
1810 offset=theta*(MagickRealType) (cossin_theta_size-1)/2.0;
1811 for (i=0; i < (ssize_t) cossin_theta_size; i++)
1813 cosThetaPtr[i]=(float)cos((double) (theta*i-offset));
1814 sinThetaPtr[i]=(float)sin((double) (theta*i-offset));
1817 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL);
1818 clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL);
1819 if (clStatus != CL_SUCCESS)
1821 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1825 /* get the OpenCL kernel */
1826 rotationalBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RotationalBlur");
1827 if (rotationalBlurKernel == NULL)
1829 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1834 /* set the kernel arguments */
1836 clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1837 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1839 GetPixelInfo(image,&bias);
1840 biasPixel.s[0] = bias.red;
1841 biasPixel.s[1] = bias.green;
1842 biasPixel.s[2] = bias.blue;
1843 biasPixel.s[3] = bias.alpha;
1844 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel);
1845 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &channel);
1847 matte = (image->alpha_trait != BlendPixelTrait)?1:0;
1848 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &matte);
1850 clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter);
1852 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer);
1853 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer);
1854 clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size);
1855 if (clStatus != CL_SUCCESS)
1857 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1862 global_work_size[0] = image->columns;
1863 global_work_size[1] = image->rows;
1864 /* launch the kernel */
1865 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, rotationalBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
1866 if (clStatus != CL_SUCCESS)
1868 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1871 clEnv->library->clFlush(queue);
1873 if (ALIGNED(filteredPixels,CLPixelPacket))
1875 length = image->columns * image->rows;
1876 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1880 length = image->columns * image->rows;
1881 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1883 if (clStatus != CL_SUCCESS)
1885 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1888 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
1891 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1893 image_view=DestroyCacheView(image_view);
1894 if (filteredImage_view != NULL)
1895 filteredImage_view=DestroyCacheView(filteredImage_view);
1897 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
1898 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
1899 if (sinThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(sinThetaBuffer);
1900 if (cosThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(cosThetaBuffer);
1901 if (rotationalBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, rotationalBlurKernel);
1902 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1903 if (outputReady == MagickFalse)
1905 if (filteredImage != NULL)
1907 DestroyImage(filteredImage);
1908 filteredImage = NULL;
1911 return filteredImage;
1914 MagickExport Image* AccelerateRotationalBlurImage(const Image *image,
1915 const ChannelType channel,const double angle,ExceptionInfo *exception)
1920 assert(image != NULL);
1921 assert(exception != (ExceptionInfo *) NULL);
1923 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1924 (checkAccelerateCondition(image, channel) == MagickFalse))
1927 filteredImage=ComputeRotationalBlurImage(image, channel, angle, exception);
1928 return filteredImage;
1932 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1936 % U n s h a r p M a s k I m a g e w i t h O p e n C L %
1940 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1942 % UnsharpMaskImage() sharpens one or more image channels. We convolve the
1943 % image with a Gaussian operator of the given radius and standard deviation
1944 % (sigma). For reasonable results, radius should be larger than sigma. Use a
1945 % radius of 0 and UnsharpMaskImage() selects a suitable radius for you.
1947 % The format of the UnsharpMaskImage method is:
1949 % Image *UnsharpMaskImage(const Image *image,const double radius,
1950 % const double sigma,const double amount,const double threshold,
1951 % ExceptionInfo *exception)
1952 % Image *UnsharpMaskImageChannel(const Image *image,
1953 % const ChannelType channel,const double radius,const double sigma,
1954 % const double gain,const double threshold,ExceptionInfo *exception)
1956 % A description of each parameter follows:
1958 % o image: the image.
1960 % o channel: the channel type.
1962 % o radius: the radius of the Gaussian, in pixels, not counting the center
1965 % o sigma: the standard deviation of the Gaussian, in pixels.
1967 % o gain: the percentage of the difference between the original and the
1968 % blur image that is added back into the original.
1970 % o threshold: the threshold in pixels needed to apply the diffence gain.
1972 % o exception: return any errors or warnings in this structure.
1976 static Image *ComputeUnsharpMaskImage(const Image *image,
1977 const ChannelType channel,const double radius,const double sigma,
1978 const double gain,const double threshold,ExceptionInfo *exception)
1981 *filteredImage_view,
1985 geometry[MaxTextExtent];
1998 unsharpMaskBlurColumnKernel;
2001 filteredImageBuffer,
2046 filteredImage = NULL;
2047 filteredImage_view = NULL;
2051 filteredImageBuffer = NULL;
2052 tempImageBuffer = NULL;
2053 imageKernelBuffer = NULL;
2054 blurRowKernel = NULL;
2055 unsharpMaskBlurColumnKernel = NULL;
2057 outputReady = MagickFalse;
2059 clEnv = GetDefaultOpenCLEnv();
2060 context = GetOpenCLContext(clEnv);
2061 queue = AcquireOpenCLCommandQueue(clEnv);
2063 /* Create and initialize OpenCL buffers. */
2065 image_view=AcquireVirtualCacheView(image,exception);
2066 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
2067 if (inputPixels == (const void *) NULL)
2069 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
2073 /* If the host pointer is aligned to the size of CLPixelPacket,
2074 then use the host buffer directly from the GPU; otherwise,
2075 create a buffer on the GPU and copy the data over */
2076 if (ALIGNED(inputPixels,CLPixelPacket))
2078 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2082 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2084 /* create a CL buffer from image pixel buffer */
2085 length = image->columns * image->rows;
2086 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2087 if (clStatus != CL_SUCCESS)
2089 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2096 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
2097 assert(filteredImage != NULL);
2098 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
2100 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2103 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
2104 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
2105 if (filteredPixels == (void *) NULL)
2107 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2111 if (ALIGNED(filteredPixels,CLPixelPacket))
2113 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2114 hostPtr = filteredPixels;
2118 mem_flags = CL_MEM_WRITE_ONLY;
2122 /* create a CL buffer from image pixel buffer */
2123 length = image->columns * image->rows;
2124 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2125 if (clStatus != CL_SUCCESS)
2127 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2132 /* create the blur kernel */
2134 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
2135 kernel=AcquireKernelInfo(geometry,exception);
2136 if (kernel == (KernelInfo *) NULL)
2138 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
2142 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
2143 if (clStatus != CL_SUCCESS)
2145 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2150 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
2151 if (clStatus != CL_SUCCESS)
2153 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
2156 for (i = 0; i < kernel->width; i++)
2158 kernelBufferPtr[i] = (float) kernel->values[i];
2160 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
2161 if (clStatus != CL_SUCCESS)
2163 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
2169 /* create temp buffer */
2171 length = image->columns * image->rows;
2172 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
2173 if (clStatus != CL_SUCCESS)
2175 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2180 /* get the opencl kernel */
2182 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
2183 if (blurRowKernel == NULL)
2185 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2189 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn");
2190 if (unsharpMaskBlurColumnKernel == NULL)
2192 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2200 imageColumns = (unsigned int) image->columns;
2201 imageRows = (unsigned int) image->rows;
2203 kernelWidth = (unsigned int) kernel->width;
2205 /* set the kernel arguments */
2207 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2208 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2209 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
2210 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2211 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2212 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2213 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2214 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
2215 if (clStatus != CL_SUCCESS)
2217 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2222 /* launch the kernel */
2227 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
2228 gsize[1] = image->rows;
2229 wsize[0] = chunkSize;
2232 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2233 if (clStatus != CL_SUCCESS)
2235 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2238 clEnv->library->clFlush(queue);
2244 imageColumns = (unsigned int) image->columns;
2245 imageRows = (unsigned int) image->rows;
2246 kernelWidth = (unsigned int) kernel->width;
2247 fGain = (float) gain;
2248 fThreshold = (float) threshold;
2251 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2252 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2253 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
2254 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2255 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2256 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
2257 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
2258 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
2259 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2260 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2261 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
2262 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
2264 if (clStatus != CL_SUCCESS)
2266 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2271 /* launch the kernel */
2276 gsize[0] = image->columns;
2277 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
2279 wsize[1] = chunkSize;
2281 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2282 if (clStatus != CL_SUCCESS)
2284 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2287 clEnv->library->clFlush(queue);
2293 if (ALIGNED(filteredPixels,CLPixelPacket))
2295 length = image->columns * image->rows;
2296 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2300 length = image->columns * image->rows;
2301 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2303 if (clStatus != CL_SUCCESS)
2305 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2309 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
2312 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2314 image_view=DestroyCacheView(image_view);
2315 if (filteredImage_view != NULL)
2316 filteredImage_view=DestroyCacheView(filteredImage_view);
2318 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
2319 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
2320 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
2321 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
2322 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
2323 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
2324 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
2325 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2326 if (outputReady == MagickFalse)
2328 if (filteredImage != NULL)
2330 DestroyImage(filteredImage);
2331 filteredImage = NULL;
2334 return(filteredImage);
2337 static Image *ComputeUnsharpMaskImageSection(const Image *image,
2338 const ChannelType channel,const double radius,const double sigma,
2339 const double gain,const double threshold,ExceptionInfo *exception)
2342 *filteredImage_view,
2346 geometry[MaxTextExtent];
2359 unsharpMaskBlurColumnKernel;
2362 filteredImageBuffer,
2407 filteredImage = NULL;
2408 filteredImage_view = NULL;
2412 filteredImageBuffer = NULL;
2413 tempImageBuffer = NULL;
2414 imageKernelBuffer = NULL;
2415 blurRowKernel = NULL;
2416 unsharpMaskBlurColumnKernel = NULL;
2418 outputReady = MagickFalse;
2420 clEnv = GetDefaultOpenCLEnv();
2421 context = GetOpenCLContext(clEnv);
2422 queue = AcquireOpenCLCommandQueue(clEnv);
2424 /* Create and initialize OpenCL buffers. */
2426 image_view=AcquireVirtualCacheView(image,exception);
2427 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
2428 if (inputPixels == (const void *) NULL)
2430 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
2434 /* If the host pointer is aligned to the size of CLPixelPacket,
2435 then use the host buffer directly from the GPU; otherwise,
2436 create a buffer on the GPU and copy the data over */
2437 if (ALIGNED(inputPixels,CLPixelPacket))
2439 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2443 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2445 /* create a CL buffer from image pixel buffer */
2446 length = image->columns * image->rows;
2447 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2448 if (clStatus != CL_SUCCESS)
2450 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2457 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
2458 assert(filteredImage != NULL);
2459 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
2461 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2464 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
2465 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
2466 if (filteredPixels == (void *) NULL)
2468 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2472 if (ALIGNED(filteredPixels,CLPixelPacket))
2474 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2475 hostPtr = filteredPixels;
2479 mem_flags = CL_MEM_WRITE_ONLY;
2483 /* create a CL buffer from image pixel buffer */
2484 length = image->columns * image->rows;
2485 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2486 if (clStatus != CL_SUCCESS)
2488 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2493 /* create the blur kernel */
2495 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
2496 kernel=AcquireKernelInfo(geometry,exception);
2497 if (kernel == (KernelInfo *) NULL)
2499 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
2503 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
2504 if (clStatus != CL_SUCCESS)
2506 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2511 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
2512 if (clStatus != CL_SUCCESS)
2514 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
2517 for (i = 0; i < kernel->width; i++)
2519 kernelBufferPtr[i] = (float) kernel->values[i];
2521 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
2522 if (clStatus != CL_SUCCESS)
2524 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
2530 unsigned int offsetRows;
2533 /* create temp buffer */
2535 length = image->columns * (image->rows / 2 + 1 + (kernel->width-1) / 2);
2536 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
2537 if (clStatus != CL_SUCCESS)
2539 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2544 /* get the opencl kernel */
2546 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection");
2547 if (blurRowKernel == NULL)
2549 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2553 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumnSection");
2554 if (unsharpMaskBlurColumnKernel == NULL)
2556 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2561 for (sec = 0; sec < 2; sec++)
2566 imageColumns = (unsigned int) image->columns;
2568 imageRows = (unsigned int) (image->rows / 2 + (kernel->width-1) / 2);
2570 imageRows = (unsigned int) ((image->rows - image->rows / 2) + (kernel->width-1) / 2);
2572 offsetRows = (unsigned int) (sec * image->rows / 2);
2574 kernelWidth = (unsigned int) kernel->width;
2576 /* set the kernel arguments */
2578 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2579 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2580 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
2581 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2582 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2583 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2584 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2585 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
2586 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
2587 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec);
2588 if (clStatus != CL_SUCCESS)
2590 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2594 /* launch the kernel */
2599 gsize[0] = chunkSize*((imageColumns+chunkSize-1)/chunkSize);
2600 gsize[1] = imageRows;
2601 wsize[0] = chunkSize;
2604 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2605 if (clStatus != CL_SUCCESS)
2607 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2610 clEnv->library->clFlush(queue);
2617 imageColumns = (unsigned int) image->columns;
2619 imageRows = (unsigned int) (image->rows / 2);
2621 imageRows = (unsigned int) (image->rows - image->rows / 2);
2623 offsetRows = (unsigned int) (sec * image->rows / 2);
2625 kernelWidth = (unsigned int) kernel->width;
2627 fGain = (float) gain;
2628 fThreshold = (float) threshold;
2631 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2632 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2633 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
2634 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2635 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2636 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
2637 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
2638 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
2639 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2640 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2641 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
2642 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
2643 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
2644 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&sec);
2646 if (clStatus != CL_SUCCESS)
2648 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2653 /* launch the kernel */
2658 gsize[0] = imageColumns;
2659 gsize[1] = chunkSize*((imageRows+chunkSize-1)/chunkSize);
2661 wsize[1] = chunkSize;
2663 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2664 if (clStatus != CL_SUCCESS)
2666 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2669 clEnv->library->clFlush(queue);
2675 if (ALIGNED(filteredPixels,CLPixelPacket))
2677 length = image->columns * image->rows;
2678 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2682 length = image->columns * image->rows;
2683 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2685 if (clStatus != CL_SUCCESS)
2687 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2691 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
2694 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2696 image_view=DestroyCacheView(image_view);
2697 if (filteredImage_view != NULL)
2698 filteredImage_view=DestroyCacheView(filteredImage_view);
2700 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
2701 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
2702 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
2703 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
2704 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
2705 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
2706 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
2707 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2708 if (outputReady == MagickFalse)
2710 if (filteredImage != NULL)
2712 DestroyImage(filteredImage);
2713 filteredImage = NULL;
2716 return filteredImage;
2719 static Image *ComputeUnsharpMaskImageSingle(const Image *image,
2720 const ChannelType channel,const double radius,const double sigma,
2721 const double gain,const double threshold,int blurOnly, ExceptionInfo *exception)
2724 *filteredImage_view,
2728 geometry[MaxTextExtent];
2744 filteredImageBuffer,
2785 filteredImage = NULL;
2786 filteredImage_view = NULL;
2790 filteredImageBuffer = NULL;
2791 imageKernelBuffer = NULL;
2792 unsharpMaskKernel = NULL;
2794 outputReady = MagickFalse;
2796 clEnv = GetDefaultOpenCLEnv();
2797 context = GetOpenCLContext(clEnv);
2798 queue = AcquireOpenCLCommandQueue(clEnv);
2800 /* Create and initialize OpenCL buffers. */
2802 image_view=AcquireVirtualCacheView(image,exception);
2803 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
2804 if (inputPixels == (const void *) NULL)
2806 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
2810 /* If the host pointer is aligned to the size of CLPixelPacket,
2811 then use the host buffer directly from the GPU; otherwise,
2812 create a buffer on the GPU and copy the data over */
2813 if (ALIGNED(inputPixels,CLPixelPacket))
2815 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2819 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2821 /* create a CL buffer from image pixel buffer */
2822 length = image->columns * image->rows;
2823 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2824 if (clStatus != CL_SUCCESS)
2826 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2833 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
2834 assert(filteredImage != NULL);
2835 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
2837 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2840 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
2841 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
2842 if (filteredPixels == (void *) NULL)
2844 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2848 if (ALIGNED(filteredPixels,CLPixelPacket))
2850 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2851 hostPtr = filteredPixels;
2855 mem_flags = CL_MEM_WRITE_ONLY;
2859 /* create a CL buffer from image pixel buffer */
2860 length = image->columns * image->rows;
2861 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2862 if (clStatus != CL_SUCCESS)
2864 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2869 /* create the blur kernel */
2871 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
2872 kernel=AcquireKernelInfo(geometry,exception);
2873 if (kernel == (KernelInfo *) NULL)
2875 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
2879 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
2880 if (clStatus != CL_SUCCESS)
2882 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2887 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
2888 if (clStatus != CL_SUCCESS)
2890 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
2893 for (i = 0; i < kernel->width; i++)
2895 kernelBufferPtr[i] = (float) kernel->values[i];
2897 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
2898 if (clStatus != CL_SUCCESS)
2900 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
2906 /* get the opencl kernel */
2908 unsharpMaskKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMask");
2909 if (unsharpMaskKernel == NULL)
2911 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2917 imageColumns = (unsigned int) image->columns;
2918 imageRows = (unsigned int) image->rows;
2919 kernelWidth = (unsigned int) kernel->width;
2920 fGain = (float) gain;
2921 fThreshold = (float) threshold;
2922 justBlur = blurOnly;
2924 /* set the kernel arguments */
2926 clStatus=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2927 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
2928 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2929 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2930 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2931 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2932 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernel->width)),(void *) NULL);
2933 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain);
2934 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold);
2935 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&justBlur);
2936 if (clStatus != CL_SUCCESS)
2938 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2943 /* launch the kernel */
2948 gsize[0] = ((image->columns + 7) / 8) * 8;
2949 gsize[1] = ((image->rows + 31) / 32) * 32;
2953 clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2954 if (clStatus != CL_SUCCESS)
2956 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2959 clEnv->library->clFlush(queue);
2964 if (ALIGNED(filteredPixels,CLPixelPacket))
2966 length = image->columns * image->rows;
2967 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2971 length = image->columns * image->rows;
2972 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2974 if (clStatus != CL_SUCCESS)
2976 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2980 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
2983 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2985 image_view=DestroyCacheView(image_view);
2986 if (filteredImage_view != NULL)
2987 filteredImage_view=DestroyCacheView(filteredImage_view);
2989 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
2990 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
2991 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
2992 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
2993 if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskKernel);
2994 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2995 if (outputReady == MagickFalse)
2997 if (filteredImage != NULL)
2999 DestroyImage(filteredImage);
3000 filteredImage = NULL;
3003 return(filteredImage);
3007 MagickExport Image *AccelerateUnsharpMaskImage(const Image *image,
3008 const ChannelType channel,const double radius,const double sigma,
3009 const double gain,const double threshold,ExceptionInfo *exception)
3014 assert(image != NULL);
3015 assert(exception != (ExceptionInfo *) NULL);
3017 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3018 (checkAccelerateCondition(image, channel) == MagickFalse))
3022 filteredImage = ComputeUnsharpMaskImageSingle(image,channel,radius,sigma,gain,threshold, 0, exception);
3023 else if (splitImage(image) && (image->rows / 2 > radius))
3024 filteredImage = ComputeUnsharpMaskImageSection(image,channel,radius,sigma,gain,threshold,exception);
3026 filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
3027 return(filteredImage);
3031 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3035 % A c c e l e r a t e R e s i z e I m a g e %
3039 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3041 % AccelerateResizeImage() is an OpenCL implementation of ResizeImage()
3043 % AccelerateResizeImage() scales an image to the desired dimensions, using the given
3044 % filter (see AcquireFilterInfo()).
3046 % If an undefined filter is given the filter defaults to Mitchell for a
3047 % colormapped image, a image with a matte channel, or if the image is
3048 % enlarged. Otherwise the filter defaults to a Lanczos.
3050 % AccelerateResizeImage() was inspired by Paul Heckbert's "zoom" program.
3052 % The format of the AccelerateResizeImage method is:
3054 % Image *ResizeImage(Image *image,const size_t columns,
3055 % const size_t rows, const ResizeFilter* filter,
3056 % ExceptionInfo *exception)
3058 % A description of each parameter follows:
3060 % o image: the image.
3062 % o columns: the number of columns in the scaled image.
3064 % o rows: the number of rows in the scaled image.
3066 % o filter: Image filter to use.
3068 % o exception: return any errors or warnings in this structure.
3072 static MagickBooleanType resizeHorizontalFilter(cl_mem image,
3073 const unsigned int imageColumns,const unsigned int imageRows,
3074 const unsigned int matte,cl_mem resizedImage,
3075 const unsigned int resizedColumns,const unsigned int resizedRows,
3076 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3077 const float xFactor,MagickCLEnv clEnv,cl_command_queue queue,
3078 ExceptionInfo *exception)
3086 workgroupSize = 256;
3090 resizeFilterSupport,
3091 resizeFilterWindowSupport,
3104 status = MagickFalse;
3107 deviceLocalMemorySize,
3108 gammaAccumulatorLocalMemorySize,
3109 global_work_size[2],
3110 imageCacheLocalMemorySize,
3111 pixelAccumulatorLocalMemorySize,
3113 totalLocalMemorySize,
3114 weightAccumulatorLocalMemorySize;
3121 horizontalKernel = NULL;
3122 status = MagickFalse;
3125 Apply filter to resize vertically from image to resize image.
3127 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
3128 support=scale*GetResizeFilterSupport(resizeFilter);
3132 Support too small even for nearest neighbour: Reduce to point
3135 support=(MagickRealType) 0.5;
3138 scale=PerceptibleReciprocal(scale);
3140 if (resizedColumns < workgroupSize)
3143 pixelPerWorkgroup = 32;
3147 chunkSize = workgroupSize;
3148 pixelPerWorkgroup = workgroupSize;
3151 /* get the local memory size supported by the device */
3152 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
3154 DisableMSCWarning(4127)
3158 /* calculate the local memory size needed per workgroup */
3159 cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
3160 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
3161 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
3162 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
3163 totalLocalMemorySize = imageCacheLocalMemorySize;
3165 /* local size for the pixel accumulator */
3166 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
3167 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
3169 /* local memory size for the weight accumulator */
3170 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
3171 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
3173 /* local memory size for the gamma accumulator */
3175 gammaAccumulatorLocalMemorySize = sizeof(float);
3177 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
3178 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
3180 if (totalLocalMemorySize <= deviceLocalMemorySize)
3184 pixelPerWorkgroup = pixelPerWorkgroup/2;
3185 chunkSize = chunkSize/2;
3186 if (pixelPerWorkgroup == 0
3189 /* quit, fallback to CPU */
3195 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
3196 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
3199 if (resizeFilterType == SincFastWeightingFunction
3200 && resizeWindowType == SincFastWeightingFunction)
3202 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilterSinc");
3206 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter");
3208 if (horizontalKernel == NULL)
3210 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3215 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&image);
3216 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns);
3217 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows);
3218 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
3219 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
3220 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
3222 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
3223 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
3225 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
3226 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
3227 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
3229 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
3230 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
3232 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
3233 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
3235 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
3236 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
3238 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
3239 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
3242 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
3243 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
3244 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
3245 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
3248 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
3249 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
3250 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
3252 if (clStatus != CL_SUCCESS)
3254 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3258 global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
3259 global_work_size[1] = resizedRows;
3261 local_work_size[0] = workgroupSize;
3262 local_work_size[1] = 1;
3263 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
3264 (void) local_work_size;
3265 if (clStatus != CL_SUCCESS)
3267 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3270 clEnv->library->clFlush(queue);
3271 status = MagickTrue;
3275 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3277 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
3282 static MagickBooleanType resizeVerticalFilter(cl_mem image,
3283 const unsigned int imageColumns,const unsigned int imageRows,
3284 const unsigned int matte,cl_mem resizedImage,
3285 const unsigned int resizedColumns,const unsigned int resizedRows,
3286 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3287 const float yFactor,MagickCLEnv clEnv,cl_command_queue queue,
3288 ExceptionInfo *exception)
3296 workgroupSize = 256;
3300 resizeFilterSupport,
3301 resizeFilterWindowSupport,
3314 status = MagickFalse;
3317 deviceLocalMemorySize,
3318 gammaAccumulatorLocalMemorySize,
3319 global_work_size[2],
3320 imageCacheLocalMemorySize,
3321 pixelAccumulatorLocalMemorySize,
3323 totalLocalMemorySize,
3324 weightAccumulatorLocalMemorySize;
3331 horizontalKernel = NULL;
3332 status = MagickFalse;
3335 Apply filter to resize vertically from image to resize image.
3337 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
3338 support=scale*GetResizeFilterSupport(resizeFilter);
3342 Support too small even for nearest neighbour: Reduce to point
3345 support=(MagickRealType) 0.5;
3348 scale=PerceptibleReciprocal(scale);
3350 if (resizedRows < workgroupSize)
3353 pixelPerWorkgroup = 32;
3357 chunkSize = workgroupSize;
3358 pixelPerWorkgroup = workgroupSize;
3361 /* get the local memory size supported by the device */
3362 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
3364 DisableMSCWarning(4127)
3368 /* calculate the local memory size needed per workgroup */
3369 cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
3370 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
3371 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
3372 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
3373 totalLocalMemorySize = imageCacheLocalMemorySize;
3375 /* local size for the pixel accumulator */
3376 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
3377 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
3379 /* local memory size for the weight accumulator */
3380 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
3381 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
3383 /* local memory size for the gamma accumulator */
3385 gammaAccumulatorLocalMemorySize = sizeof(float);
3387 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
3388 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
3390 if (totalLocalMemorySize <= deviceLocalMemorySize)
3394 pixelPerWorkgroup = pixelPerWorkgroup/2;
3395 chunkSize = chunkSize/2;
3396 if (pixelPerWorkgroup == 0
3399 /* quit, fallback to CPU */
3405 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
3406 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
3408 if (resizeFilterType == SincFastWeightingFunction
3409 && resizeWindowType == SincFastWeightingFunction)
3410 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilterSinc");
3412 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter");
3414 if (horizontalKernel == NULL)
3416 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3421 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&image);
3422 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns);
3423 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows);
3424 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
3425 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&yFactor);
3426 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
3428 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
3429 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
3431 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
3432 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
3433 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
3435 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
3436 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
3438 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
3439 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
3441 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
3442 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
3444 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
3445 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
3448 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
3449 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
3450 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
3451 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
3454 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
3455 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
3456 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
3458 if (clStatus != CL_SUCCESS)
3460 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3464 global_work_size[0] = resizedColumns;
3465 global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
3467 local_work_size[0] = 1;
3468 local_work_size[1] = workgroupSize;
3469 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
3470 if (clStatus != CL_SUCCESS)
3472 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3475 clEnv->library->clFlush(queue);
3476 status = MagickTrue;
3480 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3482 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
3487 static Image *ComputeResizeImage(const Image* image,
3488 const size_t resizedColumns,const size_t resizedRows,
3489 const ResizeFilter *resizeFilter,ExceptionInfo *exception)
3492 *filteredImage_view,
3505 cubicCoefficientsBuffer,
3506 filteredImageBuffer,
3514 *resizeFilterCoefficient;
3520 *mappedCoefficientBuffer,
3544 outputReady = MagickFalse;
3545 filteredImage = NULL;
3546 filteredImage_view = NULL;
3550 tempImageBuffer = NULL;
3551 filteredImageBuffer = NULL;
3552 cubicCoefficientsBuffer = NULL;
3555 clEnv = GetDefaultOpenCLEnv();
3556 context = GetOpenCLContext(clEnv);
3558 /* Create and initialize OpenCL buffers. */
3559 image_view=AcquireVirtualCacheView(image,exception);
3560 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
3561 if (inputPixels == (const void *) NULL)
3563 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
3567 /* If the host pointer is aligned to the size of CLPixelPacket,
3568 then use the host buffer directly from the GPU; otherwise,
3569 create a buffer on the GPU and copy the data over */
3570 if (ALIGNED(inputPixels,CLPixelPacket))
3572 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
3576 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
3578 /* create a CL buffer from image pixel buffer */
3579 length = image->columns * image->rows;
3580 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3581 if (clStatus != CL_SUCCESS)
3583 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3587 cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus);
3588 if (clStatus != CL_SUCCESS)
3590 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3593 queue = AcquireOpenCLCommandQueue(clEnv);
3594 mappedCoefficientBuffer = (float*)clEnv->library->clEnqueueMapBuffer(queue, cubicCoefficientsBuffer, CL_TRUE, CL_MAP_WRITE, 0, 7 * sizeof(float)
3595 , 0, NULL, NULL, &clStatus);
3596 if (clStatus != CL_SUCCESS)
3598 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
3601 resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter);
3602 for (i = 0; i < 7; i++)
3604 mappedCoefficientBuffer[i] = (float) resizeFilterCoefficient[i];
3606 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL);
3607 if (clStatus != CL_SUCCESS)
3609 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
3613 filteredImage = CloneImage(image,resizedColumns,resizedRows,MagickTrue,exception);
3614 if (filteredImage == NULL)
3617 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
3619 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
3622 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
3623 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
3624 if (filteredPixels == (void *) NULL)
3626 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
3630 if (ALIGNED(filteredPixels,CLPixelPacket))
3632 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
3633 hostPtr = filteredPixels;
3637 mem_flags = CL_MEM_WRITE_ONLY;
3641 /* create a CL buffer from image pixel buffer */
3642 length = filteredImage->columns * filteredImage->rows;
3643 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
3644 if (clStatus != CL_SUCCESS)
3646 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3650 xFactor=(float) resizedColumns/(float) image->columns;
3651 yFactor=(float) resizedRows/(float) image->rows;
3652 if (xFactor > yFactor)
3655 length = resizedColumns*image->rows;
3656 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
3657 if (clStatus != CL_SUCCESS)
3659 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3663 status = resizeHorizontalFilter(imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, (image->alpha_trait != BlendPixelTrait)?1:0
3664 , tempImageBuffer, (unsigned int) resizedColumns, (unsigned int) image->rows
3665 , resizeFilter, cubicCoefficientsBuffer
3666 , xFactor, clEnv, queue, exception);
3667 if (status != MagickTrue)
3670 status = resizeVerticalFilter(tempImageBuffer, (unsigned int) resizedColumns, (unsigned int) image->rows, (image->alpha_trait != BlendPixelTrait)?1:0
3671 , filteredImageBuffer, (unsigned int) resizedColumns, (unsigned int) resizedRows
3672 , resizeFilter, cubicCoefficientsBuffer
3673 , yFactor, clEnv, queue, exception);
3674 if (status != MagickTrue)
3679 length = image->columns*resizedRows;
3680 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
3681 if (clStatus != CL_SUCCESS)
3683 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3687 status = resizeVerticalFilter(imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, (image->alpha_trait != BlendPixelTrait)?1:0
3688 , tempImageBuffer, (unsigned int) image->columns, (unsigned int) resizedRows
3689 , resizeFilter, cubicCoefficientsBuffer
3690 , yFactor, clEnv, queue, exception);
3691 if (status != MagickTrue)
3694 status = resizeHorizontalFilter(tempImageBuffer, (unsigned int) image->columns, (unsigned int) resizedRows, (image->alpha_trait != BlendPixelTrait)?1:0
3695 , filteredImageBuffer, (unsigned int) resizedColumns, (unsigned int) resizedRows
3696 , resizeFilter, cubicCoefficientsBuffer
3697 , xFactor, clEnv, queue, exception);
3698 if (status != MagickTrue)
3701 length = resizedColumns*resizedRows;
3702 if (ALIGNED(filteredPixels,CLPixelPacket))
3704 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3708 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
3710 if (clStatus != CL_SUCCESS)
3712 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3715 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
3718 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3720 image_view=DestroyCacheView(image_view);
3721 if (filteredImage_view != NULL)
3722 filteredImage_view=DestroyCacheView(filteredImage_view);
3724 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
3725 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
3726 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
3727 if (cubicCoefficientsBuffer!=NULL) clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer);
3728 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3729 if (outputReady == MagickFalse && filteredImage != NULL)
3730 filteredImage=DestroyImage(filteredImage);
3731 return(filteredImage);
3734 const ResizeWeightingFunctionType supportedResizeWeighting[] =
3736 BoxWeightingFunction,
3737 TriangleWeightingFunction,
3738 HanningWeightingFunction,
3739 HammingWeightingFunction,
3740 BlackmanWeightingFunction,
3741 CubicBCWeightingFunction,
3742 SincWeightingFunction,
3743 SincFastWeightingFunction,
3744 LastWeightingFunction
3747 static MagickBooleanType gpuSupportedResizeWeighting(
3748 ResizeWeightingFunctionType f)
3755 if (supportedResizeWeighting[i] == LastWeightingFunction)
3757 if (supportedResizeWeighting[i] == f)
3760 return(MagickFalse);
3763 MagickExport Image *AccelerateResizeImage(const Image *image,
3764 const size_t resizedColumns,const size_t resizedRows,
3765 const ResizeFilter *resizeFilter,ExceptionInfo *exception)
3770 assert(image != NULL);
3771 assert(exception != (ExceptionInfo *) NULL);
3773 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3774 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
3777 if (gpuSupportedResizeWeighting(GetResizeFilterWeightingType(resizeFilter)) == MagickFalse ||
3778 gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(resizeFilter)) == MagickFalse)
3781 filteredImage=ComputeResizeImage(image,resizedColumns,resizedRows,resizeFilter,exception);
3782 return(filteredImage);
3786 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3790 % C o n t r a s t I m a g e w i t h O p e n C L %
3794 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3796 % ContrastImage() enhances the intensity differences between the lighter and
3797 % darker elements of the image. Set sharpen to a MagickTrue to increase the
3798 % image contrast otherwise the contrast is reduced.
3800 % The format of the ContrastImage method is:
3802 % MagickBooleanType ContrastImage(Image *image,
3803 % const MagickBooleanType sharpen)
3805 % A description of each parameter follows:
3807 % o image: the image.
3809 % o sharpen: Increase or decrease image contrast.
3813 static MagickBooleanType ComputeContrastImage(Image *image,
3814 const MagickBooleanType sharpen,ExceptionInfo *exception)
3847 global_work_size[2];
3856 outputReady = MagickFalse;
3861 filterKernel = NULL;
3864 clEnv = GetDefaultOpenCLEnv();
3865 context = GetOpenCLContext(clEnv);
3867 /* Create and initialize OpenCL buffers. */
3868 image_view=AcquireAuthenticCacheView(image,exception);
3869 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
3870 if (inputPixels == (void *) NULL)
3872 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
3876 /* If the host pointer is aligned to the size of CLPixelPacket,
3877 then use the host buffer directly from the GPU; otherwise,
3878 create a buffer on the GPU and copy the data over */
3879 if (ALIGNED(inputPixels,CLPixelPacket))
3881 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3885 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3887 /* create a CL buffer from image pixel buffer */
3888 length = image->columns * image->rows;
3889 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3890 if (clStatus != CL_SUCCESS)
3892 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3896 filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast");
3897 if (filterKernel == NULL)
3899 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3904 clStatus=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
3906 uSharpen = (sharpen == MagickFalse)?0:1;
3907 clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen);
3908 if (clStatus != CL_SUCCESS)
3910 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3914 global_work_size[0] = image->columns;
3915 global_work_size[1] = image->rows;
3916 /* launch the kernel */
3917 queue = AcquireOpenCLCommandQueue(clEnv);
3918 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3919 if (clStatus != CL_SUCCESS)
3921 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3924 clEnv->library->clFlush(queue);
3926 if (ALIGNED(inputPixels,CLPixelPacket))
3928 length = image->columns * image->rows;
3929 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3933 length = image->columns * image->rows;
3934 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3936 if (clStatus != CL_SUCCESS)
3938 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3941 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
3944 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3946 image_view=DestroyCacheView(image_view);
3948 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
3949 if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel);
3950 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3951 return(outputReady);
3954 MagickExport MagickBooleanType AccelerateContrastImage(Image *image,
3955 const MagickBooleanType sharpen,ExceptionInfo *exception)
3960 assert(image != NULL);
3961 assert(exception != (ExceptionInfo *) NULL);
3963 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3964 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
3965 return(MagickFalse);
3967 status = ComputeContrastImage(image,sharpen,exception);
3972 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3976 % M o d u l a t e I m a g e w i t h O p e n C L %
3980 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3982 % ModulateImage() lets you control the brightness, saturation, and hue
3983 % of an image. Modulate represents the brightness, saturation, and hue
3984 % as one parameter (e.g. 90,150,100). If the image colorspace is HSL, the
3985 % modulation is lightness, saturation, and hue. For HWB, use blackness,
3986 % whiteness, and hue. And for HCL, use chrome, luma, and hue.
3988 % The format of the ModulateImage method is:
3990 % MagickBooleanType ModulateImage(Image *image,const char *modulate)
3992 % A description of each parameter follows:
3994 % o image: the image.
3996 % o percent_*: Define the percent change in brightness, saturation, and
4001 MagickBooleanType ComputeModulateImage(Image *image,
4002 double percent_brightness,double percent_hue,double percent_saturation,
4003 ColorspaceType colorspace,ExceptionInfo *exception)
4049 modulateKernel = NULL;
4051 assert(image != (Image *) NULL);
4052 assert(image->signature == MagickSignature);
4053 if (image->debug != MagickFalse)
4054 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
4057 * initialize opencl env
4059 clEnv = GetDefaultOpenCLEnv();
4060 context = GetOpenCLContext(clEnv);
4061 queue = AcquireOpenCLCommandQueue(clEnv);
4063 outputReady = MagickFalse;
4065 /* Create and initialize OpenCL buffers.
4066 inputPixels = AcquirePixelCachePixels(image, &length, exception);
4067 assume this will get a writable image
4069 image_view=AcquireAuthenticCacheView(image,exception);
4070 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
4071 if (inputPixels == (void *) NULL)
4073 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
4077 /* If the host pointer is aligned to the size of CLPixelPacket,
4078 then use the host buffer directly from the GPU; otherwise,
4079 create a buffer on the GPU and copy the data over
4081 if (ALIGNED(inputPixels,CLPixelPacket))
4083 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4087 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4089 /* create a CL buffer from image pixel buffer */
4090 length = image->columns * image->rows;
4091 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4092 if (clStatus != CL_SUCCESS)
4094 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4098 modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate");
4099 if (modulateKernel == NULL)
4101 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4105 bright=percent_brightness;
4107 saturation=percent_saturation;
4111 clStatus=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4112 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&bright);
4113 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&hue);
4114 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&saturation);
4115 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color);
4116 if (clStatus != CL_SUCCESS)
4118 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4119 printf("no kernel\n");
4124 size_t global_work_size[2];
4125 global_work_size[0] = image->columns;
4126 global_work_size[1] = image->rows;
4127 /* launch the kernel */
4128 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4129 if (clStatus != CL_SUCCESS)
4131 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4134 clEnv->library->clFlush(queue);
4137 if (ALIGNED(inputPixels,CLPixelPacket))
4139 length = image->columns * image->rows;
4140 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4144 length = image->columns * image->rows;
4145 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
4147 if (clStatus != CL_SUCCESS)
4149 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4153 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
4156 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4158 image_view=DestroyCacheView(image_view);
4160 if (imageBuffer!=NULL)
4161 clEnv->library->clReleaseMemObject(imageBuffer);
4162 if (modulateKernel!=NULL)
4163 RelinquishOpenCLKernel(clEnv, modulateKernel);
4165 RelinquishOpenCLCommandQueue(clEnv, queue);
4171 MagickExport MagickBooleanType AccelerateModulateImage(Image *image,
4172 double percent_brightness,double percent_hue,double percent_saturation,
4173 ColorspaceType colorspace,ExceptionInfo *exception)
4178 assert(image != NULL);
4179 assert(exception != (ExceptionInfo *) NULL);
4181 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
4182 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
4183 return(MagickFalse);
4185 if ((colorspace != HSLColorspace && colorspace != UndefinedColorspace))
4186 return(MagickFalse);
4188 status = ComputeModulateImage(image,percent_brightness, percent_hue, percent_saturation, colorspace, exception);
4193 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4197 % G r a y s c a l e I m a g e w i t h O p e n C L %
4201 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4203 % GrayscaleImage() converts the colors in the reference image to gray.
4205 % The format of the GrayscaleImageChannel method is:
4207 % MagickBooleanType GrayscaleImage(Image *image,
4208 % const PixelIntensityMethod method)
4210 % A description of each parameter follows:
4212 % o image: the image.
4214 % o channel: the channel.
4218 MagickBooleanType ComputeGrayscaleImage(Image *image,
4219 const PixelIntensityMethod method,ExceptionInfo *exception)
4263 grayscaleKernel = NULL;
4265 assert(image != (Image *) NULL);
4266 assert(image->signature == MagickSignature);
4267 if (image->debug != MagickFalse)
4268 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
4271 * initialize opencl env
4273 clEnv = GetDefaultOpenCLEnv();
4274 context = GetOpenCLContext(clEnv);
4275 queue = AcquireOpenCLCommandQueue(clEnv);
4277 outputReady = MagickFalse;
4279 /* Create and initialize OpenCL buffers.
4280 inputPixels = AcquirePixelCachePixels(image, &length, exception);
4281 assume this will get a writable image
4283 image_view=AcquireAuthenticCacheView(image,exception);
4284 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
4285 if (inputPixels == (void *) NULL)
4287 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
4291 /* If the host pointer is aligned to the size of CLPixelPacket,
4292 then use the host buffer directly from the GPU; otherwise,
4293 create a buffer on the GPU and copy the data over
4295 if (ALIGNED(inputPixels,CLPixelPacket))
4297 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4301 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4303 /* create a CL buffer from image pixel buffer */
4304 length = image->columns * image->rows;
4305 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4306 if (clStatus != CL_SUCCESS)
4308 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4312 intensityMethod = method;
4313 colorspace = image->colorspace;
4315 grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Grayscale");
4316 if (grayscaleKernel == NULL)
4318 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4323 clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4324 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&intensityMethod);
4325 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&colorspace);
4326 if (clStatus != CL_SUCCESS)
4328 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4329 printf("no kernel\n");
4334 size_t global_work_size[2];
4335 global_work_size[0] = image->columns;
4336 global_work_size[1] = image->rows;
4337 /* launch the kernel */
4338 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, grayscaleKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4339 if (clStatus != CL_SUCCESS)
4341 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4344 clEnv->library->clFlush(queue);
4347 if (ALIGNED(inputPixels,CLPixelPacket))
4349 length = image->columns * image->rows;
4350 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4354 length = image->columns * image->rows;
4355 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
4357 if (clStatus != CL_SUCCESS)
4359 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4363 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
4366 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4368 image_view=DestroyCacheView(image_view);
4370 if (imageBuffer!=NULL)
4371 clEnv->library->clReleaseMemObject(imageBuffer);
4372 if (grayscaleKernel!=NULL)
4373 RelinquishOpenCLKernel(clEnv, grayscaleKernel);
4375 RelinquishOpenCLCommandQueue(clEnv, queue);
4377 return( outputReady);
4380 MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image,
4381 const PixelIntensityMethod method,ExceptionInfo *exception)
4386 assert(image != NULL);
4387 assert(exception != (ExceptionInfo *) NULL);
4389 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
4390 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
4391 return(MagickFalse);
4393 if (method == Rec601LuminancePixelIntensityMethod || method == Rec709LuminancePixelIntensityMethod)
4394 return(MagickFalse);
4396 if (image->colorspace != sRGBColorspace)
4397 return(MagickFalse);
4399 status=ComputeGrayscaleImage(image,method,exception);
4404 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4408 % E q u a l i z e I m a g e w i t h O p e n C L %
4412 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4414 % EqualizeImage() applies a histogram equalization to the image.
4416 % The format of the EqualizeImage method is:
4418 % MagickBooleanType EqualizeImage(Image *image)
4419 % MagickBooleanType EqualizeImageChannel(Image *image,
4420 % const ChannelType channel)
4422 % A description of each parameter follows:
4424 % o image: the image.
4426 % o channel: the channel.
4430 static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
4431 cl_command_queue queue,cl_mem imageBuffer,cl_mem histogramBuffer,
4432 Image *image,const ChannelType channel,ExceptionInfo *exception)
4449 global_work_size[2];
4451 histogramKernel = NULL;
4453 outputReady = MagickFalse;
4454 method = image->intensity;
4455 colorspace = image->colorspace;
4457 /* get the OpenCL kernel */
4458 histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram");
4459 if (histogramKernel == NULL)
4461 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4465 /* set the kernel arguments */
4467 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4468 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel);
4469 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&method);
4470 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&colorspace);
4471 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer);
4472 if (clStatus != CL_SUCCESS)
4474 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4478 /* launch the kernel */
4479 global_work_size[0] = image->columns;
4480 global_work_size[1] = image->rows;
4482 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4484 if (clStatus != CL_SUCCESS)
4486 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4489 clEnv->library->clFlush(queue);
4491 outputReady = MagickTrue;
4494 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4496 if (histogramKernel!=NULL)
4497 RelinquishOpenCLKernel(clEnv, histogramKernel);
4499 return(outputReady);
4502 MagickExport MagickBooleanType ComputeEqualizeImage(Image *image,
4503 const ChannelType channel,ExceptionInfo *exception)
4505 #define EqualizeImageTag "Equalize/Image"
4557 global_work_size[2];
4568 histogramBuffer = NULL;
4569 equalizeMapBuffer = NULL;
4570 histogramKernel = NULL;
4571 equalizeKernel = NULL;
4574 outputReady = MagickFalse;
4576 assert(image != (Image *) NULL);
4577 assert(image->signature == MagickSignature);
4578 if (image->debug != MagickFalse)
4579 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
4582 * initialize opencl env
4584 clEnv = GetDefaultOpenCLEnv();
4585 context = GetOpenCLContext(clEnv);
4586 queue = AcquireOpenCLCommandQueue(clEnv);
4589 Allocate and initialize histogram arrays.
4591 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram));
4592 if (histogram == (cl_uint4 *) NULL)
4593 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
4595 /* reset histogram */
4596 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram));
4598 /* Create and initialize OpenCL buffers. */
4599 /* inputPixels = AcquirePixelCachePixels(image, &length, exception); */
4600 /* assume this will get a writable image */
4601 image_view=AcquireAuthenticCacheView(image,exception);
4602 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
4604 if (inputPixels == (void *) NULL)
4606 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
4609 /* If the host pointer is aligned to the size of CLPixelPacket,
4610 then use the host buffer directly from the GPU; otherwise,
4611 create a buffer on the GPU and copy the data over */
4612 if (ALIGNED(inputPixels,CLPixelPacket))
4614 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4618 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
4620 /* create a CL buffer from image pixel buffer */
4621 length = image->columns * image->rows;
4622 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4623 if (clStatus != CL_SUCCESS)
4625 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4629 /* If the host pointer is aligned to the size of cl_uint,
4630 then use the host buffer directly from the GPU; otherwise,
4631 create a buffer on the GPU and copy the data over */
4632 if (ALIGNED(histogram,cl_uint4))
4634 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4635 hostPtr = histogram;
4639 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4640 hostPtr = histogram;
4642 /* create a CL buffer for histogram */
4643 length = (MaxMap+1);
4644 histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
4645 if (clStatus != CL_SUCCESS)
4647 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4651 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, channel, exception);
4652 if (status == MagickFalse)
4655 /* read from the kenel output */
4656 if (ALIGNED(histogram,cl_uint4))
4658 length = (MaxMap+1);
4659 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
4663 length = (MaxMap+1);
4664 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
4666 if (clStatus != CL_SUCCESS)
4668 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4672 /* unmap, don't block gpu to use this buffer again. */
4673 if (ALIGNED(histogram,cl_uint4))
4675 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
4676 if (clStatus != CL_SUCCESS)
4678 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
4683 /* recreate input buffer later, in case image updated */
4684 #ifdef RECREATEBUFFER
4685 if (imageBuffer!=NULL)
4686 clEnv->library->clReleaseMemObject(imageBuffer);
4690 equalize_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*equalize_map));
4691 if (equalize_map == (PixelPacket *) NULL)
4692 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
4694 map=(FloatPixelPacket *) AcquireQuantumMemory(MaxMap+1UL,sizeof(*map));
4695 if (map == (FloatPixelPacket *) NULL)
4696 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
4699 Integrate the histogram to get the equalization map.
4701 (void) ResetMagickMemory(&intensity,0,sizeof(intensity));
4702 for (i=0; i <= (ssize_t) MaxMap; i++)
4704 if ((channel & SyncChannels) != 0)
4706 intensity.red+=histogram[i].s[2];
4710 if ((channel & RedChannel) != 0)
4711 intensity.red+=histogram[i].s[2];
4712 if ((channel & GreenChannel) != 0)
4713 intensity.green+=histogram[i].s[1];
4714 if ((channel & BlueChannel) != 0)
4715 intensity.blue+=histogram[i].s[0];
4716 if ((channel & OpacityChannel) != 0)
4717 intensity.alpha+=histogram[i].s[3];
4719 if (((channel & IndexChannel) != 0) &&
4720 (image->colorspace == CMYKColorspace))
4722 intensity.index+=histogram[i].index;
4728 white=map[(int) MaxMap];
4729 (void) ResetMagickMemory(equalize_map,0,(MaxMap+1)*sizeof(*equalize_map));
4730 for (i=0; i <= (ssize_t) MaxMap; i++)
4732 if ((channel & SyncChannels) != 0)
4734 if (white.red != black.red)
4735 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4736 (map[i].red-black.red))/(white.red-black.red)));
4739 if (((channel & RedChannel) != 0) && (white.red != black.red))
4740 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4741 (map[i].red-black.red))/(white.red-black.red)));
4742 if (((channel & GreenChannel) != 0) && (white.green != black.green))
4743 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4744 (map[i].green-black.green))/(white.green-black.green)));
4745 if (((channel & BlueChannel) != 0) && (white.blue != black.blue))
4746 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4747 (map[i].blue-black.blue))/(white.blue-black.blue)));
4748 if (((channel & OpacityChannel) != 0) && (white.alpha != black.alpha))
4749 equalize_map[i].alpha=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4750 (map[i].alpha-black.alpha))/(white.alpha-black.alpha)));
4752 if ((((channel & IndexChannel) != 0) &&
4753 (image->colorspace == CMYKColorspace)) &&
4754 (white.index != black.index))
4755 equalize_map[i].index=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4756 (map[i].index-black.index))/(white.index-black.index)));
4760 if (image->storage_class == PseudoClass)
4765 for (i=0; i < (ssize_t) image->colors; i++)
4767 if ((channel & SyncChannels) != 0)
4769 if (white.red != black.red)
4771 image->colormap[i].red=equalize_map[
4772 ScaleQuantumToMap(image->colormap[i].red)].red;
4773 image->colormap[i].green=equalize_map[
4774 ScaleQuantumToMap(image->colormap[i].green)].red;
4775 image->colormap[i].blue=equalize_map[
4776 ScaleQuantumToMap(image->colormap[i].blue)].red;
4777 image->colormap[i].alpha=equalize_map[
4778 ScaleQuantumToMap(image->colormap[i].alpha)].red;
4782 if (((channel & RedChannel) != 0) && (white.red != black.red))
4783 image->colormap[i].red=equalize_map[
4784 ScaleQuantumToMap(image->colormap[i].red)].red;
4785 if (((channel & GreenChannel) != 0) && (white.green != black.green))
4786 image->colormap[i].green=equalize_map[
4787 ScaleQuantumToMap(image->colormap[i].green)].green;
4788 if (((channel & BlueChannel) != 0) && (white.blue != black.blue))
4789 image->colormap[i].blue=equalize_map[
4790 ScaleQuantumToMap(image->colormap[i].blue)].blue;
4791 if (((channel & OpacityChannel) != 0) &&
4792 (white.alpha != black.alpha))
4793 image->colormap[i].alpha=equalize_map[
4794 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
4802 /* GPU can work on this again, image and equalize map as input
4803 image: uchar4 (CLPixelPacket)
4804 equalize_map: uchar4 (PixelPacket)
4805 black, white: float4 (FloatPixelPacket) */
4807 #ifdef RECREATEBUFFER
4808 /* If the host pointer is aligned to the size of CLPixelPacket,
4809 then use the host buffer directly from the GPU; otherwise,
4810 create a buffer on the GPU and copy the data over */
4811 if (ALIGNED(inputPixels,CLPixelPacket))
4813 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4817 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4819 /* create a CL buffer from image pixel buffer */
4820 length = image->columns * image->rows;
4821 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4822 if (clStatus != CL_SUCCESS)
4824 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4829 /* Create and initialize OpenCL buffers. */
4830 if (ALIGNED(equalize_map, PixelPacket))
4832 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4833 hostPtr = equalize_map;
4837 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4838 hostPtr = equalize_map;
4840 /* create a CL buffer for eqaulize_map */
4841 length = (MaxMap+1);
4842 equalizeMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
4843 if (clStatus != CL_SUCCESS)
4845 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4849 /* get the OpenCL kernel */
4850 equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize");
4851 if (equalizeKernel == NULL)
4853 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4857 /* set the kernel arguments */
4859 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4860 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(ChannelType),&channel);
4861 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&equalizeMapBuffer);
4862 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&white);
4863 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black);
4864 if (clStatus != CL_SUCCESS)
4866 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4870 /* launch the kernel */
4871 global_work_size[0] = image->columns;
4872 global_work_size[1] = image->rows;
4874 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4876 if (clStatus != CL_SUCCESS)
4878 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4881 clEnv->library->clFlush(queue);
4883 /* read the data back */
4884 if (ALIGNED(inputPixels,CLPixelPacket))
4886 length = image->columns * image->rows;
4887 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4891 length = image->columns * image->rows;
4892 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
4894 if (clStatus != CL_SUCCESS)
4896 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4900 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
4903 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4905 image_view=DestroyCacheView(image_view);
4907 if (imageBuffer!=NULL)
4908 clEnv->library->clReleaseMemObject(imageBuffer);
4911 map=(FloatPixelPacket *) RelinquishMagickMemory(map);
4913 if (equalizeMapBuffer!=NULL)
4914 clEnv->library->clReleaseMemObject(equalizeMapBuffer);
4915 if (equalize_map!=NULL)
4916 equalize_map=(PixelPacket *) RelinquishMagickMemory(equalize_map);
4918 if (histogramBuffer!=NULL)
4919 clEnv->library->clReleaseMemObject(histogramBuffer);
4920 if (histogram!=NULL)
4921 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
4923 if (histogramKernel!=NULL)
4924 RelinquishOpenCLKernel(clEnv, histogramKernel);
4925 if (equalizeKernel!=NULL)
4926 RelinquishOpenCLKernel(clEnv, equalizeKernel);
4929 RelinquishOpenCLCommandQueue(clEnv, queue);
4931 return(outputReady);
4934 MagickExport MagickBooleanType AccelerateEqualizeImage(Image *image,
4935 const ChannelType channel,ExceptionInfo *exception)
4940 assert(image != NULL);
4941 assert(exception != (ExceptionInfo *) NULL);
4943 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
4944 (checkAccelerateCondition(image, channel) == MagickFalse) ||
4945 (checkHistogramCondition(image, channel) == MagickFalse))
4946 return(MagickFalse);
4948 status=ComputeEqualizeImage(image,channel,exception);
4953 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4957 % C o n t r a s t S t r e t c h I m a g e w i t h O p e n C L %
4961 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4963 % ContrastStretchImage() is a simple image enhancement technique that attempts
4964 % to improve the contrast in an image by `stretching' the range of intensity
4965 % values it contains to span a desired range of values. It differs from the
4966 % more sophisticated histogram equalization in that it can only apply a
4967 % linear scaling function to the image pixel values. As a result the
4968 % `enhancement' is less harsh.
4970 % The format of the ContrastStretchImage method is:
4972 % MagickBooleanType ContrastStretchImage(Image *image,
4973 % const char *levels)
4974 % MagickBooleanType ContrastStretchImageChannel(Image *image,
4975 % const size_t channel,const double black_point,
4976 % const double white_point)
4978 % A description of each parameter follows:
4980 % o image: the image.
4982 % o channel: the channel.
4984 % o black_point: the black point.
4986 % o white_point: the white point.
4988 % o levels: Specify the levels where the black and white points have the
4989 % range of 0 to number-of-pixels (e.g. 1%, 10x90%, etc.).
4993 MagickExport MagickBooleanType ComputeContrastStretchImageChannel(Image *image,
4994 const ChannelType channel,const double black_point,const double white_point,
4995 ExceptionInfo *exception)
4997 #define ContrastStretchImageTag "ContrastStretch/Image"
4998 #define MaxRange(color) ((MagickRealType) ScaleQuantumToMap((Quantum) (color)))
5051 global_work_size[2];
5061 histogramBuffer = NULL;
5062 stretchMapBuffer = NULL;
5063 histogramKernel = NULL;
5064 stretchKernel = NULL;
5067 outputReady = MagickFalse;
5070 assert(image != (Image *) NULL);
5071 assert(image->signature == MagickSignature);
5072 if (image->debug != MagickFalse)
5073 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
5075 //exception=(&image->exception);
5078 * initialize opencl env
5080 clEnv = GetDefaultOpenCLEnv();
5081 context = GetOpenCLContext(clEnv);
5082 queue = AcquireOpenCLCommandQueue(clEnv);
5085 Allocate and initialize histogram arrays.
5087 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram));
5089 if (histogram == (cl_uint4 *) NULL)
5090 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed", image->filename);
5092 /* reset histogram */
5093 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram));
5096 if (IsGrayImage(image,exception) != MagickFalse)
5097 (void) SetImageColorspace(image,GRAYColorspace);
5106 /* Create and initialize OpenCL buffers. */
5107 /* inputPixels = AcquirePixelCachePixels(image, &length, exception); */
5108 /* assume this will get a writable image */
5109 image_view=AcquireAuthenticCacheView(image,exception);
5110 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
5112 if (inputPixels == (void *) NULL)
5114 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
5117 /* If the host pointer is aligned to the size of CLPixelPacket,
5118 then use the host buffer directly from the GPU; otherwise,
5119 create a buffer on the GPU and copy the data over */
5120 if (ALIGNED(inputPixels,CLPixelPacket))
5122 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
5126 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
5128 /* create a CL buffer from image pixel buffer */
5129 length = image->columns * image->rows;
5130 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
5131 if (clStatus != CL_SUCCESS)
5133 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5137 /* If the host pointer is aligned to the size of cl_uint,
5138 then use the host buffer directly from the GPU; otherwise,
5139 create a buffer on the GPU and copy the data over */
5140 if (ALIGNED(histogram,cl_uint4))
5142 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
5143 hostPtr = histogram;
5147 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
5148 hostPtr = histogram;
5150 /* create a CL buffer for histogram */
5151 length = (MaxMap+1);
5152 histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
5153 if (clStatus != CL_SUCCESS)
5155 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5159 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, channel, exception);
5160 if (status == MagickFalse)
5163 /* read from the kenel output */
5164 if (ALIGNED(histogram,cl_uint4))
5166 length = (MaxMap+1);
5167 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
5171 length = (MaxMap+1);
5172 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
5174 if (clStatus != CL_SUCCESS)
5176 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
5180 /* unmap, don't block gpu to use this buffer again. */
5181 if (ALIGNED(histogram,cl_uint4))
5183 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
5184 if (clStatus != CL_SUCCESS)
5186 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
5191 /* recreate input buffer later, in case image updated */
5192 #ifdef RECREATEBUFFER
5193 if (imageBuffer!=NULL)
5194 clEnv->library->clReleaseMemObject(imageBuffer);
5199 Find the histogram boundaries by locating the black/white levels.
5202 white.red=MaxRange(QuantumRange);
5203 if ((channel & RedChannel) != 0)
5206 for (i=0; i <= (ssize_t) MaxMap; i++)
5208 intensity+=histogram[i].s[2];
5209 if (intensity > black_point)
5212 black.red=(MagickRealType) i;
5214 for (i=(ssize_t) MaxMap; i != 0; i--)
5216 intensity+=histogram[i].s[2];
5217 if (intensity > ((double) image->columns*image->rows-white_point))
5220 white.red=(MagickRealType) i;
5223 white.green=MaxRange(QuantumRange);
5224 if ((channel & GreenChannel) != 0)
5227 for (i=0; i <= (ssize_t) MaxMap; i++)
5229 intensity+=histogram[i].s[2];
5230 if (intensity > black_point)
5233 black.green=(MagickRealType) i;
5235 for (i=(ssize_t) MaxMap; i != 0; i--)
5237 intensity+=histogram[i].s[2];
5238 if (intensity > ((double) image->columns*image->rows-white_point))
5241 white.green=(MagickRealType) i;
5244 white.blue=MaxRange(QuantumRange);
5245 if ((channel & BlueChannel) != 0)
5248 for (i=0; i <= (ssize_t) MaxMap; i++)
5250 intensity+=histogram[i].s[2];
5251 if (intensity > black_point)
5254 black.blue=(MagickRealType) i;
5256 for (i=(ssize_t) MaxMap; i != 0; i--)
5258 intensity+=histogram[i].s[2];
5259 if (intensity > ((double) image->columns*image->rows-white_point))
5262 white.blue=(MagickRealType) i;
5265 white.alpha=MaxRange(QuantumRange);
5266 if ((channel & OpacityChannel) != 0)
5269 for (i=0; i <= (ssize_t) MaxMap; i++)
5271 intensity+=histogram[i].s[2];
5272 if (intensity > black_point)
5275 black.alpha=(MagickRealType) i;
5277 for (i=(ssize_t) MaxMap; i != 0; i--)
5279 intensity+=histogram[i].s[2];
5280 if (intensity > ((double) image->columns*image->rows-white_point))
5283 white.alpha=(MagickRealType) i;
5287 white.index=MaxRange(QuantumRange);
5288 if (((channel & IndexChannel) != 0) && (image->colorspace == CMYKColorspace))
5291 for (i=0; i <= (ssize_t) MaxMap; i++)
5293 intensity+=histogram[i].index;
5294 if (intensity > black_point)
5297 black.index=(MagickRealType) i;
5299 for (i=(ssize_t) MaxMap; i != 0; i--)
5301 intensity+=histogram[i].index;
5302 if (intensity > ((double) image->columns*image->rows-white_point))
5305 white.index=(MagickRealType) i;
5310 stretch_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL,
5311 sizeof(*stretch_map));
5313 if (stretch_map == (PixelPacket *) NULL)
5314 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed",
5318 Stretch the histogram to create the stretched image mapping.
5320 (void) ResetMagickMemory(stretch_map,0,(MaxMap+1)*sizeof(*stretch_map));
5321 for (i=0; i <= (ssize_t) MaxMap; i++)
5323 if ((channel & RedChannel) != 0)
5325 if (i < (ssize_t) black.red)
5326 stretch_map[i].red=(Quantum) 0;
5328 if (i > (ssize_t) white.red)
5329 stretch_map[i].red=QuantumRange;
5331 if (black.red != white.red)
5332 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap*
5333 (i-black.red)/(white.red-black.red)));
5335 if ((channel & GreenChannel) != 0)
5337 if (i < (ssize_t) black.green)
5338 stretch_map[i].green=0;
5340 if (i > (ssize_t) white.green)
5341 stretch_map[i].green=QuantumRange;
5343 if (black.green != white.green)
5344 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap*
5345 (i-black.green)/(white.green-black.green)));
5347 if ((channel & BlueChannel) != 0)
5349 if (i < (ssize_t) black.blue)
5350 stretch_map[i].blue=0;
5352 if (i > (ssize_t) white.blue)
5353 stretch_map[i].blue= QuantumRange;
5355 if (black.blue != white.blue)
5356 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap*
5357 (i-black.blue)/(white.blue-black.blue)));
5359 if ((channel & OpacityChannel) != 0)
5361 if (i < (ssize_t) black.alpha)
5362 stretch_map[i].alpha=0;
5364 if (i > (ssize_t) white.alpha)
5365 stretch_map[i].alpha=QuantumRange;
5367 if (black.alpha != white.alpha)
5368 stretch_map[i].alpha=ScaleMapToQuantum((MagickRealType) (MaxMap*
5369 (i-black.alpha)/(white.alpha-black.alpha)));
5372 if (((channel & IndexChannel) != 0) &&
5373 (image->colorspace == CMYKColorspace))
5375 if (i < (ssize_t) black.index)
5376 stretch_map[i].index=0;
5378 if (i > (ssize_t) white.index)
5379 stretch_map[i].index=QuantumRange;
5381 if (black.index != white.index)
5382 stretch_map[i].index=ScaleMapToQuantum((MagickRealType) (MaxMap*
5383 (i-black.index)/(white.index-black.index)));
5391 if (((channel & OpacityChannel) != 0) || (((channel & IndexChannel) != 0) &&
5392 (image->colorspace == CMYKColorspace)))
5393 image->storage_class=DirectClass;
5394 if (image->storage_class == PseudoClass)
5399 for (i=0; i < (ssize_t) image->colors; i++)
5401 if ((channel & RedChannel) != 0)
5403 if (black.red != white.red)
5404 image->colormap[i].red=stretch_map[
5405 ScaleQuantumToMap(image->colormap[i].red)].red;
5407 if ((channel & GreenChannel) != 0)
5409 if (black.green != white.green)
5410 image->colormap[i].green=stretch_map[
5411 ScaleQuantumToMap(image->colormap[i].green)].green;
5413 if ((channel & BlueChannel) != 0)
5415 if (black.blue != white.blue)
5416 image->colormap[i].blue=stretch_map[
5417 ScaleQuantumToMap(image->colormap[i].blue)].blue;
5419 if ((channel & OpacityChannel) != 0)
5421 if (black.alpha != white.alpha)
5422 image->colormap[i].alpha=stretch_map[
5423 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
5433 /* GPU can work on this again, image and equalize map as input
5434 image: uchar4 (CLPixelPacket)
5435 stretch_map: uchar4 (PixelPacket)
5436 black, white: float4 (FloatPixelPacket) */
5438 #ifdef RECREATEBUFFER
5439 /* If the host pointer is aligned to the size of CLPixelPacket,
5440 then use the host buffer directly from the GPU; otherwise,
5441 create a buffer on the GPU and copy the data over */
5442 if (ALIGNED(inputPixels,CLPixelPacket))
5444 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
5448 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
5450 /* create a CL buffer from image pixel buffer */
5451 length = image->columns * image->rows;
5452 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
5453 if (clStatus != CL_SUCCESS)
5455 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5460 /* Create and initialize OpenCL buffers. */
5461 if (ALIGNED(stretch_map, PixelPacket))
5463 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
5464 hostPtr = stretch_map;
5468 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
5469 hostPtr = stretch_map;
5471 /* create a CL buffer for stretch_map */
5472 length = (MaxMap+1);
5473 stretchMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
5474 if (clStatus != CL_SUCCESS)
5476 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5480 /* get the OpenCL kernel */
5481 stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Stretch");
5482 if (stretchKernel == NULL)
5484 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
5488 /* set the kernel arguments */
5490 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
5491 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(ChannelType),&channel);
5492 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&stretchMapBuffer);
5493 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&white);
5494 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black);
5495 if (clStatus != CL_SUCCESS)
5497 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5501 /* launch the kernel */
5502 global_work_size[0] = image->columns;
5503 global_work_size[1] = image->rows;
5505 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5507 if (clStatus != CL_SUCCESS)
5509 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5512 clEnv->library->clFlush(queue);
5514 /* read the data back */
5515 if (ALIGNED(inputPixels,CLPixelPacket))
5517 length = image->columns * image->rows;
5518 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
5522 length = image->columns * image->rows;
5523 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
5525 if (clStatus != CL_SUCCESS)
5527 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
5531 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
5534 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5536 image_view=DestroyCacheView(image_view);
5538 if (imageBuffer!=NULL)
5539 clEnv->library->clReleaseMemObject(imageBuffer);
5541 if (stretchMapBuffer!=NULL)
5542 clEnv->library->clReleaseMemObject(stretchMapBuffer);
5543 if (stretch_map!=NULL)
5544 stretch_map=(PixelPacket *) RelinquishMagickMemory(stretch_map);
5547 if (histogramBuffer!=NULL)
5548 clEnv->library->clReleaseMemObject(histogramBuffer);
5549 if (histogram!=NULL)
5550 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
5553 if (histogramKernel!=NULL)
5554 RelinquishOpenCLKernel(clEnv, histogramKernel);
5555 if (stretchKernel!=NULL)
5556 RelinquishOpenCLKernel(clEnv, stretchKernel);
5559 RelinquishOpenCLCommandQueue(clEnv, queue);
5561 return(outputReady);
5564 MagickExport MagickBooleanType AccelerateContrastStretchImageChannel(
5565 Image *image,const ChannelType channel,const double black_point,
5566 const double white_point,ExceptionInfo *exception)
5571 assert(image != NULL);
5572 assert(exception != (ExceptionInfo *) NULL);
5574 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
5575 (checkAccelerateCondition(image, channel) == MagickFalse) ||
5576 (checkHistogramCondition(image, channel) == MagickFalse))
5577 return(MagickFalse);
5579 status=ComputeContrastStretchImageChannel(image,channel, black_point, white_point, exception);
5584 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
5588 % D e s p e c k l e I m a g e w i t h O p e n C L %
5592 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
5594 % DespeckleImage() reduces the speckle noise in an image while perserving the
5595 % edges of the original image. A speckle removing filter uses a complementary
5596 % hulling technique (raising pixels that are darker than their surrounding
5597 % neighbors, then complementarily lowering pixels that are brighter than their
5598 % surrounding neighbors) to reduce the speckle index of that image (reference
5599 % Crimmins speckle removal).
5601 % The format of the DespeckleImage method is:
5603 % Image *DespeckleImage(const Image *image,ExceptionInfo *exception)
5605 % A description of each parameter follows:
5607 % o image: the image.
5609 % o exception: return any errors or warnings in this structure.
5613 static Image *ComputeDespeckleImage(const Image *image,
5614 ExceptionInfo*exception)
5617 X[4] = {0, 1, 1,-1},
5618 Y[4] = {1, 0, 1, 1};
5621 *filteredImage_view,
5641 filteredImageBuffer,
5665 global_work_size[2];
5675 outputReady = MagickFalse;
5678 filteredImage = NULL;
5679 filteredImage_view = NULL;
5680 filteredPixels = NULL;
5683 filteredImageBuffer = NULL;
5687 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
5688 clEnv = GetDefaultOpenCLEnv();
5689 context = GetOpenCLContext(clEnv);
5690 queue = AcquireOpenCLCommandQueue(clEnv);
5692 image_view=AcquireVirtualCacheView(image,exception);
5693 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
5694 if (inputPixels == (void *) NULL)
5696 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
5700 if (ALIGNED(inputPixels,CLPixelPacket))
5702 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
5706 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
5708 /* create a CL buffer from image pixel buffer */
5709 length = image->columns * image->rows;
5710 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
5711 if (clStatus != CL_SUCCESS)
5713 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5717 mem_flags = CL_MEM_READ_WRITE;
5718 length = image->columns * image->rows;
5719 for (k = 0; k < 2; k++)
5721 tempImageBuffer[k] = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus);
5722 if (clStatus != CL_SUCCESS)
5724 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5729 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
5730 assert(filteredImage != NULL);
5731 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
5733 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
5736 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
5737 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
5738 if (filteredPixels == (void *) NULL)
5740 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
5744 if (ALIGNED(filteredPixels,CLPixelPacket))
5746 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
5747 hostPtr = filteredPixels;
5751 mem_flags = CL_MEM_WRITE_ONLY;
5754 /* create a CL buffer from image pixel buffer */
5755 length = image->columns * image->rows;
5756 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
5757 if (clStatus != CL_SUCCESS)
5759 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5763 hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1");
5764 hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2");
5766 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&imageBuffer);
5767 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1));
5768 imageWidth = (unsigned int) image->columns;
5769 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth);
5770 imageHeight = (unsigned int) image->rows;
5771 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight);
5772 matte = (image->alpha_trait != BlendPixelTrait)?0:1;
5773 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
5774 if (clStatus != CL_SUCCESS)
5776 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5780 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1));
5781 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer);
5782 imageWidth = (unsigned int) image->columns;
5783 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth);
5784 imageHeight = (unsigned int) image->rows;
5785 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight);
5786 matte = (image->alpha_trait != BlendPixelTrait)?0:1;
5787 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
5788 if (clStatus != CL_SUCCESS)
5790 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5795 global_work_size[0] = image->columns;
5796 global_work_size[1] = image->rows;
5799 for (k = 0; k < 4; k++)
5808 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5809 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5810 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5811 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5812 if (clStatus != CL_SUCCESS)
5814 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5817 /* launch the kernel */
5818 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5819 if (clStatus != CL_SUCCESS)
5821 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5824 /* launch the kernel */
5825 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5826 if (clStatus != CL_SUCCESS)
5828 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5834 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
5835 offset.s[0] = -X[k];
5836 offset.s[1] = -Y[k];
5838 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5839 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5840 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5841 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5842 if (clStatus != CL_SUCCESS)
5844 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5847 /* launch the kernel */
5848 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5849 if (clStatus != CL_SUCCESS)
5851 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5854 /* launch the kernel */
5855 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5856 if (clStatus != CL_SUCCESS)
5858 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5862 offset.s[0] = -X[k];
5863 offset.s[1] = -Y[k];
5865 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5866 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5867 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5868 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5869 if (clStatus != CL_SUCCESS)
5871 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5874 /* launch the kernel */
5875 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5876 if (clStatus != CL_SUCCESS)
5878 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5881 /* launch the kernel */
5882 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5883 if (clStatus != CL_SUCCESS)
5885 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5892 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5893 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5894 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5895 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5898 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer);
5900 if (clStatus != CL_SUCCESS)
5902 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5905 /* launch the kernel */
5906 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5907 if (clStatus != CL_SUCCESS)
5909 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5912 /* launch the kernel */
5913 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5914 if (clStatus != CL_SUCCESS)
5916 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5921 if (ALIGNED(filteredPixels,CLPixelPacket))
5923 length = image->columns * image->rows;
5924 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
5928 length = image->columns * image->rows;
5929 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
5931 if (clStatus != CL_SUCCESS)
5933 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
5937 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
5940 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5942 image_view=DestroyCacheView(image_view);
5943 if (filteredImage_view != NULL)
5944 filteredImage_view=DestroyCacheView(filteredImage_view);
5946 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
5947 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
5948 for (k = 0; k < 2; k++)
5950 if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
5952 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
5953 if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
5954 if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
5955 if (outputReady == MagickFalse && filteredImage != NULL)
5956 filteredImage=DestroyImage(filteredImage);
5957 return(filteredImage);
5960 MagickExport Image *AccelerateDespeckleImage(const Image* image,
5961 ExceptionInfo* exception)
5966 assert(image != NULL);
5967 assert(exception != (ExceptionInfo *) NULL);
5969 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
5970 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
5973 filteredImage=ComputeDespeckleImage(image,exception);
5974 return(filteredImage);
5977 static Image *ComputeAddNoiseImage(const Image *image,
5978 const ChannelType channel,const NoiseType noise_type,
5979 ExceptionInfo *exception)
5982 *filteredImage_view,
6007 filteredImageBuffer,
6032 **restrict random_info;
6035 global_work_size[1],
6040 numRandomNumberPerPixel;
6042 #if defined(MAGICKCORE_OPENMP_SUPPORT)
6051 outputReady = MagickFalse;
6054 filteredImage = NULL;
6055 filteredImage_view = NULL;
6056 filteredPixels = NULL;
6059 filteredImageBuffer = NULL;
6061 addNoiseKernel = NULL;
6063 clEnv = GetDefaultOpenCLEnv();
6064 context = GetOpenCLContext(clEnv);
6065 queue = AcquireOpenCLCommandQueue(clEnv);
6067 image_view=AcquireVirtualCacheView(image,exception);
6068 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
6069 if (inputPixels == (void *) NULL)
6071 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
6075 if (ALIGNED(inputPixels,CLPixelPacket))
6077 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
6081 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
6083 /* create a CL buffer from image pixel buffer */
6084 length = image->columns * image->rows;
6085 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6086 if (clStatus != CL_SUCCESS)
6088 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6093 filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
6094 assert(filteredImage != NULL);
6095 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
6097 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
6100 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
6101 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
6102 if (filteredPixels == (void *) NULL)
6104 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
6108 if (ALIGNED(filteredPixels,CLPixelPacket))
6110 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
6111 hostPtr = filteredPixels;
6115 mem_flags = CL_MEM_WRITE_ONLY;
6118 /* create a CL buffer from image pixel buffer */
6119 length = image->columns * image->rows;
6120 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
6121 if (clStatus != CL_SUCCESS)
6123 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6127 /* find out how many random numbers needed by pixel */
6128 numRandomNumberPerPixel = 0;
6130 unsigned int numRandPerChannel = 0;
6135 case LaplacianNoise:
6138 numRandPerChannel = 1;
6141 case MultiplicativeGaussianNoise:
6143 numRandPerChannel = 2;
6147 if ((channel & RedChannel) != 0)
6148 numRandomNumberPerPixel+=numRandPerChannel;
6149 if ((channel & GreenChannel) != 0)
6150 numRandomNumberPerPixel+=numRandPerChannel;
6151 if ((channel & BlueChannel) != 0)
6152 numRandomNumberPerPixel+=numRandPerChannel;
6153 if ((channel & OpacityChannel) != 0)
6154 numRandomNumberPerPixel+=numRandPerChannel;
6157 /* set up the random number generators */
6159 option=GetImageArtifact(image,"attenuate");
6160 if (option != (char *) NULL)
6161 attenuate=StringToDouble(option,(char **) NULL);
6162 random_info=AcquireRandomInfoThreadSet();
6163 #if defined(MAGICKCORE_OPENMP_SUPPORT)
6164 key=GetRandomSecretKey(random_info[0]);
6168 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"GenerateNoiseImage");
6171 cl_uint computeUnitCount;
6172 cl_uint workItemCount;
6173 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &computeUnitCount, NULL);
6174 workItemCount = computeUnitCount * 2 * 256; // 256 work items per group, 2 groups per CU
6175 inputPixelCount = (cl_int) (image->columns * image->rows);
6176 pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount;
6177 pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4;
6179 local_work_size[0] = 256;
6180 global_work_size[0] = workItemCount;
6183 RandomInfo* randomInfo = AcquireRandomInfo();
6184 const unsigned long* s = GetRandomInfoSeed(randomInfo);
6186 GetPseudoRandomValue(randomInfo);
6188 randomInfo = DestroyRandomInfo(randomInfo);
6192 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&imageBuffer);
6193 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
6194 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&inputPixelCount);
6195 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&pixelsPerWorkitem);
6196 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
6197 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
6199 option=GetImageArtifact(image,"attenuate");
6200 if (option != (char *) NULL)
6201 attenuate=(float)StringToDouble(option,(char **) NULL);
6202 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
6203 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed0);
6204 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed1);
6205 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
6207 clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,NULL,0,NULL,NULL);
6209 if (ALIGNED(filteredPixels,CLPixelPacket))
6211 length = image->columns * image->rows;
6212 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
6216 length = image->columns * image->rows;
6217 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
6219 if (clStatus != CL_SUCCESS)
6221 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
6225 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
6228 OpenCLLogException(__FUNCTION__,__LINE__,exception);
6230 image_view=DestroyCacheView(image_view);
6231 if (filteredImage_view != NULL)
6232 filteredImage_view=DestroyCacheView(filteredImage_view);
6234 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
6235 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
6236 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
6237 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
6238 if (outputReady == MagickFalse && filteredImage != NULL)
6239 filteredImage=DestroyImage(filteredImage);
6241 return(filteredImage);
6245 MagickExport Image *AccelerateAddNoiseImage(const Image *image,
6246 const ChannelType channel,const NoiseType noise_type,
6247 ExceptionInfo *exception)
6252 assert(image != NULL);
6253 assert(exception != (ExceptionInfo *) NULL);
6255 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
6256 (checkAccelerateCondition(image, channel) == MagickFalse))
6259 filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
6261 return(filteredImage);
6264 static MagickBooleanType LaunchRandomImageKernel(MagickCLEnv clEnv,
6265 cl_command_queue queue,cl_mem imageBuffer,const unsigned int imageColumns,
6266 const unsigned int imageRows,cl_mem seedBuffer,
6267 const unsigned int numGenerators,ExceptionInfo *exception)
6285 status = MagickFalse;
6286 randomImageKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RandomImage");
6289 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&imageBuffer);
6290 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageColumns);
6291 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageRows);
6292 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&seedBuffer);
6294 const float randNormNumerator = 1.0f;
6295 const unsigned int randNormDenominator = (unsigned int)(~0UL);
6296 clEnv->library->clSetKernelArg(randomImageKernel,k++,
6297 sizeof(float),(void*)&randNormNumerator);
6298 clEnv->library->clSetKernelArg(randomImageKernel,k++,
6299 sizeof(cl_uint),(void*)&randNormDenominator);
6303 global_work_size = numGenerators;
6304 local_work_size = 64;
6306 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue,randomImageKernel,1,NULL,&global_work_size,
6307 &local_work_size,0,NULL,NULL);
6309 if (clStatus != CL_SUCCESS)
6311 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
6312 "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
6315 status = MagickTrue;
6318 if (randomImageKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomImageKernel);
6322 static MagickBooleanType ComputeRandomImage(Image* image,
6323 ExceptionInfo* exception)
6337 /* Don't release this buffer in this function !!! */
6339 randomNumberSeedsBuffer;
6360 status = MagickFalse;
6361 outputReady = MagickFalse;
6367 clEnv = GetDefaultOpenCLEnv();
6368 context = GetOpenCLContext(clEnv);
6370 /* Create and initialize OpenCL buffers. */
6371 image_view=AcquireAuthenticCacheView(image,exception);
6372 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
6373 if (inputPixels == (void *) NULL)
6375 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
6379 /* If the host pointer is aligned to the size of CLPixelPacket,
6380 then use the host buffer directly from the GPU; otherwise,
6381 create a buffer on the GPU and copy the data over */
6382 if (ALIGNED(inputPixels,CLPixelPacket))
6384 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
6388 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
6390 /* create a CL buffer from image pixel buffer */
6391 length = image->columns * image->rows;
6392 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6393 if (clStatus != CL_SUCCESS)
6395 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6399 queue = AcquireOpenCLCommandQueue(clEnv);
6401 randomNumberSeedsBuffer = GetAndLockRandSeedBuffer(clEnv);
6402 if (randomNumberSeedsBuffer==NULL)
6404 (void) OpenCLThrowMagickException(exception, GetMagickModule(),
6405 ResourceLimitWarning, "Failed to get GPU random number generators.",
6410 status = LaunchRandomImageKernel(clEnv,queue,
6412 (unsigned int) image->columns,
6413 (unsigned int) image->rows,
6414 randomNumberSeedsBuffer,
6415 GetNumRandGenerators(clEnv),
6417 if (status==MagickFalse)
6422 if (ALIGNED(inputPixels,CLPixelPacket))
6424 length = image->columns * image->rows;
6425 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
6429 length = image->columns * image->rows;
6430 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
6432 if (clStatus != CL_SUCCESS)
6434 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
6437 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
6440 OpenCLLogException(__FUNCTION__,__LINE__,exception);
6442 image_view=DestroyCacheView(image_view);
6444 UnlockRandSeedBuffer(clEnv);
6445 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
6446 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
6450 MagickExport MagickBooleanType AccelerateRandomImage(Image *image,
6451 ExceptionInfo* exception)
6456 assert(image != NULL);
6457 assert(exception != (ExceptionInfo *) NULL);
6459 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
6460 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
6461 return(MagickFalse);
6463 status=ComputeRandomImage(image,exception);
6467 static Image* ComputeMotionBlurImage(const Image *image,
6468 const ChannelType channel,const double *kernel,const size_t width,
6469 const OffsetInfo *offset,ExceptionInfo *exception)
6472 *filteredImage_view,
6491 filteredImageBuffer,
6524 global_work_size[2],
6537 outputReady = MagickFalse;
6539 filteredImage = NULL;
6540 filteredImage_view = NULL;
6542 filteredImageBuffer = NULL;
6543 imageKernelBuffer = NULL;
6544 motionBlurKernel = NULL;
6547 clEnv = GetDefaultOpenCLEnv();
6548 context = GetOpenCLContext(clEnv);
6550 /* Create and initialize OpenCL buffers. */
6552 image_view=AcquireVirtualCacheView(image,exception);
6553 inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
6554 if (inputPixels == (const void *) NULL)
6556 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
6557 "UnableToReadPixelCache.","`%s'",image->filename);
6561 // If the host pointer is aligned to the size of CLPixelPacket,
6562 // then use the host buffer directly from the GPU; otherwise,
6563 // create a buffer on the GPU and copy the data over
6564 if (ALIGNED(inputPixels,CLPixelPacket))
6566 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
6570 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
6572 // create a CL buffer from image pixel buffer
6573 length = image->columns * image->rows;
6574 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6575 length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6576 if (clStatus != CL_SUCCESS)
6578 (void) ThrowMagickException(exception, GetMagickModule(),
6579 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6584 filteredImage = CloneImage(image,image->columns,image->rows,
6585 MagickTrue,exception);
6586 assert(filteredImage != NULL);
6587 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
6589 (void) ThrowMagickException(exception, GetMagickModule(),
6590 ResourceLimitError, "CloneImage failed.", "'%s'", ".");
6593 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
6594 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
6595 if (filteredPixels == (void *) NULL)
6597 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
6598 "UnableToReadPixelCache.","`%s'",filteredImage->filename);
6602 if (ALIGNED(filteredPixels,CLPixelPacket))
6604 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
6605 hostPtr = filteredPixels;
6609 mem_flags = CL_MEM_WRITE_ONLY;
6612 // create a CL buffer from image pixel buffer
6613 length = image->columns * image->rows;
6614 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6615 length * sizeof(CLPixelPacket), hostPtr, &clStatus);
6616 if (clStatus != CL_SUCCESS)
6618 (void) ThrowMagickException(exception, GetMagickModule(),
6619 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6624 imageKernelBuffer = clEnv->library->clCreateBuffer(context,
6625 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(float), NULL,
6627 if (clStatus != CL_SUCCESS)
6629 (void) ThrowMagickException(exception, GetMagickModule(),
6630 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6634 queue = AcquireOpenCLCommandQueue(clEnv);
6635 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
6636 CL_TRUE, CL_MAP_WRITE, 0, width * sizeof(float), 0, NULL, NULL, &clStatus);
6637 if (clStatus != CL_SUCCESS)
6639 (void) ThrowMagickException(exception, GetMagickModule(),
6640 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
6643 for (i = 0; i < width; i++)
6645 kernelBufferPtr[i] = (float) kernel[i];
6647 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr,
6649 if (clStatus != CL_SUCCESS)
6651 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6652 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
6656 offsetBuffer = clEnv->library->clCreateBuffer(context,
6657 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(cl_int2), NULL,
6659 if (clStatus != CL_SUCCESS)
6661 (void) ThrowMagickException(exception, GetMagickModule(),
6662 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6666 offsetBufferPtr = (int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
6667 CL_MAP_WRITE, 0, width * sizeof(cl_int2), 0, NULL, NULL, &clStatus);
6668 if (clStatus != CL_SUCCESS)
6670 (void) ThrowMagickException(exception, GetMagickModule(),
6671 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
6674 for (i = 0; i < width; i++)
6676 offsetBufferPtr[2*i] = (int)offset[i].x;
6677 offsetBufferPtr[2*i+1] = (int)offset[i].y;
6679 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
6681 if (clStatus != CL_SUCCESS)
6683 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6684 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
6689 // get the OpenCL kernel
6690 motionBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
6692 if (motionBlurKernel == NULL)
6694 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6695 "AcquireOpenCLKernel failed.", "'%s'", ".");
6699 // set the kernel arguments
6701 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6702 (void *)&imageBuffer);
6703 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6704 (void *)&filteredImageBuffer);
6705 imageWidth = (unsigned int) image->columns;
6706 imageHeight = (unsigned int) image->rows;
6707 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
6709 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
6711 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6712 (void *)&imageKernelBuffer);
6713 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
6715 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6716 (void *)&offsetBuffer);
6718 GetPixelInfo(image,&bias);
6719 biasPixel.s[0] = bias.red;
6720 biasPixel.s[1] = bias.green;
6721 biasPixel.s[2] = bias.blue;
6722 biasPixel.s[3] = bias.alpha;
6723 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_float4), &biasPixel);
6725 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(ChannelType), &channel);
6726 matte = (image->alpha_trait != BlendPixelTrait)?1:0;
6727 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), &matte);
6728 if (clStatus != CL_SUCCESS)
6730 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6731 "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
6735 // launch the kernel
6736 local_work_size[0] = 16;
6737 local_work_size[1] = 16;
6738 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
6739 (unsigned int) image->columns,(unsigned int) local_work_size[0]);
6740 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
6741 (unsigned int) image->rows,(unsigned int) local_work_size[1]);
6742 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
6743 global_work_size, local_work_size, 0, NULL, NULL);
6745 if (clStatus != CL_SUCCESS)
6747 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6748 "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
6751 clEnv->library->clFlush(queue);
6753 if (ALIGNED(filteredPixels,CLPixelPacket))
6755 length = image->columns * image->rows;
6756 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE,
6757 CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL,
6762 length = image->columns * image->rows;
6763 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0,
6764 length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
6766 if (clStatus != CL_SUCCESS)
6768 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6769 "Reading output image from CL buffer failed.", "'%s'", ".");
6772 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
6776 image_view=DestroyCacheView(image_view);
6777 if (filteredImage_view != NULL)
6778 filteredImage_view=DestroyCacheView(filteredImage_view);
6780 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
6781 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
6782 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
6783 if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, motionBlurKernel);
6784 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
6785 if (outputReady == MagickFalse && filteredImage != NULL)
6786 filteredImage=DestroyImage(filteredImage);
6788 return(filteredImage);
6791 MagickExport Image *AccelerateMotionBlurImage(const Image *image,
6792 const ChannelType channel,const double* kernel,const size_t width,
6793 const OffsetInfo *offset,ExceptionInfo *exception)
6798 assert(image != NULL);
6799 assert(kernel != (double *) NULL);
6800 assert(offset != (OffsetInfo *) NULL);
6801 assert(exception != (ExceptionInfo *) NULL);
6803 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
6804 (checkAccelerateCondition(image, channel) == MagickFalse))
6807 filteredImage=ComputeMotionBlurImage(image, channel, kernel, width,
6809 return(filteredImage);
6812 static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv,
6813 cl_command_queue queue,cl_mem imageBuffer,const unsigned int inputWidth,
6814 const unsigned int inputHeight,const unsigned int matte,
6815 const ChannelType channel,const CompositeOperator compose,
6816 const cl_mem compositeImageBuffer,const unsigned int compositeWidth,
6817 const unsigned int compositeHeight,const float destination_dissolve,
6818 const float source_dissolve,ExceptionInfo *magick_unused(exception))
6830 global_work_size[2],
6836 magick_unreferenced(exception);
6838 compositeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
6842 clStatus=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&imageBuffer);
6843 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputWidth);
6844 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputHeight);
6845 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&compositeImageBuffer);
6846 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeWidth);
6847 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeHeight);
6848 composeOp = (unsigned int)compose;
6849 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&composeOp);
6850 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(ChannelType),(void*)&channel);
6851 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&matte);
6852 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&destination_dissolve);
6853 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&source_dissolve);
6855 if (clStatus!=CL_SUCCESS)
6858 local_work_size[0] = 64;
6859 local_work_size[1] = 1;
6861 global_work_size[0] = padGlobalWorkgroupSizeToLocalWorkgroupSize(inputWidth,
6862 (unsigned int) local_work_size[0]);
6863 global_work_size[1] = inputHeight;
6864 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL,
6865 global_work_size, local_work_size, 0, NULL, NULL);
6868 RelinquishOpenCLKernel(clEnv, compositeKernel);
6870 return((clStatus==CL_SUCCESS) ? MagickTrue : MagickFalse);
6873 static MagickBooleanType ComputeCompositeImage(Image *image,
6874 const ChannelType channel,const CompositeOperator compose,
6875 const Image *compositeImage,const ssize_t magick_unused(x_offset),
6876 const ssize_t magick_unused(y_offset),const float destination_dissolve,
6877 const float source_dissolve,ExceptionInfo *exception)
6895 compositeImageBuffer,
6914 magick_unreferenced(x_offset);
6915 magick_unreferenced(y_offset);
6917 status = MagickFalse;
6918 outputReady = MagickFalse;
6919 composePixels = NULL;
6921 compositeImageBuffer = NULL;
6923 clEnv = GetDefaultOpenCLEnv();
6924 context = GetOpenCLContext(clEnv);
6925 queue = AcquireOpenCLCommandQueue(clEnv);
6927 /* Create and initialize OpenCL buffers. */
6928 image_view=AcquireAuthenticCacheView(image,exception);
6929 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
6930 if (inputPixels == (void *) NULL)
6932 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,
6933 "UnableToReadPixelCache.","`%s'",image->filename);
6937 /* If the host pointer is aligned to the size of CLPixelPacket,
6938 then use the host buffer directly from the GPU; otherwise,
6939 create a buffer on the GPU and copy the data over */
6940 if (ALIGNED(inputPixels,CLPixelPacket))
6942 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
6946 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
6948 /* create a CL buffer from image pixel buffer */
6949 length = image->columns * image->rows;
6950 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6951 length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6952 if (clStatus != CL_SUCCESS)
6954 (void) OpenCLThrowMagickException(exception, GetMagickModule(),
6955 ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6960 /* Create and initialize OpenCL buffers. */
6961 composePixels = AcquirePixelCachePixels(compositeImage, &length, exception);
6962 if (composePixels == (void *) NULL)
6964 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,
6965 "UnableToReadPixelCache.","`%s'",compositeImage->filename);
6969 /* If the host pointer is aligned to the size of CLPixelPacket,
6970 then use the host buffer directly from the GPU; otherwise,
6971 create a buffer on the GPU and copy the data over */
6972 if (ALIGNED(composePixels,CLPixelPacket))
6974 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
6978 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
6980 /* create a CL buffer from image pixel buffer */
6981 length = compositeImage->columns * compositeImage->rows;
6982 compositeImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6983 length * sizeof(CLPixelPacket), (void*)composePixels, &clStatus);
6984 if (clStatus != CL_SUCCESS)
6986 (void) OpenCLThrowMagickException(exception, GetMagickModule(),
6987 ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6991 status = LaunchCompositeKernel(clEnv,queue,imageBuffer,
6992 (unsigned int) image->columns,
6993 (unsigned int) image->rows,
6994 (unsigned int) (image->alpha_trait != BlendPixelTrait) ? 1 : 0,
6995 channel, compose, compositeImageBuffer,
6996 (unsigned int) compositeImage->columns,
6997 (unsigned int) compositeImage->rows,
6998 destination_dissolve,source_dissolve,
7001 if (status==MagickFalse)
7004 length = image->columns * image->rows;
7005 if (ALIGNED(inputPixels,CLPixelPacket))
7007 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE,
7008 CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL,
7013 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0,
7014 length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
7016 if (clStatus==CL_SUCCESS)
7017 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
7021 image_view=DestroyCacheView(image_view);
7022 if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
7023 if (compositeImageBuffer!=NULL) clEnv->library->clReleaseMemObject(compositeImageBuffer);
7024 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
7026 return(outputReady);
7029 MagickExport MagickBooleanType AccelerateCompositeImage(Image *image,
7030 const ChannelType channel,const CompositeOperator compose,
7031 const Image *composite,const ssize_t x_offset,const ssize_t y_offset,
7032 const float destination_dissolve,const float source_dissolve,
7033 ExceptionInfo *exception)
7038 assert(image != NULL);
7039 assert(exception != (ExceptionInfo *) NULL);
7041 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
7042 (checkAccelerateCondition(image, channel) == MagickFalse))
7043 return(MagickFalse);
7045 /* only support zero offset and
7046 images with the size for now */
7049 || image->columns!=composite->columns
7050 || image->rows!=composite->rows)
7054 case ColorDodgeCompositeOp:
7055 case BlendCompositeOp:
7058 // unsupported compose operator, quit
7062 status = ComputeCompositeImage(image,channel,compose,composite,
7063 x_offset,y_offset,destination_dissolve,source_dissolve,exception);
7068 #else /* MAGICKCORE_OPENCL_SUPPORT */
7070 MagickExport Image *AccelerateConvolveImageChannel(
7071 const Image *magick_unused(image),const ChannelType magick_unused(channel),
7072 const KernelInfo *magick_unused(kernel),
7073 ExceptionInfo *magick_unused(exception))
7075 magick_unreferenced(image);
7076 magick_unreferenced(channel);
7077 magick_unreferenced(kernel);
7078 magick_unreferenced(exception);
7083 MagickExport MagickBooleanType AccelerateFunctionImage(
7084 Image *magick_unused(image),const ChannelType magick_unused(channel),
7085 const MagickFunction magick_unused(function),
7086 const size_t magick_unused(number_parameters),
7087 const double *magick_unused(parameters),
7088 ExceptionInfo *magick_unused(exception))
7090 magick_unreferenced(image);
7091 magick_unreferenced(channel);
7092 magick_unreferenced(function);
7093 magick_unreferenced(number_parameters);
7094 magick_unreferenced(parameters);
7095 magick_unreferenced(exception);
7100 MagickExport Image *AccelerateBlurImage(const Image *magick_unused(image),
7101 const ChannelType magick_unused(channel),const double magick_unused(radius),
7102 const double magick_unused(sigma),ExceptionInfo *magick_unused(exception))
7104 magick_unreferenced(image);
7105 magick_unreferenced(channel);
7106 magick_unreferenced(radius);
7107 magick_unreferenced(sigma);
7108 magick_unreferenced(exception);
7113 MagickExport Image *AccelerateRotationalBlurImage(
7114 const Image *magick_unused(image),const ChannelType magick_unused(channel),
7115 const double magick_unused(angle),ExceptionInfo *magick_unused(exception))
7117 magick_unreferenced(image);
7118 magick_unreferenced(channel);
7119 magick_unreferenced(angle);
7120 magick_unreferenced(exception);
7126 MagickExport Image *AccelerateUnsharpMaskImage(
7127 const Image *magick_unused(image),const ChannelType magick_unused(channel),
7128 const double magick_unused(radius),const double magick_unused(sigma),
7129 const double magick_unused(gain),const double magick_unused(threshold),
7130 ExceptionInfo *magick_unused(exception))
7132 magick_unreferenced(image);
7133 magick_unreferenced(channel);
7134 magick_unreferenced(radius);
7135 magick_unreferenced(sigma);
7136 magick_unreferenced(gain);
7137 magick_unreferenced(threshold);
7138 magick_unreferenced(exception);
7144 MagickBooleanType AccelerateCompositeImage(Image *image,
7145 const ChannelType channel,const CompositeOperator compose,
7146 const Image *composite,const ssize_t x_offset,const ssize_t y_offset,
7147 const float destination_dissolve,const float source_dissolve,
7148 ExceptionInfo *exception)
7150 magick_unreferenced(image);
7151 magick_unreferenced(channel);
7152 magick_unreferenced(compose);
7153 magick_unreferenced(composite);
7154 magick_unreferenced(x_offset);
7155 magick_unreferenced(y_offset);
7156 magick_unreferenced(destination_dissolve);
7157 magick_unreferenced(source_dissolve);
7158 magick_unreferenced(exception);
7164 MagickExport MagickBooleanType AccelerateContrastImage(
7165 Image* magick_unused(image),const MagickBooleanType magick_unused(sharpen),
7166 ExceptionInfo* magick_unused(exception))
7168 magick_unreferenced(image);
7169 magick_unreferenced(sharpen);
7170 magick_unreferenced(exception);
7175 MagickExport MagickBooleanType AccelerateContrastStretchImageChannel(
7176 Image * image, const ChannelType channel, const double black_point, const double white_point,
7177 ExceptionInfo* magick_unused(exception))
7179 magick_unreferenced(image);
7180 magick_unreferenced(channel);
7181 magick_unreferenced(black_point);
7182 magick_unreferenced(white_point);
7183 magick_unreferenced(exception);
7188 MagickExport MagickBooleanType AccelerateEqualizeImage(
7189 Image* magick_unused(image), const ChannelType magick_unused(channel),
7190 ExceptionInfo* magick_unused(exception))
7192 magick_unreferenced(image);
7193 magick_unreferenced(channel);
7194 magick_unreferenced(exception);
7199 MagickExport Image *AccelerateDespeckleImage(const Image* magick_unused(image),
7200 ExceptionInfo* magick_unused(exception))
7202 magick_unreferenced(image);
7203 magick_unreferenced(exception);
7208 MagickExport Image *AccelerateResizeImage(const Image* magick_unused(image),
7209 const size_t magick_unused(resizedColumns),
7210 const size_t magick_unused(resizedRows),
7211 const ResizeFilter* magick_unused(resizeFilter),
7212 ExceptionInfo *magick_unused(exception))
7214 magick_unreferenced(image);
7215 magick_unreferenced(resizedColumns);
7216 magick_unreferenced(resizedRows);
7217 magick_unreferenced(resizeFilter);
7218 magick_unreferenced(exception);
7224 MagickBooleanType AccelerateModulateImage(
7225 Image* image, double percent_brightness, double percent_hue,
7226 double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
7228 magick_unreferenced(image);
7229 magick_unreferenced(percent_brightness);
7230 magick_unreferenced(percent_hue);
7231 magick_unreferenced(percent_saturation);
7232 magick_unreferenced(colorspace);
7233 magick_unreferenced(exception);
7234 return(MagickFalse);
7238 MagickBooleanType AccelerateGrayscaleImage(
7239 Image* image, const PixelIntensityMethod method, ExceptionInfo* exception)
7241 magick_unreferenced(image);
7242 magick_unreferenced(method);
7243 magick_unreferenced(exception);
7244 return(MagickFalse);
7247 MagickExport Image *AccelerateAddNoiseImage(const Image *image,
7248 const ChannelType channel, const NoiseType noise_type,ExceptionInfo *exception)
7250 magick_unreferenced(image);
7251 magick_unreferenced(channel);
7252 magick_unreferenced(noise_type);
7253 magick_unreferenced(exception);
7258 MagickExport MagickBooleanType AccelerateRandomImage(Image* image, ExceptionInfo* exception)
7260 magick_unreferenced(image);
7261 magick_unreferenced(exception);
7266 Image* AccelerateMotionBlurImage(const Image *image, const ChannelType channel,
7267 const double* kernel, const size_t width,
7268 const OffsetInfo *offset,
7269 ExceptionInfo *exception)
7271 magick_unreferenced(image);
7272 magick_unreferenced(channel);
7273 magick_unreferenced(kernel);
7274 magick_unreferenced(width);
7275 magick_unreferenced(offset);
7276 magick_unreferenced(exception);
7280 #endif /* MAGICKCORE_OPENCL_SUPPORT */