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-2014 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-private.h"
66 #include "MagickCore/prepress.h"
67 #include "MagickCore/quantize.h"
68 #include "MagickCore/random_.h"
69 #include "MagickCore/random-private.h"
70 #include "MagickCore/registry.h"
71 #include "MagickCore/resize.h"
72 #include "MagickCore/resize-private.h"
73 #include "MagickCore/semaphore.h"
74 #include "MagickCore/splay-tree.h"
75 #include "MagickCore/statistic.h"
76 #include "MagickCore/string_.h"
77 #include "MagickCore/string-private.h"
78 #include "MagickCore/token.h"
80 #ifdef MAGICKCORE_CLPERFMARKER
81 #include "CLPerfMarker.h"
84 #define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
85 #define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
87 #if defined(MAGICKCORE_OPENCL_SUPPORT)
89 #define ALIGNED(pointer,type) ((((long)(pointer)) & (sizeof(type)-1)) == 0)
90 /*#define ALIGNED(pointer,type) (0) */
92 static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception)
94 MagickBooleanType flag;
97 clEnv = GetDefaultOpenCLEnv();
99 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
100 , sizeof(MagickBooleanType), &flag, exception);
101 if (flag != MagickFalse)
104 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED
105 , sizeof(MagickBooleanType), &flag, exception);
106 if (flag == MagickFalse)
108 if(InitOpenCLEnv(clEnv, exception) == MagickFalse)
111 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
112 , sizeof(MagickBooleanType), &flag, exception);
113 if (flag != MagickFalse)
121 static MagickBooleanType checkAccelerateCondition(const Image* image, const ChannelType channel)
123 /* check if the image's colorspace is supported */
124 if (image->colorspace != RGBColorspace
125 && image->colorspace != sRGBColorspace)
128 /* check if the channel is supported */
129 if (((channel&RedChannel) == 0)
130 || ((channel&GreenChannel) == 0)
131 || ((channel&BlueChannel) == 0))
137 /* check if if the virtual pixel method is compatible with the OpenCL implementation */
138 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod)&&
139 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
146 static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType channel, const KernelInfo *kernel, ExceptionInfo *exception)
148 MagickBooleanType outputReady;
152 size_t global_work_size[2];
153 size_t localGroupSize[2];
154 size_t localMemoryRequirement;
155 Image* filteredImage;
156 MagickSizeType length;
157 const void *inputPixels;
158 void *filteredPixels;
159 cl_mem_flags mem_flags;
160 float* kernelBufferPtr;
164 unsigned int matte, filterWidth, filterHeight, imageWidth, imageHeight;
168 cl_mem inputImageBuffer, filteredImageBuffer, convolutionKernel;
169 cl_ulong deviceLocalMemorySize;
172 cl_command_queue queue;
174 /* intialize all CL objects to NULL */
176 inputImageBuffer = NULL;
177 filteredImageBuffer = NULL;
178 convolutionKernel = NULL;
183 filteredImage = NULL;
184 outputReady = MagickFalse;
186 clEnv = GetDefaultOpenCLEnv();
187 context = GetOpenCLContext(clEnv);
190 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
191 if (inputPixels == (const void *) NULL)
193 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
197 /* Create and initialize OpenCL buffers. */
199 /* If the host pointer is aligned to the size of CLPixelPacket,
200 then use the host buffer directly from the GPU; otherwise,
201 create a buffer on the GPU and copy the data over */
202 if (ALIGNED(inputPixels,CLPixelPacket))
204 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
208 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
210 /* create a CL buffer from image pixel buffer */
211 length = inputImage->columns * inputImage->rows;
212 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
213 if (clStatus != CL_SUCCESS)
215 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
219 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
220 assert(filteredImage != NULL);
221 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
223 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
226 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
227 if (filteredPixels == (void *) NULL)
229 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
233 if (ALIGNED(filteredPixels,CLPixelPacket))
235 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
236 hostPtr = filteredPixels;
240 mem_flags = CL_MEM_WRITE_ONLY;
243 /* create a CL buffer from image pixel buffer */
244 length = inputImage->columns * inputImage->rows;
245 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
246 if (clStatus != CL_SUCCESS)
248 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
252 kernelSize = kernel->width * kernel->height;
253 convolutionKernel = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus);
254 if (clStatus != CL_SUCCESS)
256 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
260 queue = AcquireOpenCLCommandQueue(clEnv);
262 kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize * sizeof(float)
263 , 0, NULL, NULL, &clStatus);
264 if (clStatus != CL_SUCCESS)
266 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
269 for (i = 0; i < kernelSize; i++)
271 kernelBufferPtr[i] = (float) kernel->values[i];
273 clStatus = clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
274 if (clStatus != CL_SUCCESS)
276 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
281 /* Compute the local memory requirement for a 16x16 workgroup.
282 If it's larger than 16k, reduce the workgroup size to 8x8 */
283 localGroupSize[0] = 16;
284 localGroupSize[1] = 16;
285 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
286 + kernel->width*kernel->height*sizeof(float);
287 if (localMemoryRequirement > 16384)
291 localGroupSize[0] = 8;
292 localGroupSize[1] = 8;
294 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
295 + kernel->width*kernel->height*sizeof(float);
298 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE, sizeof(cl_device_id), &device, exception);
299 clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &deviceLocalMemorySize, NULL);
300 if (localMemoryRequirement <= deviceLocalMemorySize)
302 /* get the OpenCL kernel */
303 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized");
304 if (clkernel == NULL)
306 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
310 /* set the kernel arguments */
312 clStatus =clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
313 clStatus|=clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
314 imageWidth = inputImage->columns;
315 imageHeight = inputImage->rows;
316 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
317 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
318 clStatus|=clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
319 filterWidth = kernel->width;
320 filterHeight = kernel->height;
321 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
322 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
323 matte = (inputImage->alpha_trait == BlendPixelTrait)?1:0;
324 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
325 clStatus|=clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
326 clStatus|=clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*sizeof(CLPixelPacket),NULL);
327 clStatus|=clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL);
328 if (clStatus != CL_SUCCESS)
330 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
334 /* pad the global size to a multiple of the local work size dimension */
335 global_work_size[0] = ((inputImage->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ;
336 global_work_size[1] = ((inputImage->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1];
338 /* launch the kernel */
339 clStatus = clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, NULL);
340 if (clStatus != CL_SUCCESS)
342 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
348 /* get the OpenCL kernel */
349 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve");
350 if (clkernel == NULL)
352 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
356 /* set the kernel arguments */
358 clStatus =clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
359 clStatus|=clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
360 clStatus|=clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
361 filterWidth = kernel->width;
362 filterHeight = kernel->height;
363 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
364 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
365 matte = (inputImage->alpha_trait == BlendPixelTrait)?1:0;
366 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
367 clStatus|=clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
368 if (clStatus != CL_SUCCESS)
370 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
374 global_work_size[0] = inputImage->columns;
375 global_work_size[1] = inputImage->rows;
377 /* launch the kernel */
378 clStatus = clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
379 if (clStatus != CL_SUCCESS)
381 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
387 if (ALIGNED(filteredPixels,CLPixelPacket))
389 length = inputImage->columns * inputImage->rows;
390 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
394 length = inputImage->columns * inputImage->rows;
395 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
397 if (clStatus != CL_SUCCESS)
399 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
403 /* everything is fine! :) */
404 outputReady = MagickTrue;
407 OpenCLLogException(__FUNCTION__,__LINE__,exception);
409 if (inputImageBuffer != NULL)
410 clReleaseMemObject(inputImageBuffer);
412 if (filteredImageBuffer != NULL)
413 clReleaseMemObject(filteredImageBuffer);
415 if (convolutionKernel != NULL)
416 clReleaseMemObject(convolutionKernel);
418 if (clkernel != NULL)
419 RelinquishOpenCLKernel(clEnv, clkernel);
422 RelinquishOpenCLCommandQueue(clEnv, queue);
424 if (outputReady == MagickFalse)
426 if (filteredImage != NULL)
428 DestroyImage(filteredImage);
429 filteredImage = NULL;
433 return filteredImage;
437 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
441 % C o n v o l v e I m a g e w i t h O p e n C L %
445 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
447 % ConvolveImage() applies a custom convolution kernel to the image.
449 % The format of the ConvolveImage method is:
451 % Image *ConvolveImage(const Image *image,const size_t order,
452 % const double *kernel,ExceptionInfo *exception)
453 % Image *ConvolveImageChannel(const Image *image,const ChannelType channel,
454 % const size_t order,const double *kernel,ExceptionInfo *exception)
456 % A description of each parameter follows:
458 % o image: the image.
460 % o channel: the channel type.
462 % o kernel: kernel info.
464 % o exception: return any errors or warnings in this structure.
468 MagickExport Image* AccelerateConvolveImageChannel(const Image *image, const ChannelType channel, const KernelInfo *kernel, ExceptionInfo *exception)
470 MagickBooleanType status;
471 Image* filteredImage = NULL;
473 assert(image != NULL);
474 assert(kernel != (KernelInfo *) NULL);
475 assert(exception != (ExceptionInfo *) NULL);
477 status = checkOpenCLEnvironment(exception);
478 if (status == MagickFalse)
481 status = checkAccelerateCondition(image, channel);
482 if (status == MagickFalse)
485 filteredImage = ComputeConvolveImage(image, channel, kernel, exception);
486 return filteredImage;
489 static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType channel,const MagickFunction function,
490 const size_t number_parameters,const double *parameters, ExceptionInfo *exception)
492 MagickBooleanType status;
496 MagickSizeType length;
498 float* parametersBufferPtr;
503 cl_command_queue queue;
504 cl_mem_flags mem_flags;
506 cl_mem parametersBuffer;
507 size_t globalWorkSize[2];
511 status = MagickFalse;
517 parametersBuffer = NULL;
519 clEnv = GetDefaultOpenCLEnv();
520 context = GetOpenCLContext(clEnv);
522 pixels = GetPixelCachePixels(image, &length, exception);
523 if (pixels == (void *) NULL)
525 (void) OpenCLThrowMagickException(exception, GetMagickModule(), CacheWarning,
526 "GetPixelCachePixels failed.",
527 "'%s'", image->filename);
532 if (ALIGNED(pixels,CLPixelPacket))
534 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
538 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
540 /* create a CL buffer from image pixel buffer */
541 length = image->columns * image->rows;
542 imageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)pixels, &clStatus);
543 if (clStatus != CL_SUCCESS)
545 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
549 parametersBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus);
550 if (clStatus != CL_SUCCESS)
552 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
556 queue = AcquireOpenCLCommandQueue(clEnv);
558 parametersBufferPtr = (float*)clEnqueueMapBuffer(queue, parametersBuffer, CL_TRUE, CL_MAP_WRITE, 0, number_parameters * sizeof(float)
559 , 0, NULL, NULL, &clStatus);
560 if (clStatus != CL_SUCCESS)
562 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
565 for (i = 0; i < number_parameters; i++)
567 parametersBufferPtr[i] = (float)parameters[i];
569 clStatus = clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL);
570 if (clStatus != CL_SUCCESS)
572 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
577 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "FunctionImage");
578 if (clkernel == NULL)
580 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
584 /* set the kernel arguments */
586 clStatus =clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
587 clStatus|=clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
588 clStatus|=clSetKernelArg(clkernel,i++,sizeof(MagickFunction),(void *)&function);
589 clStatus|=clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&number_parameters);
590 clStatus|=clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)¶metersBuffer);
591 if (clStatus != CL_SUCCESS)
593 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
597 globalWorkSize[0] = image->columns;
598 globalWorkSize[1] = image->rows;
599 /* launch the kernel */
600 clStatus = clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, NULL);
601 if (clStatus != CL_SUCCESS)
603 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
609 if (ALIGNED(pixels,CLPixelPacket))
611 length = image->columns * image->rows;
612 clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
616 length = image->columns * image->rows;
617 clStatus = clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), pixels, 0, NULL, NULL);
619 if (clStatus != CL_SUCCESS)
621 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
627 OpenCLLogException(__FUNCTION__,__LINE__,exception);
629 if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel);
630 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
631 if (imageBuffer != NULL) clReleaseMemObject(imageBuffer);
632 if (parametersBuffer != NULL) clReleaseMemObject(parametersBuffer);
639 MagickExport MagickBooleanType
640 AccelerateFunctionImage(Image *image, const ChannelType channel,const MagickFunction function,
641 const size_t number_parameters,const double *parameters, ExceptionInfo *exception)
643 MagickBooleanType status;
645 status = MagickFalse;
647 assert(image != NULL);
648 assert(exception != (ExceptionInfo *) NULL);
650 status = checkOpenCLEnvironment(exception);
651 if (status != MagickFalse)
653 status = checkAccelerateCondition(image, channel);
654 if (status != MagickFalse)
656 status = ComputeFunctionImage(image, channel, function, number_parameters, parameters, exception);
663 static MagickBooleanType splitImage(const Image* inputImage)
665 MagickBooleanType split;
668 unsigned long allocSize;
669 unsigned long tempSize;
671 clEnv = GetDefaultOpenCLEnv();
673 allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
674 tempSize = inputImage->columns * inputImage->rows * 4 * 4;
677 printf("alloc size: %lu\n", allocSize);
678 printf("temp size: %lu\n", tempSize);
681 split = ((tempSize > allocSize) ? MagickTrue:MagickFalse);
686 static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channel, const double radius, const double sigma, ExceptionInfo *exception)
688 MagickBooleanType outputReady;
689 Image* filteredImage;
694 const void *inputPixels;
695 void *filteredPixels;
696 cl_mem_flags mem_flags;
699 cl_mem inputImageBuffer, tempImageBuffer, filteredImageBuffer, imageKernelBuffer;
700 cl_kernel blurRowKernel, blurColumnKernel;
701 cl_command_queue queue;
704 float* kernelBufferPtr;
705 MagickSizeType length;
707 char geometry[MaxTextExtent];
708 KernelInfo* kernel = NULL;
709 unsigned int kernelWidth;
710 unsigned int imageColumns, imageRows;
715 filteredImage = NULL;
716 inputImageBuffer = NULL;
717 tempImageBuffer = NULL;
718 filteredImageBuffer = NULL;
719 imageKernelBuffer = NULL;
720 blurRowKernel = NULL;
721 blurColumnKernel = NULL;
724 outputReady = MagickFalse;
726 clEnv = GetDefaultOpenCLEnv();
727 context = GetOpenCLContext(clEnv);
728 queue = AcquireOpenCLCommandQueue(clEnv);
730 /* Create and initialize OpenCL buffers. */
733 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
734 if (inputPixels == (const void *) NULL)
736 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
739 /* If the host pointer is aligned to the size of CLPixelPacket,
740 then use the host buffer directly from the GPU; otherwise,
741 create a buffer on the GPU and copy the data over */
742 if (ALIGNED(inputPixels,CLPixelPacket))
744 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
748 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
750 /* create a CL buffer from image pixel buffer */
751 length = inputImage->columns * inputImage->rows;
752 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
753 if (clStatus != CL_SUCCESS)
755 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
762 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
763 assert(filteredImage != NULL);
764 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
766 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
769 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
770 if (filteredPixels == (void *) NULL)
772 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
776 if (ALIGNED(filteredPixels,CLPixelPacket))
778 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
779 hostPtr = filteredPixels;
783 mem_flags = CL_MEM_WRITE_ONLY;
786 /* create a CL buffer from image pixel buffer */
787 length = inputImage->columns * inputImage->rows;
788 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
789 if (clStatus != CL_SUCCESS)
791 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
796 /* create processing kernel */
798 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
799 kernel=AcquireKernelInfo(geometry);
800 if (kernel == (KernelInfo *) NULL)
802 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
806 imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus);
807 if (clStatus != CL_SUCCESS)
809 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
812 kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
813 if (clStatus != CL_SUCCESS)
815 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
819 for (i = 0; i < kernel->width; i++)
821 kernelBufferPtr[i] = (float) kernel->values[i];
824 clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
825 if (clStatus != CL_SUCCESS)
827 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
834 /* create temp buffer */
836 length = inputImage->columns * inputImage->rows;
837 tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
838 if (clStatus != CL_SUCCESS)
840 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
845 /* get the OpenCL kernels */
847 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
848 if (blurRowKernel == NULL)
850 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
854 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn");
855 if (blurColumnKernel == NULL)
857 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
863 /* need logic to decide this value */
867 imageColumns = inputImage->columns;
868 imageRows = inputImage->rows;
870 /* set the kernel arguments */
872 clStatus=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
873 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
874 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
875 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
876 kernelWidth = kernel->width;
877 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
878 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
879 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
880 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
881 if (clStatus != CL_SUCCESS)
883 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
888 /* launch the kernel */
893 gsize[0] = chunkSize*((inputImage->columns+chunkSize-1)/chunkSize);
894 gsize[1] = inputImage->rows;
895 wsize[0] = chunkSize;
898 clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
899 if (clStatus != CL_SUCCESS)
901 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
909 /* need logic to decide this value */
913 imageColumns = inputImage->columns;
914 imageRows = inputImage->rows;
916 /* set the kernel arguments */
918 clStatus=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
919 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
920 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
921 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
922 kernelWidth = kernel->width;
923 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
924 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
925 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
926 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *)NULL);
927 if (clStatus != CL_SUCCESS)
929 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
934 /* launch the kernel */
939 gsize[0] = inputImage->columns;
940 gsize[1] = chunkSize*((inputImage->rows+chunkSize-1)/chunkSize);
942 wsize[1] = chunkSize;
944 clStatus = clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
945 if (clStatus != CL_SUCCESS)
947 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
957 if (ALIGNED(filteredPixels,CLPixelPacket))
959 length = inputImage->columns * inputImage->rows;
960 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
964 length = inputImage->columns * inputImage->rows;
965 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
967 if (clStatus != CL_SUCCESS)
969 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
973 outputReady = MagickTrue;
976 OpenCLLogException(__FUNCTION__,__LINE__,exception);
978 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
979 if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer);
980 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
981 if (imageKernelBuffer!=NULL) clReleaseMemObject(imageKernelBuffer);
982 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
983 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
984 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
985 if (kernel!=NULL) DestroyKernelInfo(kernel);
986 if (outputReady == MagickFalse)
988 if (filteredImage != NULL)
990 DestroyImage(filteredImage);
991 filteredImage = NULL;
994 return filteredImage;
997 static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType channel, const double radius, const double sigma, ExceptionInfo *exception)
999 MagickBooleanType outputReady;
1000 Image* filteredImage;
1005 const void *inputPixels;
1006 void *filteredPixels;
1007 cl_mem_flags mem_flags;
1010 cl_mem inputImageBuffer, tempImageBuffer, filteredImageBuffer, imageKernelBuffer;
1011 cl_kernel blurRowKernel, blurColumnKernel;
1012 cl_command_queue queue;
1015 float* kernelBufferPtr;
1016 MagickSizeType length;
1018 char geometry[MaxTextExtent];
1019 KernelInfo* kernel = NULL;
1020 unsigned int kernelWidth;
1021 unsigned int imageColumns, imageRows;
1026 filteredImage = NULL;
1027 inputImageBuffer = NULL;
1028 tempImageBuffer = NULL;
1029 filteredImageBuffer = NULL;
1030 imageKernelBuffer = NULL;
1031 blurRowKernel = NULL;
1032 blurColumnKernel = NULL;
1035 outputReady = MagickFalse;
1037 clEnv = GetDefaultOpenCLEnv();
1038 context = GetOpenCLContext(clEnv);
1039 queue = AcquireOpenCLCommandQueue(clEnv);
1041 /* Create and initialize OpenCL buffers. */
1044 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
1045 if (inputPixels == (const void *) NULL)
1047 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
1050 /* If the host pointer is aligned to the size of CLPixelPacket,
1051 then use the host buffer directly from the GPU; otherwise,
1052 create a buffer on the GPU and copy the data over */
1053 if (ALIGNED(inputPixels,CLPixelPacket))
1055 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1059 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1061 /* create a CL buffer from image pixel buffer */
1062 length = inputImage->columns * inputImage->rows;
1063 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1064 if (clStatus != CL_SUCCESS)
1066 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1073 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
1074 assert(filteredImage != NULL);
1075 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1077 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1080 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
1081 if (filteredPixels == (void *) NULL)
1083 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1087 if (ALIGNED(filteredPixels,CLPixelPacket))
1089 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1090 hostPtr = filteredPixels;
1094 mem_flags = CL_MEM_WRITE_ONLY;
1097 /* create a CL buffer from image pixel buffer */
1098 length = inputImage->columns * inputImage->rows;
1099 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1100 if (clStatus != CL_SUCCESS)
1102 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1107 /* create processing kernel */
1109 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
1110 kernel=AcquireKernelInfo(geometry);
1111 if (kernel == (KernelInfo *) NULL)
1113 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
1117 imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus);
1118 if (clStatus != CL_SUCCESS)
1120 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1123 kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
1124 if (clStatus != CL_SUCCESS)
1126 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
1130 for (i = 0; i < kernel->width; i++)
1132 kernelBufferPtr[i] = (float) kernel->values[i];
1135 clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
1136 if (clStatus != CL_SUCCESS)
1138 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
1144 unsigned int offsetRows;
1147 /* create temp buffer */
1149 length = inputImage->columns * (inputImage->rows / 2 + 1 + (kernel->width-1) / 2);
1150 tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
1151 if (clStatus != CL_SUCCESS)
1153 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1158 /* get the OpenCL kernels */
1160 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection");
1161 if (blurRowKernel == NULL)
1163 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1167 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumnSection");
1168 if (blurColumnKernel == NULL)
1170 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1175 for (sec = 0; sec < 2; sec++)
1178 /* need logic to decide this value */
1179 int chunkSize = 256;
1182 imageColumns = inputImage->columns;
1184 imageRows = inputImage->rows / 2 + (kernel->width-1) / 2;
1186 imageRows = (inputImage->rows - inputImage->rows / 2) + (kernel->width-1) / 2;
1188 offsetRows = sec * inputImage->rows / 2;
1190 kernelWidth = kernel->width;
1192 /* set the kernel arguments */
1194 clStatus=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1195 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1196 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
1197 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1198 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1199 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1200 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1201 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
1202 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
1203 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec);
1204 if (clStatus != CL_SUCCESS)
1206 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
1211 /* launch the kernel */
1216 gsize[0] = chunkSize*((imageColumns+chunkSize-1)/chunkSize);
1217 gsize[1] = imageRows;
1218 wsize[0] = chunkSize;
1221 clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1222 if (clStatus != CL_SUCCESS)
1224 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
1232 /* need logic to decide this value */
1233 int chunkSize = 256;
1236 imageColumns = inputImage->columns;
1238 imageRows = inputImage->rows / 2;
1240 imageRows = (inputImage->rows - inputImage->rows / 2);
1242 offsetRows = sec * inputImage->rows / 2;
1244 kernelWidth = kernel->width;
1246 /* set the kernel arguments */
1248 clStatus=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1249 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1250 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
1251 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1252 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1253 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1254 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1255 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *)NULL);
1256 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
1257 clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&sec);
1258 if (clStatus != CL_SUCCESS)
1260 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
1265 /* launch the kernel */
1270 gsize[0] = imageColumns;
1271 gsize[1] = chunkSize*((imageRows+chunkSize-1)/chunkSize);
1273 wsize[1] = chunkSize;
1275 clStatus = clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1276 if (clStatus != CL_SUCCESS)
1278 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
1289 if (ALIGNED(filteredPixels,CLPixelPacket))
1291 length = inputImage->columns * inputImage->rows;
1292 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1296 length = inputImage->columns * inputImage->rows;
1297 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1299 if (clStatus != CL_SUCCESS)
1301 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1305 outputReady = MagickTrue;
1308 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1310 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
1311 if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer);
1312 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
1313 if (imageKernelBuffer!=NULL) clReleaseMemObject(imageKernelBuffer);
1314 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
1315 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
1316 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1317 if (kernel!=NULL) DestroyKernelInfo(kernel);
1318 if (outputReady == MagickFalse)
1320 if (filteredImage != NULL)
1322 DestroyImage(filteredImage);
1323 filteredImage = NULL;
1326 return filteredImage;
1330 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1334 % B l u r I m a g e w i t h O p e n C L %
1338 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1340 % BlurImage() blurs an image. We convolve the image with a Gaussian operator
1341 % of the given radius and standard deviation (sigma). For reasonable results,
1342 % the radius should be larger than sigma. Use a radius of 0 and BlurImage()
1343 % selects a suitable radius for you.
1345 % The format of the BlurImage method is:
1347 % Image *BlurImage(const Image *image,const double radius,
1348 % const double sigma,ExceptionInfo *exception)
1349 % Image *BlurImageChannel(const Image *image,const ChannelType channel,
1350 % const double radius,const double sigma,ExceptionInfo *exception)
1352 % A description of each parameter follows:
1354 % o image: the image.
1356 % o channel: the channel type.
1358 % o radius: the radius of the Gaussian, in pixels, not counting the center
1361 % o sigma: the standard deviation of the Gaussian, in pixels.
1363 % o exception: return any errors or warnings in this structure.
1368 Image* AccelerateBlurImage(const Image *image, const ChannelType channel, const double radius, const double sigma,ExceptionInfo *exception)
1370 MagickBooleanType status;
1371 Image* filteredImage = NULL;
1373 assert(image != NULL);
1374 assert(exception != (ExceptionInfo *) NULL);
1376 status = checkOpenCLEnvironment(exception);
1377 if (status == MagickFalse)
1380 status = checkAccelerateCondition(image, channel);
1381 if (status == MagickFalse)
1384 if (splitImage(image) && (image->rows / 2 > radius))
1385 filteredImage = ComputeBlurImageSection(image, channel, radius, sigma, exception);
1387 filteredImage = ComputeBlurImage(image, channel, radius, sigma, exception);
1389 return filteredImage;
1393 static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType channel, const double angle, ExceptionInfo *exception)
1396 MagickBooleanType outputReady;
1397 Image* filteredImage;
1401 size_t global_work_size[2];
1404 cl_mem_flags mem_flags;
1405 cl_mem inputImageBuffer, filteredImageBuffer, sinThetaBuffer, cosThetaBuffer;
1406 cl_kernel radialBlurKernel;
1407 cl_command_queue queue;
1409 const void *inputPixels;
1410 void *filteredPixels;
1414 MagickSizeType length;
1417 cl_float4 biasPixel;
1418 cl_float2 blurCenter;
1420 unsigned int cossin_theta_size;
1421 float offset, theta;
1425 outputReady = MagickFalse;
1427 filteredImage = NULL;
1428 inputImageBuffer = NULL;
1429 filteredImageBuffer = NULL;
1430 sinThetaBuffer = NULL;
1431 cosThetaBuffer = NULL;
1433 radialBlurKernel = NULL;
1436 clEnv = GetDefaultOpenCLEnv();
1437 context = GetOpenCLContext(clEnv);
1440 /* Create and initialize OpenCL buffers. */
1443 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
1444 if (inputPixels == (const void *) NULL)
1446 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
1450 /* If the host pointer is aligned to the size of CLPixelPacket,
1451 then use the host buffer directly from the GPU; otherwise,
1452 create a buffer on the GPU and copy the data over */
1453 if (ALIGNED(inputPixels,CLPixelPacket))
1455 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1459 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1461 /* create a CL buffer from image pixel buffer */
1462 length = inputImage->columns * inputImage->rows;
1463 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1464 if (clStatus != CL_SUCCESS)
1466 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1471 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
1472 assert(filteredImage != NULL);
1473 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1475 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1478 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
1479 if (filteredPixels == (void *) NULL)
1481 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1485 if (ALIGNED(filteredPixels,CLPixelPacket))
1487 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1488 hostPtr = filteredPixels;
1492 mem_flags = CL_MEM_WRITE_ONLY;
1495 /* create a CL buffer from image pixel buffer */
1496 length = inputImage->columns * inputImage->rows;
1497 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1498 if (clStatus != CL_SUCCESS)
1500 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1504 blurCenter.s[0] = (float) (inputImage->columns-1)/2.0;
1505 blurCenter.s[1] = (float) (inputImage->rows-1)/2.0;
1506 blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]);
1507 cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL);
1509 /* create a buffer for sin_theta and cos_theta */
1510 sinThetaBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
1511 if (clStatus != CL_SUCCESS)
1513 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1516 cosThetaBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
1517 if (clStatus != CL_SUCCESS)
1519 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1524 queue = AcquireOpenCLCommandQueue(clEnv);
1525 sinThetaPtr = (float*) clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
1526 if (clStatus != CL_SUCCESS)
1528 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
1532 cosThetaPtr = (float*) clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
1533 if (clStatus != CL_SUCCESS)
1535 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
1539 theta=DegreesToRadians(angle)/(MagickRealType) (cossin_theta_size-1);
1540 offset=theta*(MagickRealType) (cossin_theta_size-1)/2.0;
1541 for (i=0; i < (ssize_t) cossin_theta_size; i++)
1543 cosThetaPtr[i]=(float)cos((double) (theta*i-offset));
1544 sinThetaPtr[i]=(float)sin((double) (theta*i-offset));
1547 clStatus = clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL);
1548 clStatus |= clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL);
1549 if (clStatus != CL_SUCCESS)
1551 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
1555 /* get the OpenCL kernel */
1556 radialBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RadialBlur");
1557 if (radialBlurKernel == NULL)
1559 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1564 /* set the kernel arguments */
1566 clStatus=clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1567 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1569 GetPixelInfo(inputImage,&bias);
1570 biasPixel.s[0] = bias.red;
1571 biasPixel.s[1] = bias.green;
1572 biasPixel.s[2] = bias.blue;
1573 biasPixel.s[3] = bias.alpha;
1574 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(cl_float4), &biasPixel);
1575 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(ChannelType), &channel);
1577 matte = (inputImage->alpha_trait == BlendPixelTrait)?1:0;
1578 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &matte);
1580 clStatus=clSetKernelArg(radialBlurKernel,i++,sizeof(cl_float2), &blurCenter);
1582 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer);
1583 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer);
1584 clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size);
1585 if (clStatus != CL_SUCCESS)
1587 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
1592 global_work_size[0] = inputImage->columns;
1593 global_work_size[1] = inputImage->rows;
1594 /* launch the kernel */
1595 clStatus = clEnqueueNDRangeKernel(queue, radialBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
1596 if (clStatus != CL_SUCCESS)
1598 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
1603 if (ALIGNED(filteredPixels,CLPixelPacket))
1605 length = inputImage->columns * inputImage->rows;
1606 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1610 length = inputImage->columns * inputImage->rows;
1611 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1613 if (clStatus != CL_SUCCESS)
1615 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1618 outputReady = MagickTrue;
1621 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1623 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
1624 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
1625 if (sinThetaBuffer!=NULL) clReleaseMemObject(sinThetaBuffer);
1626 if (cosThetaBuffer!=NULL) clReleaseMemObject(cosThetaBuffer);
1627 if (radialBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, radialBlurKernel);
1628 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1629 if (outputReady == MagickFalse)
1631 if (filteredImage != NULL)
1633 DestroyImage(filteredImage);
1634 filteredImage = NULL;
1637 return filteredImage;
1641 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1645 % R a d i a l B l u r I m a g e w i t h O p e n C L %
1649 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1651 % RadialBlurImage() applies a radial blur to the image.
1653 % Andrew Protano contributed this effect.
1655 % The format of the RadialBlurImage method is:
1657 % Image *RadialBlurImage(const Image *image,const double angle,
1658 % ExceptionInfo *exception)
1659 % Image *RadialBlurImageChannel(const Image *image,const ChannelType channel,
1660 % const double angle,ExceptionInfo *exception)
1662 % A description of each parameter follows:
1664 % o image: the image.
1666 % o channel: the channel type.
1668 % o angle: the angle of the radial blur.
1670 % o exception: return any errors or warnings in this structure.
1675 Image* AccelerateRadialBlurImage(const Image *image, const ChannelType channel, const double angle, ExceptionInfo *exception)
1677 MagickBooleanType status;
1678 Image* filteredImage;
1681 assert(image != NULL);
1682 assert(exception != NULL);
1684 status = checkOpenCLEnvironment(exception);
1685 if (status == MagickFalse)
1688 status = checkAccelerateCondition(image, channel);
1689 if (status == MagickFalse)
1692 filteredImage = ComputeRadialBlurImage(image, channel, angle, exception);
1693 return filteredImage;
1698 static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType channel,const double radius,const double sigma,
1699 const double gain,const double threshold,ExceptionInfo *exception)
1701 MagickBooleanType outputReady = MagickFalse;
1702 Image* filteredImage = NULL;
1703 MagickCLEnv clEnv = NULL;
1707 const void *inputPixels;
1708 void *filteredPixels;
1709 cl_mem_flags mem_flags;
1711 KernelInfo *kernel = NULL;
1712 char geometry[MaxTextExtent];
1714 cl_context context = NULL;
1715 cl_mem inputImageBuffer = NULL;
1716 cl_mem filteredImageBuffer = NULL;
1717 cl_mem tempImageBuffer = NULL;
1718 cl_mem imageKernelBuffer = NULL;
1719 cl_kernel blurRowKernel = NULL;
1720 cl_kernel unsharpMaskBlurColumnKernel = NULL;
1721 cl_command_queue queue = NULL;
1724 float* kernelBufferPtr;
1725 MagickSizeType length;
1726 unsigned int kernelWidth;
1729 unsigned int imageColumns, imageRows;
1733 clEnv = GetDefaultOpenCLEnv();
1734 context = GetOpenCLContext(clEnv);
1735 queue = AcquireOpenCLCommandQueue(clEnv);
1737 /* Create and initialize OpenCL buffers. */
1740 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
1741 if (inputPixels == (const void *) NULL)
1743 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
1747 /* If the host pointer is aligned to the size of CLPixelPacket,
1748 then use the host buffer directly from the GPU; otherwise,
1749 create a buffer on the GPU and copy the data over */
1750 if (ALIGNED(inputPixels,CLPixelPacket))
1752 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1756 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1758 /* create a CL buffer from image pixel buffer */
1759 length = inputImage->columns * inputImage->rows;
1760 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1761 if (clStatus != CL_SUCCESS)
1763 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1770 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
1771 assert(filteredImage != NULL);
1772 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1774 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1777 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
1778 if (filteredPixels == (void *) NULL)
1780 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1784 if (ALIGNED(filteredPixels,CLPixelPacket))
1786 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1787 hostPtr = filteredPixels;
1791 mem_flags = CL_MEM_WRITE_ONLY;
1795 /* create a CL buffer from image pixel buffer */
1796 length = inputImage->columns * inputImage->rows;
1797 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1798 if (clStatus != CL_SUCCESS)
1800 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1805 /* create the blur kernel */
1807 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
1808 kernel=AcquireKernelInfo(geometry);
1809 if (kernel == (KernelInfo *) NULL)
1811 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
1815 imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
1816 if (clStatus != CL_SUCCESS)
1818 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1823 kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
1824 if (clStatus != CL_SUCCESS)
1826 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
1829 for (i = 0; i < kernel->width; i++)
1831 kernelBufferPtr[i] = (float) kernel->values[i];
1833 clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
1834 if (clStatus != CL_SUCCESS)
1836 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
1842 /* create temp buffer */
1844 length = inputImage->columns * inputImage->rows;
1845 tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
1846 if (clStatus != CL_SUCCESS)
1848 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
1853 /* get the opencl kernel */
1855 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
1856 if (blurRowKernel == NULL)
1858 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1862 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn");
1863 if (unsharpMaskBlurColumnKernel == NULL)
1865 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1873 imageColumns = inputImage->columns;
1874 imageRows = inputImage->rows;
1876 kernelWidth = kernel->width;
1878 /* set the kernel arguments */
1880 clStatus=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1881 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1882 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
1883 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1884 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1885 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1886 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1887 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
1888 if (clStatus != CL_SUCCESS)
1890 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
1895 /* launch the kernel */
1900 gsize[0] = chunkSize*((inputImage->columns+chunkSize-1)/chunkSize);
1901 gsize[1] = inputImage->rows;
1902 wsize[0] = chunkSize;
1905 clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1906 if (clStatus != CL_SUCCESS)
1908 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
1917 imageColumns = inputImage->columns;
1918 imageRows = inputImage->rows;
1919 kernelWidth = kernel->width;
1920 fGain = (float)gain;
1921 fThreshold = (float)threshold;
1924 clStatus=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1925 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1926 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1927 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1928 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1929 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
1930 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
1931 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
1932 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1933 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1934 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
1935 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
1937 if (clStatus != CL_SUCCESS)
1939 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
1944 /* launch the kernel */
1949 gsize[0] = inputImage->columns;
1950 gsize[1] = chunkSize*((inputImage->rows+chunkSize-1)/chunkSize);
1952 wsize[1] = chunkSize;
1954 clStatus = clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1955 if (clStatus != CL_SUCCESS)
1957 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
1966 if (ALIGNED(filteredPixels,CLPixelPacket))
1968 length = inputImage->columns * inputImage->rows;
1969 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1973 length = inputImage->columns * inputImage->rows;
1974 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1976 if (clStatus != CL_SUCCESS)
1978 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1982 outputReady = MagickTrue;
1985 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1987 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
1988 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
1989 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
1990 if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer);
1991 if (imageKernelBuffer!=NULL) clReleaseMemObject(imageKernelBuffer);
1992 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
1993 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
1994 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1995 if (outputReady == MagickFalse)
1997 if (filteredImage != NULL)
1999 DestroyImage(filteredImage);
2000 filteredImage = NULL;
2003 return filteredImage;
2007 static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const ChannelType channel,const double radius,const double sigma,
2008 const double gain,const double threshold,ExceptionInfo *exception)
2010 MagickBooleanType outputReady = MagickFalse;
2011 Image* filteredImage = NULL;
2012 MagickCLEnv clEnv = NULL;
2016 const void *inputPixels;
2017 void *filteredPixels;
2018 cl_mem_flags mem_flags;
2020 KernelInfo *kernel = NULL;
2021 char geometry[MaxTextExtent];
2023 cl_context context = NULL;
2024 cl_mem inputImageBuffer = NULL;
2025 cl_mem filteredImageBuffer = NULL;
2026 cl_mem tempImageBuffer = NULL;
2027 cl_mem imageKernelBuffer = NULL;
2028 cl_kernel blurRowKernel = NULL;
2029 cl_kernel unsharpMaskBlurColumnKernel = NULL;
2030 cl_command_queue queue = NULL;
2033 float* kernelBufferPtr;
2034 MagickSizeType length;
2035 unsigned int kernelWidth;
2038 unsigned int imageColumns, imageRows;
2042 clEnv = GetDefaultOpenCLEnv();
2043 context = GetOpenCLContext(clEnv);
2044 queue = AcquireOpenCLCommandQueue(clEnv);
2046 /* Create and initialize OpenCL buffers. */
2049 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
2050 if (inputPixels == (const void *) NULL)
2052 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
2056 /* If the host pointer is aligned to the size of CLPixelPacket,
2057 then use the host buffer directly from the GPU; otherwise,
2058 create a buffer on the GPU and copy the data over */
2059 if (ALIGNED(inputPixels,CLPixelPacket))
2061 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2065 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2067 /* create a CL buffer from image pixel buffer */
2068 length = inputImage->columns * inputImage->rows;
2069 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2070 if (clStatus != CL_SUCCESS)
2072 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2079 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
2080 assert(filteredImage != NULL);
2081 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
2083 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2086 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
2087 if (filteredPixels == (void *) NULL)
2089 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2093 if (ALIGNED(filteredPixels,CLPixelPacket))
2095 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2096 hostPtr = filteredPixels;
2100 mem_flags = CL_MEM_WRITE_ONLY;
2104 /* create a CL buffer from image pixel buffer */
2105 length = inputImage->columns * inputImage->rows;
2106 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2107 if (clStatus != CL_SUCCESS)
2109 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2114 /* create the blur kernel */
2116 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
2117 kernel=AcquireKernelInfo(geometry);
2118 if (kernel == (KernelInfo *) NULL)
2120 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
2124 imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
2125 if (clStatus != CL_SUCCESS)
2127 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2132 kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
2133 if (clStatus != CL_SUCCESS)
2135 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
2138 for (i = 0; i < kernel->width; i++)
2140 kernelBufferPtr[i] = (float) kernel->values[i];
2142 clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
2143 if (clStatus != CL_SUCCESS)
2145 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
2151 unsigned int offsetRows;
2154 /* create temp buffer */
2156 length = inputImage->columns * (inputImage->rows / 2 + 1 + (kernel->width-1) / 2);
2157 tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
2158 if (clStatus != CL_SUCCESS)
2160 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2165 /* get the opencl kernel */
2167 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection");
2168 if (blurRowKernel == NULL)
2170 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2174 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumnSection");
2175 if (unsharpMaskBlurColumnKernel == NULL)
2177 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2182 for (sec = 0; sec < 2; sec++)
2187 imageColumns = inputImage->columns;
2189 imageRows = inputImage->rows / 2 + (kernel->width-1) / 2;
2191 imageRows = (inputImage->rows - inputImage->rows / 2) + (kernel->width-1) / 2;
2193 offsetRows = sec * inputImage->rows / 2;
2195 kernelWidth = kernel->width;
2197 /* set the kernel arguments */
2199 clStatus=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
2200 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2201 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
2202 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2203 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2204 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2205 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2206 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
2207 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
2208 clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec);
2209 if (clStatus != CL_SUCCESS)
2211 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
2215 /* launch the kernel */
2220 gsize[0] = chunkSize*((imageColumns+chunkSize-1)/chunkSize);
2221 gsize[1] = imageRows;
2222 wsize[0] = chunkSize;
2225 clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2226 if (clStatus != CL_SUCCESS)
2228 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
2238 imageColumns = inputImage->columns;
2240 imageRows = inputImage->rows / 2;
2242 imageRows = (inputImage->rows - inputImage->rows / 2);
2244 offsetRows = sec * inputImage->rows / 2;
2246 kernelWidth = kernel->width;
2248 fGain = (float)gain;
2249 fThreshold = (float)threshold;
2252 clStatus=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
2253 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2254 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
2255 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2256 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2257 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
2258 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
2259 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
2260 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2261 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2262 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
2263 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
2264 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
2265 clStatus|=clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&sec);
2267 if (clStatus != CL_SUCCESS)
2269 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
2274 /* launch the kernel */
2279 gsize[0] = imageColumns;
2280 gsize[1] = chunkSize*((imageRows+chunkSize-1)/chunkSize);
2282 wsize[1] = chunkSize;
2284 clStatus = clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2285 if (clStatus != CL_SUCCESS)
2287 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
2296 if (ALIGNED(filteredPixels,CLPixelPacket))
2298 length = inputImage->columns * inputImage->rows;
2299 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2303 length = inputImage->columns * inputImage->rows;
2304 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2306 if (clStatus != CL_SUCCESS)
2308 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2312 outputReady = MagickTrue;
2315 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2317 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
2318 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
2319 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
2320 if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer);
2321 if (imageKernelBuffer!=NULL) clReleaseMemObject(imageKernelBuffer);
2322 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
2323 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
2324 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2325 if (outputReady == MagickFalse)
2327 if (filteredImage != NULL)
2329 DestroyImage(filteredImage);
2330 filteredImage = NULL;
2333 return filteredImage;
2338 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2342 % 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 %
2346 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2348 % UnsharpMaskImage() sharpens one or more image channels. We convolve the
2349 % image with a Gaussian operator of the given radius and standard deviation
2350 % (sigma). For reasonable results, radius should be larger than sigma. Use a
2351 % radius of 0 and UnsharpMaskImage() selects a suitable radius for you.
2353 % The format of the UnsharpMaskImage method is:
2355 % Image *UnsharpMaskImage(const Image *image,const double radius,
2356 % const double sigma,const double amount,const double threshold,
2357 % ExceptionInfo *exception)
2358 % Image *UnsharpMaskImageChannel(const Image *image,
2359 % const ChannelType channel,const double radius,const double sigma,
2360 % const double gain,const double threshold,ExceptionInfo *exception)
2362 % A description of each parameter follows:
2364 % o image: the image.
2366 % o channel: the channel type.
2368 % o radius: the radius of the Gaussian, in pixels, not counting the center
2371 % o sigma: the standard deviation of the Gaussian, in pixels.
2373 % o gain: the percentage of the difference between the original and the
2374 % blur image that is added back into the original.
2376 % o threshold: the threshold in pixels needed to apply the diffence gain.
2378 % o exception: return any errors or warnings in this structure.
2384 Image* AccelerateUnsharpMaskImage(const Image *image, const ChannelType channel,const double radius,const double sigma,
2385 const double gain,const double threshold,ExceptionInfo *exception)
2387 MagickBooleanType status;
2388 Image* filteredImage;
2391 assert(image != NULL);
2392 assert(exception != NULL);
2394 status = checkOpenCLEnvironment(exception);
2395 if (status == MagickFalse)
2398 status = checkAccelerateCondition(image, channel);
2399 if (status == MagickFalse)
2402 if (splitImage(image) && (image->rows / 2 > radius))
2403 filteredImage = ComputeUnsharpMaskImageSection(image,channel,radius,sigma,gain,threshold,exception);
2405 filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
2406 return filteredImage;
2410 static MagickBooleanType resizeHorizontalFilter(cl_mem inputImage
2411 , const unsigned int inputImageColumns, const unsigned int inputImageRows, const unsigned int matte
2412 , cl_mem resizedImage, const unsigned int resizedColumns, const unsigned int resizedRows
2413 , const ResizeFilter* resizeFilter, cl_mem resizeFilterCubicCoefficients, const float xFactor
2414 , MagickCLEnv clEnv, cl_command_queue queue, ExceptionInfo *exception)
2416 MagickBooleanType status = MagickFalse;
2418 float scale, support;
2420 cl_kernel horizontalKernel = NULL;
2422 size_t global_work_size[2];
2423 size_t local_work_size[2];
2424 int resizeFilterType, resizeWindowType;
2425 float resizeFilterScale, resizeFilterSupport, resizeFilterWindowSupport, resizeFilterBlur;
2426 size_t totalLocalMemorySize;
2427 size_t imageCacheLocalMemorySize, pixelAccumulatorLocalMemorySize
2428 , weightAccumulatorLocalMemorySize, gammaAccumulatorLocalMemorySize;
2429 size_t deviceLocalMemorySize;
2430 int cacheRangeStart, cacheRangeEnd, numCachedPixels;
2432 const unsigned int workgroupSize = 256;
2433 unsigned int pixelPerWorkgroup;
2434 unsigned int chunkSize;
2437 Apply filter to resize vertically from image to resize image.
2439 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
2440 support=scale*GetResizeFilterSupport(resizeFilter);
2444 Support too small even for nearest neighbour: Reduce to point
2447 support=(MagickRealType) 0.5;
2450 scale=PerceptibleReciprocal(scale);
2452 if (resizedColumns < workgroupSize)
2455 pixelPerWorkgroup = 32;
2459 chunkSize = workgroupSize;
2460 pixelPerWorkgroup = workgroupSize;
2463 /* get the local memory size supported by the device */
2464 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
2466 DisableMSCWarning(4127)
2470 /* calculate the local memory size needed per workgroup */
2471 cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
2472 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
2473 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
2474 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
2475 totalLocalMemorySize = imageCacheLocalMemorySize;
2477 /* local size for the pixel accumulator */
2478 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
2479 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
2481 /* local memory size for the weight accumulator */
2482 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2483 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
2485 /* local memory size for the gamma accumulator */
2487 gammaAccumulatorLocalMemorySize = sizeof(float);
2489 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2490 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
2492 if (totalLocalMemorySize <= deviceLocalMemorySize)
2496 pixelPerWorkgroup = pixelPerWorkgroup/2;
2497 chunkSize = chunkSize/2;
2498 if (pixelPerWorkgroup == 0
2501 /* quit, fallback to CPU */
2507 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
2508 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
2511 if (resizeFilterType == SincFastWeightingFunction
2512 && resizeWindowType == SincFastWeightingFunction)
2514 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilterSinc");
2518 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter");
2520 if (horizontalKernel == NULL)
2522 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2527 clStatus = clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&inputImage);
2528 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageColumns);
2529 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageRows);
2530 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
2531 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
2532 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
2534 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
2535 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
2537 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
2538 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
2539 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
2541 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
2542 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
2544 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
2545 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
2547 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
2548 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
2550 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
2551 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
2554 clStatus |= clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
2555 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
2556 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
2557 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
2560 clStatus |= clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
2561 clStatus |= clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
2562 clStatus |= clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
2564 if (clStatus != CL_SUCCESS)
2566 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
2570 global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
2571 global_work_size[1] = resizedRows;
2573 local_work_size[0] = workgroupSize;
2574 local_work_size[1] = 1;
2575 clStatus = clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
2576 if (clStatus != CL_SUCCESS)
2578 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
2582 status = MagickTrue;
2586 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2588 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
2594 static MagickBooleanType resizeVerticalFilter(cl_mem inputImage
2595 , const unsigned int inputImageColumns, const unsigned int inputImageRows, const unsigned int matte
2596 , cl_mem resizedImage, const unsigned int resizedColumns, const unsigned int resizedRows
2597 , const ResizeFilter* resizeFilter, cl_mem resizeFilterCubicCoefficients, const float yFactor
2598 , MagickCLEnv clEnv, cl_command_queue queue, ExceptionInfo *exception)
2600 MagickBooleanType status = MagickFalse;
2602 float scale, support;
2604 cl_kernel horizontalKernel = NULL;
2606 size_t global_work_size[2];
2607 size_t local_work_size[2];
2608 int resizeFilterType, resizeWindowType;
2609 float resizeFilterScale, resizeFilterSupport, resizeFilterWindowSupport, resizeFilterBlur;
2610 size_t totalLocalMemorySize;
2611 size_t imageCacheLocalMemorySize, pixelAccumulatorLocalMemorySize
2612 , weightAccumulatorLocalMemorySize, gammaAccumulatorLocalMemorySize;
2613 size_t deviceLocalMemorySize;
2614 int cacheRangeStart, cacheRangeEnd, numCachedPixels;
2616 const unsigned int workgroupSize = 256;
2617 unsigned int pixelPerWorkgroup;
2618 unsigned int chunkSize;
2621 Apply filter to resize vertically from image to resize image.
2623 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
2624 support=scale*GetResizeFilterSupport(resizeFilter);
2628 Support too small even for nearest neighbour: Reduce to point
2631 support=(MagickRealType) 0.5;
2634 scale=PerceptibleReciprocal(scale);
2636 if (resizedRows < workgroupSize)
2639 pixelPerWorkgroup = 32;
2643 chunkSize = workgroupSize;
2644 pixelPerWorkgroup = workgroupSize;
2647 /* get the local memory size supported by the device */
2648 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
2650 DisableMSCWarning(4127)
2654 /* calculate the local memory size needed per workgroup */
2655 cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
2656 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
2657 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
2658 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
2659 totalLocalMemorySize = imageCacheLocalMemorySize;
2661 /* local size for the pixel accumulator */
2662 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
2663 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
2665 /* local memory size for the weight accumulator */
2666 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2667 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
2669 /* local memory size for the gamma accumulator */
2671 gammaAccumulatorLocalMemorySize = sizeof(float);
2673 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2674 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
2676 if (totalLocalMemorySize <= deviceLocalMemorySize)
2680 pixelPerWorkgroup = pixelPerWorkgroup/2;
2681 chunkSize = chunkSize/2;
2682 if (pixelPerWorkgroup == 0
2685 /* quit, fallback to CPU */
2691 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
2692 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
2694 if (resizeFilterType == SincFastWeightingFunction
2695 && resizeWindowType == SincFastWeightingFunction)
2696 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilterSinc");
2698 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter");
2700 if (horizontalKernel == NULL)
2702 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2707 clStatus = clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&inputImage);
2708 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageColumns);
2709 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageRows);
2710 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
2711 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&yFactor);
2712 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
2714 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
2715 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
2717 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
2718 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
2719 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
2721 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
2722 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
2724 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
2725 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
2727 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
2728 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
2730 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
2731 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
2734 clStatus |= clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
2735 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
2736 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
2737 clStatus |= clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
2740 clStatus |= clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
2741 clStatus |= clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
2742 clStatus |= clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
2744 if (clStatus != CL_SUCCESS)
2746 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
2750 global_work_size[0] = resizedColumns;
2751 global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
2753 local_work_size[0] = 1;
2754 local_work_size[1] = workgroupSize;
2755 clStatus = clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
2756 if (clStatus != CL_SUCCESS)
2758 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
2762 status = MagickTrue;
2766 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2768 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
2775 static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedColumns, const size_t resizedRows
2776 , const ResizeFilter* resizeFilter, ExceptionInfo *exception)
2779 MagickBooleanType outputReady = MagickFalse;
2780 Image* filteredImage = NULL;
2781 MagickCLEnv clEnv = NULL;
2784 MagickBooleanType status;
2785 const void *inputPixels;
2786 void* filteredPixels;
2788 const MagickRealType* resizeFilterCoefficient;
2789 float* mappedCoefficientBuffer;
2790 float xFactor, yFactor;
2791 MagickSizeType length;
2793 cl_mem_flags mem_flags;
2794 cl_context context = NULL;
2795 cl_mem inputImageBuffer = NULL;
2796 cl_mem tempImageBuffer = NULL;
2797 cl_mem filteredImageBuffer = NULL;
2798 cl_mem cubicCoefficientsBuffer = NULL;
2799 cl_command_queue queue = NULL;
2803 clEnv = GetDefaultOpenCLEnv();
2804 context = GetOpenCLContext(clEnv);
2806 /* Create and initialize OpenCL buffers. */
2808 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
2809 if (inputPixels == (const void *) NULL)
2811 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
2815 /* If the host pointer is aligned to the size of CLPixelPacket,
2816 then use the host buffer directly from the GPU; otherwise,
2817 create a buffer on the GPU and copy the data over */
2818 if (ALIGNED(inputPixels,CLPixelPacket))
2820 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2824 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2826 /* create a CL buffer from image pixel buffer */
2827 length = inputImage->columns * inputImage->rows;
2828 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2829 if (clStatus != CL_SUCCESS)
2831 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2835 cubicCoefficientsBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus);
2836 if (clStatus != CL_SUCCESS)
2838 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2841 queue = AcquireOpenCLCommandQueue(clEnv);
2842 mappedCoefficientBuffer = (float*)clEnqueueMapBuffer(queue, cubicCoefficientsBuffer, CL_TRUE, CL_MAP_WRITE, 0, 7 * sizeof(float)
2843 , 0, NULL, NULL, &clStatus);
2844 if (clStatus != CL_SUCCESS)
2846 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
2849 resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter);
2850 for (i = 0; i < 7; i++)
2852 mappedCoefficientBuffer[i] = (float) resizeFilterCoefficient[i];
2854 clStatus = clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL);
2855 if (clStatus != CL_SUCCESS)
2857 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
2861 filteredImage = CloneImage(inputImage,resizedColumns,resizedRows,MagickTrue,exception);
2862 if (filteredImage == NULL)
2865 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
2867 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2870 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
2871 if (filteredPixels == (void *) NULL)
2873 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2877 if (ALIGNED(filteredPixels,CLPixelPacket))
2879 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2880 hostPtr = filteredPixels;
2884 mem_flags = CL_MEM_WRITE_ONLY;
2888 /* create a CL buffer from image pixel buffer */
2889 length = filteredImage->columns * filteredImage->rows;
2890 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2891 if (clStatus != CL_SUCCESS)
2893 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2897 xFactor=(float) resizedColumns/(float) inputImage->columns;
2898 yFactor=(float) resizedRows/(float) inputImage->rows;
2899 if (xFactor > yFactor)
2902 length = resizedColumns*inputImage->rows;
2903 tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
2904 if (clStatus != CL_SUCCESS)
2906 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2910 status = resizeHorizontalFilter(inputImageBuffer, inputImage->columns, inputImage->rows, (inputImage->alpha_trait == BlendPixelTrait)?1:0
2911 , tempImageBuffer, resizedColumns, inputImage->rows
2912 , resizeFilter, cubicCoefficientsBuffer
2913 , xFactor, clEnv, queue, exception);
2914 if (status != MagickTrue)
2917 status = resizeVerticalFilter(tempImageBuffer, resizedColumns, inputImage->rows, (inputImage->alpha_trait == BlendPixelTrait)?1:0
2918 , filteredImageBuffer, resizedColumns, resizedRows
2919 , resizeFilter, cubicCoefficientsBuffer
2920 , yFactor, clEnv, queue, exception);
2921 if (status != MagickTrue)
2926 length = inputImage->columns*resizedRows;
2927 tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
2928 if (clStatus != CL_SUCCESS)
2930 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
2934 status = resizeVerticalFilter(inputImageBuffer, inputImage->columns, inputImage->rows, (inputImage->alpha_trait == BlendPixelTrait)?1:0
2935 , tempImageBuffer, inputImage->columns, resizedRows
2936 , resizeFilter, cubicCoefficientsBuffer
2937 , yFactor, clEnv, queue, exception);
2938 if (status != MagickTrue)
2941 status = resizeHorizontalFilter(tempImageBuffer, inputImage->columns, resizedRows, (inputImage->alpha_trait == BlendPixelTrait)?1:0
2942 , filteredImageBuffer, resizedColumns, resizedRows
2943 , resizeFilter, cubicCoefficientsBuffer
2944 , xFactor, clEnv, queue, exception);
2945 if (status != MagickTrue)
2948 length = resizedColumns*resizedRows;
2949 if (ALIGNED(filteredPixels,CLPixelPacket))
2951 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2955 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2957 if (clStatus != CL_SUCCESS)
2959 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2962 outputReady = MagickTrue;
2965 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2967 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
2968 if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer);
2969 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
2970 if (cubicCoefficientsBuffer!=NULL) clReleaseMemObject(cubicCoefficientsBuffer);
2971 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2972 if (outputReady == MagickFalse)
2974 if (filteredImage != NULL)
2976 DestroyImage(filteredImage);
2977 filteredImage = NULL;
2981 return filteredImage;
2984 const ResizeWeightingFunctionType supportedResizeWeighting[] =
2986 BoxWeightingFunction
2987 ,TriangleWeightingFunction
2988 ,HanningWeightingFunction
2989 ,HammingWeightingFunction
2990 ,BlackmanWeightingFunction
2991 ,CubicBCWeightingFunction
2992 ,SincWeightingFunction
2993 ,SincFastWeightingFunction
2994 ,LastWeightingFunction
2997 static MagickBooleanType gpuSupportedResizeWeighting(ResizeWeightingFunctionType f)
2999 MagickBooleanType supported = MagickFalse;
3003 if (supportedResizeWeighting[i] == LastWeightingFunction)
3005 if (supportedResizeWeighting[i] == f)
3007 supported = MagickTrue;
3016 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3020 % A c c e l e r a t e R e s i z e I m a g e %
3024 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3026 % AccelerateResizeImage() is an OpenCL implementation of ResizeImage()
3028 % AccelerateResizeImage() scales an image to the desired dimensions, using the given
3029 % filter (see AcquireFilterInfo()).
3031 % If an undefined filter is given the filter defaults to Mitchell for a
3032 % colormapped image, a image with a matte channel, or if the image is
3033 % enlarged. Otherwise the filter defaults to a Lanczos.
3035 % AccelerateResizeImage() was inspired by Paul Heckbert's "zoom" program.
3037 % The format of the AccelerateResizeImage method is:
3039 % Image *ResizeImage(Image *image,const size_t columns,
3040 % const size_t rows, const ResizeFilter* filter,
3041 % ExceptionInfo *exception)
3043 % A description of each parameter follows:
3045 % o image: the image.
3047 % o columns: the number of columns in the scaled image.
3049 % o rows: the number of rows in the scaled image.
3051 % o filter: Image filter to use.
3053 % o exception: return any errors or warnings in this structure.
3058 Image* AccelerateResizeImage(const Image* image, const size_t resizedColumns, const size_t resizedRows
3059 , const ResizeFilter* resizeFilter, ExceptionInfo *exception)
3061 MagickBooleanType status;
3062 Image* filteredImage;
3064 assert(image != NULL);
3065 assert(resizeFilter != NULL);
3067 status = checkOpenCLEnvironment(exception);
3068 if (status == MagickFalse)
3071 status = checkAccelerateCondition(image, AllChannels);
3072 if (status == MagickFalse)
3075 if (gpuSupportedResizeWeighting(GetResizeFilterWeightingType(resizeFilter)) == MagickFalse
3076 || gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(resizeFilter)) == MagickFalse)
3079 filteredImage = ComputeResizeImage(image,resizedColumns,resizedRows,resizeFilter,exception);
3080 return filteredImage;
3085 static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBooleanType sharpen, ExceptionInfo *exception)
3087 MagickBooleanType outputReady = MagickFalse;
3088 MagickCLEnv clEnv = NULL;
3091 size_t global_work_size[2];
3093 void *inputPixels = NULL;
3094 MagickSizeType length;
3095 unsigned int uSharpen;
3098 cl_mem_flags mem_flags;
3099 cl_context context = NULL;
3100 cl_mem inputImageBuffer = NULL;
3101 cl_kernel filterKernel = NULL;
3102 cl_command_queue queue = NULL;
3104 clEnv = GetDefaultOpenCLEnv();
3105 context = GetOpenCLContext(clEnv);
3107 /* Create and initialize OpenCL buffers. */
3108 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3109 if (inputPixels == (void *) NULL)
3111 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3115 /* If the host pointer is aligned to the size of CLPixelPacket,
3116 then use the host buffer directly from the GPU; otherwise,
3117 create a buffer on the GPU and copy the data over */
3118 if (ALIGNED(inputPixels,CLPixelPacket))
3120 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3124 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3126 /* create a CL buffer from image pixel buffer */
3127 length = inputImage->columns * inputImage->rows;
3128 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3129 if (clStatus != CL_SUCCESS)
3131 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
3135 filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast");
3136 if (filterKernel == NULL)
3138 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3143 clStatus=clSetKernelArg(filterKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3145 uSharpen = (sharpen == MagickFalse)?0:1;
3146 clStatus|=clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen);
3147 if (clStatus != CL_SUCCESS)
3149 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
3153 global_work_size[0] = inputImage->columns;
3154 global_work_size[1] = inputImage->rows;
3155 /* launch the kernel */
3156 queue = AcquireOpenCLCommandQueue(clEnv);
3157 clStatus = clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3158 if (clStatus != CL_SUCCESS)
3160 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
3165 if (ALIGNED(inputPixels,CLPixelPacket))
3167 length = inputImage->columns * inputImage->rows;
3168 clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3172 length = inputImage->columns * inputImage->rows;
3173 clStatus = clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3175 if (clStatus != CL_SUCCESS)
3177 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3180 outputReady = MagickTrue;
3183 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3185 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
3186 if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel);
3187 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3192 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3196 % C o n t r a s t I m a g e w i t h O p e n C L %
3200 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3202 % ContrastImage() enhances the intensity differences between the lighter and
3203 % darker elements of the image. Set sharpen to a MagickTrue to increase the
3204 % image contrast otherwise the contrast is reduced.
3206 % The format of the ContrastImage method is:
3208 % MagickBooleanType ContrastImage(Image *image,
3209 % const MagickBooleanType sharpen)
3211 % A description of each parameter follows:
3213 % o image: the image.
3215 % o sharpen: Increase or decrease image contrast.
3220 MagickBooleanType AccelerateContrastImage(Image* image, const MagickBooleanType sharpen, ExceptionInfo* exception)
3222 MagickBooleanType status;
3224 assert(image != NULL);
3225 assert(exception != NULL);
3227 status = checkOpenCLEnvironment(exception);
3228 if (status == MagickFalse)
3231 status = checkAccelerateCondition(image, AllChannels);
3232 if (status == MagickFalse)
3235 status = ComputeContrastImage(image,sharpen,exception);
3241 MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, double percent_hue, double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
3253 MagickBooleanType outputReady;
3259 MagickSizeType length;
3262 cl_command_queue queue;
3263 cl_kernel modulateKernel;
3265 cl_mem inputImageBuffer;
3266 cl_mem_flags mem_flags;
3270 Image * inputImage = image;
3272 inputImageBuffer = NULL;
3273 modulateKernel = NULL;
3275 assert(inputImage != (Image *) NULL);
3276 assert(inputImage->signature == MagickSignature);
3277 if (inputImage->debug != MagickFalse)
3278 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",inputImage->filename);
3281 * initialize opencl env
3283 clEnv = GetDefaultOpenCLEnv();
3284 context = GetOpenCLContext(clEnv);
3285 queue = AcquireOpenCLCommandQueue(clEnv);
3287 outputReady = MagickFalse;
3289 /* Create and initialize OpenCL buffers.
3290 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
3291 assume this will get a writable image
3293 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3294 if (inputPixels == (void *) NULL)
3296 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3300 /* If the host pointer is aligned to the size of CLPixelPacket,
3301 then use the host buffer directly from the GPU; otherwise,
3302 create a buffer on the GPU and copy the data over
3304 if (ALIGNED(inputPixels,CLPixelPacket))
3306 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3310 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3312 /* create a CL buffer from image pixel buffer */
3313 length = inputImage->columns * inputImage->rows;
3314 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3315 if (clStatus != CL_SUCCESS)
3317 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
3321 modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate");
3322 if (modulateKernel == NULL)
3324 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3328 bright=percent_brightness;
3330 saturation=percent_saturation;
3334 clStatus=clSetKernelArg(modulateKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3335 clStatus|=clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&bright);
3336 clStatus|=clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&hue);
3337 clStatus|=clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&saturation);
3338 clStatus|=clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color);
3339 if (clStatus != CL_SUCCESS)
3341 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
3342 printf("no kernel\n");
3347 size_t global_work_size[2];
3348 global_work_size[0] = inputImage->columns;
3349 global_work_size[1] = inputImage->rows;
3350 /* launch the kernel */
3351 clStatus = clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3352 if (clStatus != CL_SUCCESS)
3354 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
3360 if (ALIGNED(inputPixels,CLPixelPacket))
3362 length = inputImage->columns * inputImage->rows;
3363 clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3367 length = inputImage->columns * inputImage->rows;
3368 clStatus = clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3370 if (clStatus != CL_SUCCESS)
3372 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3376 outputReady = MagickTrue;
3379 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3382 //ReleasePixelCachePixels();
3386 if (inputImageBuffer!=NULL)
3387 clReleaseMemObject(inputImageBuffer);
3388 if (modulateKernel!=NULL)
3389 RelinquishOpenCLKernel(clEnv, modulateKernel);
3391 RelinquishOpenCLCommandQueue(clEnv, queue);
3398 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3402 % M o d u l a t e I m a g e w i t h O p e n C L %
3406 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3408 % ModulateImage() lets you control the brightness, saturation, and hue
3409 % of an image. Modulate represents the brightness, saturation, and hue
3410 % as one parameter (e.g. 90,150,100). If the image colorspace is HSL, the
3411 % modulation is lightness, saturation, and hue. For HWB, use blackness,
3412 % whiteness, and hue. And for HCL, use chrome, luma, and hue.
3414 % The format of the ModulateImage method is:
3416 % MagickBooleanType ModulateImage(Image *image,const char *modulate)
3418 % A description of each parameter follows:
3420 % o image: the image.
3422 % o percent_*: Define the percent change in brightness, saturation, and
3428 MagickBooleanType AccelerateModulateImage(Image* image, double percent_brightness, double percent_hue, double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
3430 MagickBooleanType status;
3432 assert(image != NULL);
3433 assert(exception != NULL);
3435 status = checkOpenCLEnvironment(exception);
3436 if (status == MagickFalse)
3439 status = checkAccelerateCondition(image, AllChannels);
3440 if (status == MagickFalse)
3443 if ((colorspace != HSLColorspace && colorspace != UndefinedColorspace))
3447 status = ComputeModulateImage(image,percent_brightness, percent_hue, percent_saturation, colorspace, exception);
3452 MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const ChannelType channel, ExceptionInfo * _exception)
3454 #define EqualizeImageTag "Equalize/Image"
3457 *exception=_exception;
3474 Image * image = inputImage;
3476 MagickBooleanType outputReady;
3480 size_t global_work_size[2];
3483 cl_mem_flags mem_flags;
3486 cl_mem inputImageBuffer;
3487 cl_mem histogramBuffer;
3488 cl_mem equalizeMapBuffer;
3489 cl_kernel histogramKernel;
3490 cl_kernel equalizeKernel;
3491 cl_command_queue queue;
3496 MagickSizeType length;
3499 inputImageBuffer = NULL;
3500 histogramBuffer = NULL;
3501 histogramKernel = NULL;
3502 equalizeKernel = NULL;
3505 outputReady = MagickFalse;
3507 assert(inputImage != (Image *) NULL);
3508 assert(inputImage->signature == MagickSignature);
3509 if (inputImage->debug != MagickFalse)
3510 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",inputImage->filename);
3513 Allocate and initialize histogram arrays.
3515 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram));
3516 if (histogram == (cl_uint4 *) NULL)
3517 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
3519 /* reset histogram */
3520 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram));
3523 * initialize opencl env
3525 clEnv = GetDefaultOpenCLEnv();
3526 context = GetOpenCLContext(clEnv);
3527 queue = AcquireOpenCLCommandQueue(clEnv);
3529 /* Create and initialize OpenCL buffers. */
3530 /* inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); */
3531 /* assume this will get a writable image */
3532 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3534 if (inputPixels == (void *) NULL)
3536 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3539 /* If the host pointer is aligned to the size of CLPixelPacket,
3540 then use the host buffer directly from the GPU; otherwise,
3541 create a buffer on the GPU and copy the data over */
3542 if (ALIGNED(inputPixels,CLPixelPacket))
3544 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
3548 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
3550 /* create a CL buffer from image pixel buffer */
3551 length = inputImage->columns * inputImage->rows;
3552 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3553 if (clStatus != CL_SUCCESS)
3555 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
3559 /* If the host pointer is aligned to the size of cl_uint,
3560 then use the host buffer directly from the GPU; otherwise,
3561 create a buffer on the GPU and copy the data over */
3562 if (ALIGNED(histogram,cl_uint4))
3564 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3565 hostPtr = histogram;
3569 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3570 hostPtr = histogram;
3572 /* create a CL buffer for histogram */
3573 length = (MaxMap+1);
3574 histogramBuffer = clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
3575 if (clStatus != CL_SUCCESS)
3577 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
3581 switch (inputImage->colorspace)
3586 case sRGBColorspace:
3591 /* something is wrong, as we checked in checkAccelerateCondition */
3595 /* get the OpenCL kernel */
3596 histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram");
3597 if (histogramKernel == NULL)
3599 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3603 /* set the kernel arguments */
3605 clStatus=clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3606 clStatus|=clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel);
3607 clStatus|=clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&colorspace);
3608 clStatus|=clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer);
3609 if (clStatus != CL_SUCCESS)
3611 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
3615 /* launch the kernel */
3616 global_work_size[0] = inputImage->columns;
3617 global_work_size[1] = inputImage->rows;
3619 clStatus = clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3621 if (clStatus != CL_SUCCESS)
3623 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
3628 /* read from the kenel output */
3629 if (ALIGNED(histogram,cl_uint4))
3631 length = (MaxMap+1);
3632 clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
3636 length = (MaxMap+1);
3637 clStatus = clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
3639 if (clStatus != CL_SUCCESS)
3641 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3645 /* unmap, don't block gpu to use this buffer again. */
3646 if (ALIGNED(histogram,cl_uint4))
3648 clStatus = clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
3649 if (clStatus != CL_SUCCESS)
3651 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", ".");
3656 if (getenv("TEST")) {
3658 for (i=0; i<(MaxMap+1UL); i++)
3660 printf("histogram %d: red %d\n", i, histogram[i].s[2]);
3661 printf("histogram %d: green %d\n", i, histogram[i].s[1]);
3662 printf("histogram %d: blue %d\n", i, histogram[i].s[0]);
3663 printf("histogram %d: alpha %d\n", i, histogram[i].s[3]);
3668 equalize_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*equalize_map));
3669 if (equalize_map == (PixelPacket *) NULL)
3670 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
3672 map=(FloatPixelPacket *) AcquireQuantumMemory(MaxMap+1UL,sizeof(*map));
3673 if (map == (FloatPixelPacket *) NULL)
3674 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
3677 Integrate the histogram to get the equalization map.
3679 (void) ResetMagickMemory(&intensity,0,sizeof(intensity));
3680 for (i=0; i <= (ssize_t) MaxMap; i++)
3682 if ((channel & SyncChannels) != 0)
3684 intensity.red+=histogram[i].s[2];
3688 if ((channel & RedChannel) != 0)
3689 intensity.red+=histogram[i].s[2];
3690 if ((channel & GreenChannel) != 0)
3691 intensity.green+=histogram[i].s[1];
3692 if ((channel & BlueChannel) != 0)
3693 intensity.blue+=histogram[i].s[0];
3694 if ((channel & OpacityChannel) != 0)
3695 intensity.alpha+=histogram[i].s[3];
3696 if (((channel & IndexChannel) != 0) &&
3697 (image->colorspace == CMYKColorspace))
3699 printf("something here\n");
3700 /*intensity.index+=histogram[i].index; */
3705 white=map[(int) MaxMap];
3706 (void) ResetMagickMemory(equalize_map,0,(MaxMap+1)*sizeof(*equalize_map));
3707 for (i=0; i <= (ssize_t) MaxMap; i++)
3709 if ((channel & SyncChannels) != 0)
3711 if (white.red != black.red)
3712 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
3713 (map[i].red-black.red))/(white.red-black.red)));
3716 if (((channel & RedChannel) != 0) && (white.red != black.red))
3717 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
3718 (map[i].red-black.red))/(white.red-black.red)));
3719 if (((channel & GreenChannel) != 0) && (white.green != black.green))
3720 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
3721 (map[i].green-black.green))/(white.green-black.green)));
3722 if (((channel & BlueChannel) != 0) && (white.blue != black.blue))
3723 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
3724 (map[i].blue-black.blue))/(white.blue-black.blue)));
3725 if (((channel & OpacityChannel) != 0) && (white.alpha != black.alpha))
3726 equalize_map[i].alpha=ScaleMapToQuantum((MagickRealType) ((MaxMap*
3727 (map[i].alpha-black.alpha))/(white.alpha-black.alpha)));
3729 if ((((channel & IndexChannel) != 0) &&
3730 (image->colorspace == CMYKColorspace)) &&
3731 (white.index != black.index))
3732 equalize_map[i].index=ScaleMapToQuantum((MagickRealType) ((MaxMap*
3733 (map[i].index-black.index))/(white.index-black.index)));
3737 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
3738 map=(FloatPixelPacket *) RelinquishMagickMemory(map);
3740 if (image->storage_class == PseudoClass)
3745 for (i=0; i < (ssize_t) image->colors; i++)
3747 if ((channel & SyncChannels) != 0)
3749 if (white.red != black.red)
3751 image->colormap[i].red=equalize_map[
3752 ScaleQuantumToMap(image->colormap[i].red)].red;
3753 image->colormap[i].green=equalize_map[
3754 ScaleQuantumToMap(image->colormap[i].green)].red;
3755 image->colormap[i].blue=equalize_map[
3756 ScaleQuantumToMap(image->colormap[i].blue)].red;
3757 image->colormap[i].alpha=equalize_map[
3758 ScaleQuantumToMap(image->colormap[i].alpha)].red;
3762 if (((channel & RedChannel) != 0) && (white.red != black.red))
3763 image->colormap[i].red=equalize_map[
3764 ScaleQuantumToMap(image->colormap[i].red)].red;
3765 if (((channel & GreenChannel) != 0) && (white.green != black.green))
3766 image->colormap[i].green=equalize_map[
3767 ScaleQuantumToMap(image->colormap[i].green)].green;
3768 if (((channel & BlueChannel) != 0) && (white.blue != black.blue))
3769 image->colormap[i].blue=equalize_map[
3770 ScaleQuantumToMap(image->colormap[i].blue)].blue;
3771 if (((channel & OpacityChannel) != 0) &&
3772 (white.alpha != black.alpha))
3773 image->colormap[i].alpha=equalize_map[
3774 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
3782 /* GPU can work on this again, image and equalize map as input
3783 image: uchar4 (CLPixelPacket)
3784 equalize_map: uchar4 (PixelPacket)
3785 black, white: float4 (FloatPixelPacket) */
3787 if (inputImageBuffer!=NULL)
3788 clReleaseMemObject(inputImageBuffer);
3790 /* If the host pointer is aligned to the size of CLPixelPacket,
3791 then use the host buffer directly from the GPU; otherwise,
3792 create a buffer on the GPU and copy the data over */
3793 if (ALIGNED(inputPixels,CLPixelPacket))
3795 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3799 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3801 /* create a CL buffer from image pixel buffer */
3802 length = inputImage->columns * inputImage->rows;
3803 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3804 if (clStatus != CL_SUCCESS)
3806 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
3810 /* Create and initialize OpenCL buffers. */
3811 if (ALIGNED(equalize_map, PixelPacket))
3813 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
3814 hostPtr = equalize_map;
3818 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3819 hostPtr = equalize_map;
3821 /* create a CL buffer for eqaulize_map */
3822 length = (MaxMap+1);
3823 equalizeMapBuffer = clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
3824 if (clStatus != CL_SUCCESS)
3826 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
3830 /* get the OpenCL kernel */
3831 equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize");
3832 if (equalizeKernel == NULL)
3834 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3838 /* set the kernel arguments */
3840 clStatus=clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3841 clStatus|=clSetKernelArg(equalizeKernel,i++,sizeof(ChannelType),&channel);
3842 clStatus|=clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&equalizeMapBuffer);
3843 clStatus|=clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&white);
3844 clStatus|=clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black);
3845 if (clStatus != CL_SUCCESS)
3847 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
3851 /* launch the kernel */
3852 global_work_size[0] = inputImage->columns;
3853 global_work_size[1] = inputImage->rows;
3855 clStatus = clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3857 if (clStatus != CL_SUCCESS)
3859 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
3864 /* read the data back */
3865 if (ALIGNED(inputPixels,CLPixelPacket))
3867 length = inputImage->columns * inputImage->rows;
3868 clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3872 length = inputImage->columns * inputImage->rows;
3873 clStatus = clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3875 if (clStatus != CL_SUCCESS)
3877 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3881 outputReady = MagickTrue;
3883 equalize_map=(PixelPacket *) RelinquishMagickMemory(equalize_map);
3886 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3889 /*ReleasePixelCachePixels();*/
3893 if (inputImageBuffer!=NULL)
3894 clReleaseMemObject(inputImageBuffer);
3895 if (histogramBuffer!=NULL)
3896 clReleaseMemObject(histogramBuffer);
3897 if (histogramKernel!=NULL)
3898 RelinquishOpenCLKernel(clEnv, histogramKernel);
3900 RelinquishOpenCLCommandQueue(clEnv, queue);
3906 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3910 % E q u a l i z e I m a g e w i t h O p e n C L %
3914 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3916 % EqualizeImage() applies a histogram equalization to the image.
3918 % The format of the EqualizeImage method is:
3920 % MagickBooleanType EqualizeImage(Image *image)
3921 % MagickBooleanType EqualizeImageChannel(Image *image,
3922 % const ChannelType channel)
3924 % A description of each parameter follows:
3926 % o image: the image.
3928 % o channel: the channel.
3934 MagickBooleanType AccelerateEqualizeImage(Image* image, const ChannelType channel, ExceptionInfo* exception)
3936 MagickBooleanType status;
3938 assert(image != NULL);
3939 assert(exception != NULL);
3941 status = checkOpenCLEnvironment(exception);
3942 if (status == MagickFalse)
3945 status = checkAccelerateCondition(image, channel);
3946 if (status == MagickFalse)
3949 /* ensure this is the only pass get in for now. */
3950 if ((channel & SyncChannels) == 0)
3953 if (image->colorspace != sRGBColorspace)
3956 status = ComputeEqualizeImage(image,channel,exception);
3961 static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exception)
3964 MagickBooleanType outputReady = MagickFalse;
3965 MagickCLEnv clEnv = NULL;
3968 size_t global_work_size[2];
3970 const void *inputPixels = NULL;
3971 Image* filteredImage = NULL;
3972 void *filteredPixels = NULL;
3974 MagickSizeType length;
3976 cl_mem_flags mem_flags;
3977 cl_context context = NULL;
3978 cl_mem inputImageBuffer = NULL;
3979 cl_mem tempImageBuffer[2];
3980 cl_mem filteredImageBuffer = NULL;
3981 cl_command_queue queue = NULL;
3982 cl_kernel hullPass1 = NULL;
3983 cl_kernel hullPass2 = NULL;
3985 unsigned int imageWidth, imageHeight;
3990 X[4] = {0, 1, 1,-1},
3991 Y[4] = {1, 0, 1, 1};
3993 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
3994 clEnv = GetDefaultOpenCLEnv();
3995 context = GetOpenCLContext(clEnv);
3996 queue = AcquireOpenCLCommandQueue(clEnv);
3998 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
3999 if (inputPixels == (void *) NULL)
4001 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
4005 if (ALIGNED(inputPixels,CLPixelPacket))
4007 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4011 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
4013 /* create a CL buffer from image pixel buffer */
4014 length = inputImage->columns * inputImage->rows;
4015 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4016 if (clStatus != CL_SUCCESS)
4018 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4022 mem_flags = CL_MEM_READ_WRITE;
4023 length = inputImage->columns * inputImage->rows;
4024 for (k = 0; k < 2; k++)
4026 tempImageBuffer[k] = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus);
4027 if (clStatus != CL_SUCCESS)
4029 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4034 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
4035 assert(filteredImage != NULL);
4036 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
4038 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
4041 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
4042 if (filteredPixels == (void *) NULL)
4044 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
4048 if (ALIGNED(filteredPixels,CLPixelPacket))
4050 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
4051 hostPtr = filteredPixels;
4055 mem_flags = CL_MEM_WRITE_ONLY;
4058 /* create a CL buffer from image pixel buffer */
4059 length = inputImage->columns * inputImage->rows;
4060 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
4061 if (clStatus != CL_SUCCESS)
4063 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4067 hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1");
4068 hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2");
4070 clStatus =clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&inputImageBuffer);
4071 clStatus |=clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1));
4072 imageWidth = inputImage->columns;
4073 clStatus |=clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth);
4074 imageHeight = inputImage->rows;
4075 clStatus |=clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight);
4076 matte = (inputImage->matte==MagickFalse)?0:1;
4077 clStatus |=clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
4078 if (clStatus != CL_SUCCESS)
4080 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
4084 clStatus = clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1));
4085 clStatus |=clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer);
4086 imageWidth = inputImage->columns;
4087 clStatus |=clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth);
4088 imageHeight = inputImage->rows;
4089 clStatus |=clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight);
4090 matte = (inputImage->alpha_trait == BlendPixelTrait)?1:0;
4091 clStatus |=clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
4092 if (clStatus != CL_SUCCESS)
4094 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
4099 global_work_size[0] = inputImage->columns;
4100 global_work_size[1] = inputImage->rows;
4103 for (k = 0; k < 4; k++)
4112 clStatus = clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
4113 clStatus|= clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
4114 clStatus|=clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
4115 clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
4116 if (clStatus != CL_SUCCESS)
4118 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
4121 /* launch the kernel */
4122 clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4123 if (clStatus != CL_SUCCESS)
4125 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4128 /* launch the kernel */
4129 clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4130 if (clStatus != CL_SUCCESS)
4132 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4138 clStatus =clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
4139 offset.s[0] = -X[k];
4140 offset.s[1] = -Y[k];
4142 clStatus = clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
4143 clStatus|= clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
4144 clStatus|=clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
4145 clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
4146 if (clStatus != CL_SUCCESS)
4148 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
4151 /* launch the kernel */
4152 clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4153 if (clStatus != CL_SUCCESS)
4155 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4158 /* launch the kernel */
4159 clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4160 if (clStatus != CL_SUCCESS)
4162 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4166 offset.s[0] = -X[k];
4167 offset.s[1] = -Y[k];
4169 clStatus = clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
4170 clStatus|= clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
4171 clStatus|=clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
4172 clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
4173 if (clStatus != CL_SUCCESS)
4175 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
4178 /* launch the kernel */
4179 clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4180 if (clStatus != CL_SUCCESS)
4182 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4185 /* launch the kernel */
4186 clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4187 if (clStatus != CL_SUCCESS)
4189 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4196 clStatus = clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
4197 clStatus|= clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
4198 clStatus|=clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
4199 clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
4202 clStatus |=clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer);
4204 if (clStatus != CL_SUCCESS)
4206 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", ".");
4209 /* launch the kernel */
4210 clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4211 if (clStatus != CL_SUCCESS)
4213 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4216 /* launch the kernel */
4217 clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4218 if (clStatus != CL_SUCCESS)
4220 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", ".");
4225 if (ALIGNED(filteredPixels,CLPixelPacket))
4227 length = inputImage->columns * inputImage->rows;
4228 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4232 length = inputImage->columns * inputImage->rows;
4233 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
4235 if (clStatus != CL_SUCCESS)
4237 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4241 outputReady = MagickTrue;
4244 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4246 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4247 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
4248 for (k = 0; k < 2; k++)
4250 if (tempImageBuffer[k]!=NULL) clReleaseMemObject(tempImageBuffer[k]);
4252 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
4253 if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
4254 if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
4255 if (outputReady == MagickFalse)
4257 if (filteredImage != NULL)
4259 DestroyImage(filteredImage);
4260 filteredImage = NULL;
4263 return filteredImage;
4267 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4271 % D e s p e c k l e I m a g e w i t h O p e n C L %
4275 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4277 % DespeckleImage() reduces the speckle noise in an image while perserving the
4278 % edges of the original image. A speckle removing filter uses a complementary
4279 % hulling technique (raising pixels that are darker than their surrounding
4280 % neighbors, then complementarily lowering pixels that are brighter than their
4281 % surrounding neighbors) to reduce the speckle index of that image (reference
4282 % Crimmins speckle removal).
4284 % The format of the DespeckleImage method is:
4286 % Image *DespeckleImage(const Image *image,ExceptionInfo *exception)
4288 % A description of each parameter follows:
4290 % o image: the image.
4292 % o exception: return any errors or warnings in this structure.
4297 Image* AccelerateDespeckleImage(const Image* image, ExceptionInfo* exception)
4299 MagickBooleanType status;
4300 Image* newImage = NULL;
4302 assert(image != NULL);
4303 assert(exception != NULL);
4305 status = checkOpenCLEnvironment(exception);
4306 if (status == MagickFalse)
4309 status = checkAccelerateCondition(image, AllChannels);
4310 if (status == MagickFalse)
4313 newImage = ComputeDespeckleImage(image,exception);
4317 static Image* ComputeAddNoiseImage(const Image* inputImage,
4318 const ChannelType channel, const NoiseType noise_type,
4319 ExceptionInfo *exception)
4321 MagickBooleanType outputReady = MagickFalse;
4322 MagickCLEnv clEnv = NULL;
4325 size_t global_work_size[2];
4327 const void *inputPixels = NULL;
4328 Image* filteredImage = NULL;
4329 void *filteredPixels = NULL;
4331 unsigned int inputColumns, inputRows;
4333 float *randomNumberBufferPtr = NULL;
4334 MagickSizeType length;
4335 unsigned int numRandomNumberPerPixel;
4336 unsigned int numRowsPerKernelLaunch;
4337 unsigned int numRandomNumberPerBuffer;
4342 RandomInfo **restrict random_info;
4344 #if defined(MAGICKCORE_OPENMP_SUPPORT)
4348 cl_mem_flags mem_flags;
4349 cl_context context = NULL;
4350 cl_mem inputImageBuffer = NULL;
4351 cl_mem randomNumberBuffer = NULL;
4352 cl_mem filteredImageBuffer = NULL;
4353 cl_command_queue queue = NULL;
4354 cl_kernel addNoiseKernel = NULL;
4357 clEnv = GetDefaultOpenCLEnv();
4358 context = GetOpenCLContext(clEnv);
4359 queue = AcquireOpenCLCommandQueue(clEnv);
4361 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
4362 if (inputPixels == (void *) NULL)
4364 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
4368 if (ALIGNED(inputPixels,CLPixelPacket))
4370 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4374 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
4376 /* create a CL buffer from image pixel buffer */
4377 length = inputImage->columns * inputImage->rows;
4378 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4379 if (clStatus != CL_SUCCESS)
4381 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4386 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
4387 assert(filteredImage != NULL);
4388 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
4390 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
4393 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
4394 if (filteredPixels == (void *) NULL)
4396 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
4400 if (ALIGNED(filteredPixels,CLPixelPacket))
4402 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
4403 hostPtr = filteredPixels;
4407 mem_flags = CL_MEM_WRITE_ONLY;
4410 /* create a CL buffer from image pixel buffer */
4411 length = inputImage->columns * inputImage->rows;
4412 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
4413 if (clStatus != CL_SUCCESS)
4415 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4419 /* find out how many random numbers needed by pixel */
4420 numRandomNumberPerPixel = 0;
4422 unsigned int numRandPerChannel = 0;
4427 case LaplacianNoise:
4430 numRandPerChannel = 1;
4433 case MultiplicativeGaussianNoise:
4435 numRandPerChannel = 2;
4439 if ((channel & RedChannel) != 0)
4440 numRandomNumberPerPixel+=numRandPerChannel;
4441 if ((channel & GreenChannel) != 0)
4442 numRandomNumberPerPixel+=numRandPerChannel;
4443 if ((channel & BlueChannel) != 0)
4444 numRandomNumberPerPixel+=numRandPerChannel;
4445 if ((channel & OpacityChannel) != 0)
4446 numRandomNumberPerPixel+=numRandPerChannel;
4449 numRowsPerKernelLaunch = 512;
4450 /* create a buffer for random numbers */
4451 numRandomNumberPerBuffer = (inputImage->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
4452 randomNumberBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, numRandomNumberPerBuffer*sizeof(float)
4456 /* set up the random number generators */
4458 option=GetImageArtifact(inputImage,"attenuate");
4459 if (option != (char *) NULL)
4460 attenuate=StringToDouble(option,(char **) NULL);
4461 random_info=AcquireRandomInfoThreadSet();
4462 #if defined(MAGICKCORE_OPENMP_SUPPORT)
4463 key=GetRandomSecretKey(random_info[0]);
4466 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
4469 clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&inputImageBuffer);
4470 clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
4471 inputColumns = inputImage->columns;
4472 clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
4473 inputRows = inputImage->rows;
4474 clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
4475 clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
4476 clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
4478 option=GetImageArtifact(inputImage,"attenuate");
4479 if (option != (char *) NULL)
4480 attenuate=(float)StringToDouble(option,(char **) NULL);
4481 clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
4482 clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
4483 clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
4485 global_work_size[0] = inputColumns;
4486 for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
4488 /* Generate random numbers in the buffer */
4489 randomNumberBufferPtr = (float*)clEnqueueMapBuffer(queue, randomNumberBuffer, CL_TRUE, CL_MAP_WRITE, 0
4490 , numRandomNumberPerBuffer*sizeof(float), 0, NULL, NULL, &clStatus);
4491 if (clStatus != CL_SUCCESS)
4493 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
4497 #if defined(MAGICKCORE_OPENMP_SUPPORT)
4498 #pragma omp parallel for schedule(static,4) \
4499 num_threads((key == ~0UL) == 0 ? 1 : (size_t) GetMagickResourceLimit(ThreadResource))
4501 for (i = 0; i < numRandomNumberPerBuffer; i++)
4503 const int id = GetOpenMPThreadId();
4504 randomNumberBufferPtr[i] = (float)GetPseudoRandomValue(random_info[id]);
4507 clStatus = clEnqueueUnmapMemObject(queue, randomNumberBuffer, randomNumberBufferPtr, 0, NULL, NULL);
4508 if (clStatus != CL_SUCCESS)
4510 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.",".");
4514 /* set the row offset */
4515 clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
4516 global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
4517 clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
4520 if (ALIGNED(filteredPixels,CLPixelPacket))
4522 length = inputImage->columns * inputImage->rows;
4523 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4527 length = inputImage->columns * inputImage->rows;
4528 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
4530 if (clStatus != CL_SUCCESS)
4532 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4536 outputReady = MagickTrue;
4539 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4541 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4542 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
4543 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
4544 if (randomNumberBuffer!=NULL) clReleaseMemObject(randomNumberBuffer);
4545 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
4546 if (outputReady == MagickFalse
4547 && filteredImage != NULL)
4549 DestroyImage(filteredImage);
4550 filteredImage = NULL;
4552 return filteredImage;
4556 static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage,
4557 const ChannelType channel, const NoiseType noise_type,
4558 ExceptionInfo *exception)
4560 MagickBooleanType outputReady = MagickFalse;
4561 MagickCLEnv clEnv = NULL;
4564 size_t global_work_size[2];
4565 size_t random_work_size;
4567 const void *inputPixels = NULL;
4568 Image* filteredImage = NULL;
4569 void *filteredPixels = NULL;
4571 unsigned int inputColumns, inputRows;
4573 MagickSizeType length;
4574 unsigned int numRandomNumberPerPixel;
4575 unsigned int numRowsPerKernelLaunch;
4576 unsigned int numRandomNumberPerBuffer;
4577 unsigned int numRandomNumberGenerators;
4578 unsigned int initRandom;
4585 cl_mem_flags mem_flags;
4586 cl_context context = NULL;
4587 cl_mem inputImageBuffer = NULL;
4588 cl_mem randomNumberBuffer = NULL;
4589 cl_mem filteredImageBuffer = NULL;
4590 cl_mem randomNumberSeedsBuffer = NULL;
4591 cl_command_queue queue = NULL;
4592 cl_kernel addNoiseKernel = NULL;
4593 cl_kernel randomNumberGeneratorKernel = NULL;
4596 clEnv = GetDefaultOpenCLEnv();
4597 context = GetOpenCLContext(clEnv);
4598 queue = AcquireOpenCLCommandQueue(clEnv);
4600 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
4601 if (inputPixels == (void *) NULL)
4603 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
4607 if (ALIGNED(inputPixels,CLPixelPacket))
4609 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4613 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
4615 /* create a CL buffer from image pixel buffer */
4616 length = inputImage->columns * inputImage->rows;
4617 inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4618 if (clStatus != CL_SUCCESS)
4620 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4625 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
4626 assert(filteredImage != NULL);
4627 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
4629 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
4632 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
4633 if (filteredPixels == (void *) NULL)
4635 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
4639 if (ALIGNED(filteredPixels,CLPixelPacket))
4641 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
4642 hostPtr = filteredPixels;
4646 mem_flags = CL_MEM_WRITE_ONLY;
4649 /* create a CL buffer from image pixel buffer */
4650 length = inputImage->columns * inputImage->rows;
4651 filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
4652 if (clStatus != CL_SUCCESS)
4654 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4658 /* find out how many random numbers needed by pixel */
4659 numRandomNumberPerPixel = 0;
4661 unsigned int numRandPerChannel = 0;
4666 case LaplacianNoise:
4669 numRandPerChannel = 1;
4672 case MultiplicativeGaussianNoise:
4674 numRandPerChannel = 2;
4678 if ((channel & RedChannel) != 0)
4679 numRandomNumberPerPixel+=numRandPerChannel;
4680 if ((channel & GreenChannel) != 0)
4681 numRandomNumberPerPixel+=numRandPerChannel;
4682 if ((channel & BlueChannel) != 0)
4683 numRandomNumberPerPixel+=numRandPerChannel;
4684 if ((channel & OpacityChannel) != 0)
4685 numRandomNumberPerPixel+=numRandPerChannel;
4688 numRowsPerKernelLaunch = 512;
4690 /* create a buffer for random numbers */
4691 numRandomNumberPerBuffer = (inputImage->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
4692 randomNumberBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, numRandomNumberPerBuffer*sizeof(float)
4696 /* setup the random number generators */
4697 unsigned long* seeds;
4698 numRandomNumberGenerators = 512;
4699 randomNumberSeedsBuffer = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR|CL_MEM_READ_WRITE
4700 , numRandomNumberGenerators * 4 * sizeof(unsigned long), NULL, &clStatus);
4701 if (clStatus != CL_SUCCESS)
4703 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.",".");
4706 seeds = (unsigned long*) clEnqueueMapBuffer(queue, randomNumberSeedsBuffer, CL_TRUE, CL_MAP_WRITE, 0
4707 , numRandomNumberGenerators*4*sizeof(unsigned long), 0, NULL, NULL, &clStatus);
4708 if (clStatus != CL_SUCCESS)
4710 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.",".");
4714 for (i = 0; i < numRandomNumberGenerators; i++) {
4715 RandomInfo* randomInfo = AcquireRandomInfo();
4716 const unsigned long* s = GetRandomInfoSeed(randomInfo);
4719 fNormalize = GetRandomInfoNormalize(randomInfo);
4722 randomInfo = DestroyRandomInfo(randomInfo);
4725 clStatus = clEnqueueUnmapMemObject(queue, randomNumberSeedsBuffer, seeds, 0, NULL, NULL);
4726 if (clStatus != CL_SUCCESS)
4728 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.",".");
4732 randomNumberGeneratorKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE
4733 ,"randomNumberGeneratorKernel");
4736 clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberSeedsBuffer);
4737 clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(float),(void *)&fNormalize);
4738 clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
4740 clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&initRandom);
4741 clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerBuffer);
4743 random_work_size = numRandomNumberGenerators;
4746 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
4748 clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&inputImageBuffer);
4749 clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
4750 inputColumns = inputImage->columns;
4751 clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
4752 inputRows = inputImage->rows;
4753 clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
4754 clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
4755 clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
4757 option=GetImageArtifact(inputImage,"attenuate");
4758 if (option != (char *) NULL)
4759 attenuate=(float)StringToDouble(option,(char **) NULL);
4760 clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
4761 clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
4762 clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
4764 global_work_size[0] = inputColumns;
4765 for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
4767 size_t generator_local_size = 64;
4768 /* Generate random numbers in the buffer */
4769 clEnqueueNDRangeKernel(queue,randomNumberGeneratorKernel,1,NULL
4770 ,&random_work_size,&generator_local_size,0,NULL,NULL);
4771 if (initRandom != 0)
4773 /* make sure we only do init once */
4775 clSetKernelArg(randomNumberGeneratorKernel,3,sizeof(unsigned int),(void *)&initRandom);
4778 /* set the row offset */
4779 clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
4780 global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
4781 clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
4784 if (ALIGNED(filteredPixels,CLPixelPacket))
4786 length = inputImage->columns * inputImage->rows;
4787 clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4791 length = inputImage->columns * inputImage->rows;
4792 clStatus = clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
4794 if (clStatus != CL_SUCCESS)
4796 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4800 outputReady = MagickTrue;
4803 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4805 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4806 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
4807 if (randomNumberGeneratorKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomNumberGeneratorKernel);
4808 if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer);
4809 if (randomNumberBuffer!=NULL) clReleaseMemObject(randomNumberBuffer);
4810 if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer);
4811 if (randomNumberSeedsBuffer!=NULL) clReleaseMemObject(randomNumberSeedsBuffer);
4812 if (outputReady == MagickFalse
4813 && filteredImage != NULL)
4815 DestroyImage(filteredImage);
4816 filteredImage = NULL;
4818 return filteredImage;
4824 Image* AccelerateAddNoiseImage(const Image *image, const ChannelType channel,
4825 const NoiseType noise_type,ExceptionInfo *exception)
4827 MagickBooleanType status;
4828 Image* filteredImage = NULL;
4830 assert(image != NULL);
4831 assert(exception != NULL);
4833 status = checkOpenCLEnvironment(exception);
4834 if (status == MagickFalse)
4837 status = checkAccelerateCondition(image, channel);
4838 if (status == MagickFalse)
4841 DisableMSCWarning(4127)
4842 if (sizeof(unsigned long) == 4)
4844 filteredImage = ComputeAddNoiseImageOptRandomNum(image,channel,noise_type,exception);
4846 filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
4848 return filteredImage;
4852 #else /* MAGICKCORE_OPENCL_SUPPORT */
4854 MagickExport Image *AccelerateConvolveImageChannel(
4855 const Image *magick_unused(image),const ChannelType magick_unused(channel),
4856 const KernelInfo *magick_unused(kernel),
4857 ExceptionInfo *magick_unused(exception))
4859 magick_unreferenced(image);
4860 magick_unreferenced(channel);
4861 magick_unreferenced(kernel);
4862 magick_unreferenced(exception);
4867 MagickExport MagickBooleanType AccelerateFunctionImage(
4868 Image *magick_unused(image),const ChannelType magick_unused(channel),
4869 const MagickFunction magick_unused(function),
4870 const size_t magick_unused(number_parameters),
4871 const double *magick_unused(parameters),
4872 ExceptionInfo *magick_unused(exception))
4874 magick_unreferenced(image);
4875 magick_unreferenced(channel);
4876 magick_unreferenced(function);
4877 magick_unreferenced(number_parameters);
4878 magick_unreferenced(parameters);
4879 magick_unreferenced(exception);
4884 MagickExport Image *AccelerateBlurImage(const Image *magick_unused(image),
4885 const ChannelType magick_unused(channel),const double magick_unused(radius),
4886 const double magick_unused(sigma),ExceptionInfo *magick_unused(exception))
4888 magick_unreferenced(image);
4889 magick_unreferenced(channel);
4890 magick_unreferenced(radius);
4891 magick_unreferenced(sigma);
4892 magick_unreferenced(exception);
4897 MagickExport Image *AccelerateRadialBlurImage(
4898 const Image *magick_unused(image),const ChannelType magick_unused(channel),
4899 const double magick_unused(angle),ExceptionInfo *magick_unused(exception))
4901 magick_unreferenced(image);
4902 magick_unreferenced(channel);
4903 magick_unreferenced(angle);
4904 magick_unreferenced(exception);
4910 MagickExport Image *AccelerateUnsharpMaskImage(
4911 const Image *magick_unused(image),const ChannelType magick_unused(channel),
4912 const double magick_unused(radius),const double magick_unused(sigma),
4913 const double magick_unused(gain),const double magick_unused(threshold),
4914 ExceptionInfo *magick_unused(exception))
4916 magick_unreferenced(image);
4917 magick_unreferenced(channel);
4918 magick_unreferenced(radius);
4919 magick_unreferenced(sigma);
4920 magick_unreferenced(gain);
4921 magick_unreferenced(threshold);
4922 magick_unreferenced(exception);
4928 MagickExport MagickBooleanType AccelerateContrastImage(
4929 Image* magick_unused(image),const MagickBooleanType magick_unused(sharpen),
4930 ExceptionInfo* magick_unused(exception))
4932 magick_unreferenced(image);
4933 magick_unreferenced(sharpen);
4934 magick_unreferenced(exception);
4939 MagickExport MagickBooleanType AccelerateEqualizeImage(
4940 Image* magick_unused(image), const ChannelType magick_unused(channel),
4941 ExceptionInfo* magick_unused(exception))
4943 magick_unreferenced(image);
4944 magick_unreferenced(channel);
4945 magick_unreferenced(exception);
4950 MagickExport Image *AccelerateDespeckleImage(const Image* magick_unused(image),
4951 ExceptionInfo* magick_unused(exception))
4953 magick_unreferenced(image);
4954 magick_unreferenced(exception);
4959 MagickExport Image *AccelerateResizeImage(const Image* magick_unused(image),
4960 const size_t magick_unused(resizedColumns),
4961 const size_t magick_unused(resizedRows),
4962 const ResizeFilter* magick_unused(resizeFilter),
4963 ExceptionInfo *magick_unused(exception))
4965 magick_unreferenced(image);
4966 magick_unreferenced(resizedColumns);
4967 magick_unreferenced(resizedRows);
4968 magick_unreferenced(resizeFilter);
4969 magick_unreferenced(exception);
4976 MagickBooleanType AccelerateModulateImage(
4977 Image* image, double percent_brightness, double percent_hue,
4978 double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
4980 magick_unreferenced(image);
4981 magick_unreferenced(percent_brightness);
4982 magick_unreferenced(percent_hue);
4983 magick_unreferenced(percent_saturation);
4984 magick_unreferenced(colorspace);
4985 magick_unreferenced(exception);
4986 return(MagickFalse);
4989 MagickExport Image *AccelerateAddNoiseImage(const Image *image,
4990 const ChannelType channel, const NoiseType noise_type,ExceptionInfo *exception)
4992 magick_unreferenced(image);
4993 magick_unreferenced(channel);
4994 magick_unreferenced(noise_type);
4995 magick_unreferenced(exception);
4999 #endif /* MAGICKCORE_OPENCL_SUPPORT */
5001 MagickExport MagickBooleanType AccelerateConvolveImage(
5002 const Image *magick_unused(image),const KernelInfo *magick_unused(kernel),
5003 Image *magick_unused(convolve_image),ExceptionInfo *magick_unused(exception))
5005 magick_unreferenced(image);
5006 magick_unreferenced(kernel);
5007 magick_unreferenced(convolve_image);
5008 magick_unreferenced(exception);
5010 /* legacy, do not use */
5011 return(MagickFalse);