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 /* pad the global workgroup size to the next multiple of
93 the local workgroup size */
94 inline static unsigned int
95 padGlobalWorkgroupSizeToLocalWorkgroupSize(const unsigned int orgGlobalSize,
96 const unsigned int localGroupSize)
98 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
101 static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception)
103 MagickBooleanType flag;
106 clEnv = GetDefaultOpenCLEnv();
108 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
109 , sizeof(MagickBooleanType), &flag, exception);
110 if (flag != MagickFalse)
113 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED
114 , sizeof(MagickBooleanType), &flag, exception);
115 if (flag == MagickFalse)
117 if(InitOpenCLEnv(clEnv, exception) == MagickFalse)
120 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
121 , sizeof(MagickBooleanType), &flag, exception);
122 if (flag != MagickFalse)
130 static MagickBooleanType checkAccelerateCondition(const Image* image, const ChannelType channel)
132 /* check if the image's colorspace is supported */
133 if (image->colorspace != RGBColorspace
134 && image->colorspace != sRGBColorspace
135 && image->colorspace != GRAYColorspace)
138 /* check if the channel is supported */
139 if (((channel&RedChannel) == 0)
140 || ((channel&GreenChannel) == 0)
141 || ((channel&BlueChannel) == 0))
147 /* check if if the virtual pixel method is compatible with the OpenCL implementation */
148 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod)&&
149 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
155 static MagickBooleanType checkHistogramCondition(Image *image, const ChannelType channel)
158 /* ensure this is the only pass get in for now. */
159 if ((channel & SyncChannels) == 0)
162 if (image->intensity == Rec601LuminancePixelIntensityMethod ||
163 image->intensity == Rec709LuminancePixelIntensityMethod)
166 if (image->colorspace != sRGBColorspace)
173 static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType channel, const KernelInfo *kernel, ExceptionInfo *exception)
175 MagickBooleanType outputReady;
179 size_t global_work_size[3];
180 size_t localGroupSize[3];
181 size_t localMemoryRequirement;
182 Image* filteredImage;
183 MagickSizeType length;
184 const void *inputPixels;
185 void *filteredPixels;
186 cl_mem_flags mem_flags;
187 float* kernelBufferPtr;
192 filterWidth, filterHeight,
193 imageWidth, imageHeight;
197 cl_mem inputImageBuffer, filteredImageBuffer, convolutionKernel;
198 cl_ulong deviceLocalMemorySize;
200 cl_command_queue queue;
202 /* intialize all CL objects to NULL */
204 inputImageBuffer = NULL;
205 filteredImageBuffer = NULL;
206 convolutionKernel = NULL;
210 filteredImage = NULL;
211 outputReady = MagickFalse;
213 clEnv = GetDefaultOpenCLEnv();
214 context = GetOpenCLContext(clEnv);
217 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
218 if (inputPixels == (const void *) NULL)
220 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
224 /* Create and initialize OpenCL buffers. */
226 /* If the host pointer is aligned to the size of CLPixelPacket,
227 then use the host buffer directly from the GPU; otherwise,
228 create a buffer on the GPU and copy the data over */
229 if (ALIGNED(inputPixels,CLPixelPacket))
231 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
235 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
237 /* create a CL buffer from image pixel buffer */
238 length = inputImage->columns * inputImage->rows;
239 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
240 if (clStatus != CL_SUCCESS)
242 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
246 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
247 assert(filteredImage != NULL);
248 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
250 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
253 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
254 if (filteredPixels == (void *) NULL)
256 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
260 if (ALIGNED(filteredPixels,CLPixelPacket))
262 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
263 hostPtr = filteredPixels;
267 mem_flags = CL_MEM_WRITE_ONLY;
270 /* create a CL buffer from image pixel buffer */
271 length = inputImage->columns * inputImage->rows;
272 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
273 if (clStatus != CL_SUCCESS)
275 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
279 kernelSize = kernel->width * kernel->height;
280 convolutionKernel = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus);
281 if (clStatus != CL_SUCCESS)
283 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
287 queue = AcquireOpenCLCommandQueue(clEnv);
289 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize * sizeof(float)
290 , 0, NULL, NULL, &clStatus);
291 if (clStatus != CL_SUCCESS)
293 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
296 for (i = 0; i < kernelSize; i++)
298 kernelBufferPtr[i] = (float) kernel->values[i];
300 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
301 if (clStatus != CL_SUCCESS)
303 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
306 clEnv->library->clFlush(queue);
308 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
310 /* Compute the local memory requirement for a 16x16 workgroup.
311 If it's larger than 16k, reduce the workgroup size to 8x8 */
312 localGroupSize[0] = 16;
313 localGroupSize[1] = 16;
314 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
315 + kernel->width*kernel->height*sizeof(float);
317 if (localMemoryRequirement > deviceLocalMemorySize)
319 localGroupSize[0] = 8;
320 localGroupSize[1] = 8;
321 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
322 + kernel->width*kernel->height*sizeof(float);
324 if (localMemoryRequirement <= deviceLocalMemorySize)
326 /* get the OpenCL kernel */
327 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized");
328 if (clkernel == NULL)
330 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
334 /* set the kernel arguments */
336 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
337 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
338 imageWidth = inputImage->columns;
339 imageHeight = inputImage->rows;
340 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
341 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
342 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
343 filterWidth = kernel->width;
344 filterHeight = kernel->height;
345 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
346 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
347 matte = (inputImage->matte==MagickTrue)?1:0;
348 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
349 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
350 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*sizeof(CLPixelPacket),NULL);
351 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL);
352 if (clStatus != CL_SUCCESS)
354 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
358 /* pad the global size to a multiple of the local work size dimension */
359 global_work_size[0] = ((inputImage->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ;
360 global_work_size[1] = ((inputImage->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1];
362 /* launch the kernel */
363 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, NULL);
364 if (clStatus != CL_SUCCESS)
366 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
372 /* get the OpenCL kernel */
373 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve");
374 if (clkernel == NULL)
376 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
380 /* set the kernel arguments */
382 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
383 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
384 imageWidth = inputImage->columns;
385 imageHeight = inputImage->rows;
386 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
387 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
388 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
389 filterWidth = kernel->width;
390 filterHeight = kernel->height;
391 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
392 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
393 matte = (inputImage->matte==MagickTrue)?1:0;
394 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
395 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
396 if (clStatus != CL_SUCCESS)
398 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
402 localGroupSize[0] = 8;
403 localGroupSize[1] = 8;
404 global_work_size[0] = (inputImage->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0];
405 global_work_size[1] = (inputImage->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1];
406 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, NULL);
408 if (clStatus != CL_SUCCESS)
410 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
414 clEnv->library->clFlush(queue);
416 if (ALIGNED(filteredPixels,CLPixelPacket))
418 length = inputImage->columns * inputImage->rows;
419 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
423 length = inputImage->columns * inputImage->rows;
424 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
426 if (clStatus != CL_SUCCESS)
428 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
432 /* everything is fine! :) */
433 outputReady = MagickTrue;
436 OpenCLLogException(__FUNCTION__,__LINE__,exception);
438 if (inputImageBuffer != NULL)
439 clEnv->library->clReleaseMemObject(inputImageBuffer);
441 if (filteredImageBuffer != NULL)
442 clEnv->library->clReleaseMemObject(filteredImageBuffer);
444 if (convolutionKernel != NULL)
445 clEnv->library->clReleaseMemObject(convolutionKernel);
447 if (clkernel != NULL)
448 RelinquishOpenCLKernel(clEnv, clkernel);
451 RelinquishOpenCLCommandQueue(clEnv, queue);
453 if (outputReady == MagickFalse)
455 if (filteredImage != NULL)
457 DestroyImage(filteredImage);
458 filteredImage = NULL;
462 return filteredImage;
466 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
470 % C o n v o l v e I m a g e w i t h O p e n C L %
474 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
476 % ConvolveImage() applies a custom convolution kernel to the image.
478 % The format of the ConvolveImage method is:
480 % Image *ConvolveImage(const Image *image,const size_t order,
481 % const double *kernel,ExceptionInfo *exception)
482 % Image *ConvolveImageChannel(const Image *image,const ChannelType channel,
483 % const size_t order,const double *kernel,ExceptionInfo *exception)
485 % A description of each parameter follows:
487 % o image: the image.
489 % o channel: the channel type.
491 % o kernel: kernel info.
493 % o exception: return any errors or warnings in this structure.
497 MagickExport Image* AccelerateConvolveImageChannel(const Image *image, const ChannelType channel, const KernelInfo *kernel, ExceptionInfo *exception)
499 MagickBooleanType status;
500 Image* filteredImage = NULL;
502 assert(image != NULL);
503 assert(kernel != (KernelInfo *) NULL);
504 assert(exception != (ExceptionInfo *) NULL);
506 status = checkOpenCLEnvironment(exception);
507 if (status == MagickFalse)
510 status = checkAccelerateCondition(image, channel);
511 if (status == MagickFalse)
514 filteredImage = ComputeConvolveImage(image, channel, kernel, exception);
515 return filteredImage;
518 static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType channel,const MagickFunction function,
519 const size_t number_parameters,const double *parameters, ExceptionInfo *exception)
521 MagickBooleanType status;
525 MagickSizeType length;
527 float* parametersBufferPtr;
532 cl_command_queue queue;
533 cl_mem_flags mem_flags;
535 cl_mem parametersBuffer;
536 size_t globalWorkSize[2];
540 status = MagickFalse;
546 parametersBuffer = NULL;
548 clEnv = GetDefaultOpenCLEnv();
549 context = GetOpenCLContext(clEnv);
551 pixels = GetPixelCachePixels(image, &length, exception);
552 if (pixels == (void *) NULL)
554 (void) OpenCLThrowMagickException(exception, GetMagickModule(), CacheWarning,
555 "GetPixelCachePixels failed.",
556 "'%s'", image->filename);
561 if (ALIGNED(pixels,CLPixelPacket))
563 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
567 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
569 /* create a CL buffer from image pixel buffer */
570 length = image->columns * image->rows;
571 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)pixels, &clStatus);
572 if (clStatus != CL_SUCCESS)
574 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
578 parametersBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus);
579 if (clStatus != CL_SUCCESS)
581 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
585 queue = AcquireOpenCLCommandQueue(clEnv);
587 parametersBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, parametersBuffer, CL_TRUE, CL_MAP_WRITE, 0, number_parameters * sizeof(float)
588 , 0, NULL, NULL, &clStatus);
589 if (clStatus != CL_SUCCESS)
591 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
594 for (i = 0; i < number_parameters; i++)
596 parametersBufferPtr[i] = (float)parameters[i];
598 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL);
599 if (clStatus != CL_SUCCESS)
601 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
604 clEnv->library->clFlush(queue);
606 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "FunctionImage");
607 if (clkernel == NULL)
609 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
613 /* set the kernel arguments */
615 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
616 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
617 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(MagickFunction),(void *)&function);
618 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&number_parameters);
619 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)¶metersBuffer);
620 if (clStatus != CL_SUCCESS)
622 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
626 globalWorkSize[0] = image->columns;
627 globalWorkSize[1] = image->rows;
628 /* launch the kernel */
629 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, NULL);
630 if (clStatus != CL_SUCCESS)
632 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
635 clEnv->library->clFlush(queue);
638 if (ALIGNED(pixels,CLPixelPacket))
640 length = image->columns * image->rows;
641 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
645 length = image->columns * image->rows;
646 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), pixels, 0, NULL, NULL);
648 if (clStatus != CL_SUCCESS)
650 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
656 OpenCLLogException(__FUNCTION__,__LINE__,exception);
658 if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel);
659 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
660 if (imageBuffer != NULL) clEnv->library->clReleaseMemObject(imageBuffer);
661 if (parametersBuffer != NULL) clEnv->library->clReleaseMemObject(parametersBuffer);
668 MagickExport MagickBooleanType
669 AccelerateFunctionImage(Image *image, const ChannelType channel,const MagickFunction function,
670 const size_t number_parameters,const double *parameters, ExceptionInfo *exception)
672 MagickBooleanType status;
674 status = MagickFalse;
676 assert(image != NULL);
677 assert(exception != (ExceptionInfo *) NULL);
679 status = checkOpenCLEnvironment(exception);
680 if (status != MagickFalse)
682 status = checkAccelerateCondition(image, channel);
683 if (status != MagickFalse)
685 status = ComputeFunctionImage(image, channel, function, number_parameters, parameters, exception);
692 static MagickBooleanType splitImage(const Image* inputImage)
694 MagickBooleanType split;
697 unsigned long allocSize;
698 unsigned long tempSize;
700 clEnv = GetDefaultOpenCLEnv();
702 allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
703 tempSize = inputImage->columns * inputImage->rows * 4 * 4;
706 printf("alloc size: %lu\n", allocSize);
707 printf("temp size: %lu\n", tempSize);
710 split = ((tempSize > allocSize) ? MagickTrue:MagickFalse);
715 static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channel, const double radius, const double sigma, ExceptionInfo *exception)
717 MagickBooleanType outputReady;
718 Image* filteredImage;
723 const void *inputPixels;
724 void *filteredPixels;
725 cl_mem_flags mem_flags;
728 cl_mem inputImageBuffer, tempImageBuffer, filteredImageBuffer, imageKernelBuffer;
729 cl_kernel blurRowKernel, blurColumnKernel;
730 cl_command_queue queue;
733 float* kernelBufferPtr;
734 MagickSizeType length;
736 char geometry[MaxTextExtent];
737 KernelInfo* kernel = NULL;
738 unsigned int kernelWidth;
739 unsigned int imageColumns, imageRows;
744 filteredImage = NULL;
745 inputImageBuffer = NULL;
746 tempImageBuffer = NULL;
747 filteredImageBuffer = NULL;
748 imageKernelBuffer = NULL;
749 blurRowKernel = NULL;
750 blurColumnKernel = NULL;
753 outputReady = MagickFalse;
755 clEnv = GetDefaultOpenCLEnv();
756 context = GetOpenCLContext(clEnv);
757 queue = AcquireOpenCLCommandQueue(clEnv);
759 /* Create and initialize OpenCL buffers. */
762 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
763 if (inputPixels == (const void *) NULL)
765 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
768 /* If the host pointer is aligned to the size of CLPixelPacket,
769 then use the host buffer directly from the GPU; otherwise,
770 create a buffer on the GPU and copy the data over */
771 if (ALIGNED(inputPixels,CLPixelPacket))
773 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
777 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
779 /* create a CL buffer from image pixel buffer */
780 length = inputImage->columns * inputImage->rows;
781 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
782 if (clStatus != CL_SUCCESS)
784 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
791 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
792 assert(filteredImage != NULL);
793 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
795 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
798 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
799 if (filteredPixels == (void *) NULL)
801 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
805 if (ALIGNED(filteredPixels,CLPixelPacket))
807 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
808 hostPtr = filteredPixels;
812 mem_flags = CL_MEM_WRITE_ONLY;
815 /* create a CL buffer from image pixel buffer */
816 length = inputImage->columns * inputImage->rows;
817 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
818 if (clStatus != CL_SUCCESS)
820 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
825 /* create processing kernel */
827 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
828 kernel=AcquireKernelInfo(geometry);
829 if (kernel == (KernelInfo *) NULL)
831 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
835 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus);
836 if (clStatus != CL_SUCCESS)
838 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
841 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
842 if (clStatus != CL_SUCCESS)
844 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
848 for (i = 0; i < kernel->width; i++)
850 kernelBufferPtr[i] = (float) kernel->values[i];
853 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
854 if (clStatus != CL_SUCCESS)
856 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
863 /* create temp buffer */
865 length = inputImage->columns * inputImage->rows;
866 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
867 if (clStatus != CL_SUCCESS)
869 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
874 /* get the OpenCL kernels */
876 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
877 if (blurRowKernel == NULL)
879 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
883 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn");
884 if (blurColumnKernel == NULL)
886 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
892 /* need logic to decide this value */
896 imageColumns = inputImage->columns;
897 imageRows = inputImage->rows;
899 /* set the kernel arguments */
901 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
902 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
903 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
904 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
905 kernelWidth = kernel->width;
906 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
907 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
908 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
909 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
910 if (clStatus != CL_SUCCESS)
912 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
917 /* launch the kernel */
922 gsize[0] = chunkSize*((inputImage->columns+chunkSize-1)/chunkSize);
923 gsize[1] = inputImage->rows;
924 wsize[0] = chunkSize;
927 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
928 if (clStatus != CL_SUCCESS)
930 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
933 clEnv->library->clFlush(queue);
938 /* need logic to decide this value */
942 imageColumns = inputImage->columns;
943 imageRows = inputImage->rows;
945 /* set the kernel arguments */
947 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
948 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
949 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
950 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
951 kernelWidth = kernel->width;
952 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
953 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
954 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
955 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *)NULL);
956 if (clStatus != CL_SUCCESS)
958 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
963 /* launch the kernel */
968 gsize[0] = inputImage->columns;
969 gsize[1] = chunkSize*((inputImage->rows+chunkSize-1)/chunkSize);
971 wsize[1] = chunkSize;
973 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
974 if (clStatus != CL_SUCCESS)
976 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
979 clEnv->library->clFlush(queue);
986 if (ALIGNED(filteredPixels,CLPixelPacket))
988 length = inputImage->columns * inputImage->rows;
989 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
993 length = inputImage->columns * inputImage->rows;
994 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
996 if (clStatus != CL_SUCCESS)
998 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1002 outputReady = MagickTrue;
1005 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1007 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
1008 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
1009 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
1010 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
1011 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
1012 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
1013 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1014 if (kernel!=NULL) DestroyKernelInfo(kernel);
1015 if (outputReady == MagickFalse)
1017 if (filteredImage != NULL)
1019 DestroyImage(filteredImage);
1020 filteredImage = NULL;
1023 return filteredImage;
1026 static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType channel, const double radius, const double sigma, ExceptionInfo *exception)
1028 MagickBooleanType outputReady;
1029 Image* filteredImage;
1034 const void *inputPixels;
1035 void *filteredPixels;
1036 cl_mem_flags mem_flags;
1039 cl_mem inputImageBuffer, tempImageBuffer, filteredImageBuffer, imageKernelBuffer;
1040 cl_kernel blurRowKernel, blurColumnKernel;
1041 cl_command_queue queue;
1044 float* kernelBufferPtr;
1045 MagickSizeType length;
1047 char geometry[MaxTextExtent];
1048 KernelInfo* kernel = NULL;
1049 unsigned int kernelWidth;
1050 unsigned int imageColumns, imageRows;
1055 filteredImage = NULL;
1056 inputImageBuffer = NULL;
1057 tempImageBuffer = NULL;
1058 filteredImageBuffer = NULL;
1059 imageKernelBuffer = NULL;
1060 blurRowKernel = NULL;
1061 blurColumnKernel = NULL;
1064 outputReady = MagickFalse;
1066 clEnv = GetDefaultOpenCLEnv();
1067 context = GetOpenCLContext(clEnv);
1068 queue = AcquireOpenCLCommandQueue(clEnv);
1070 /* Create and initialize OpenCL buffers. */
1073 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
1074 if (inputPixels == (const void *) NULL)
1076 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
1079 /* If the host pointer is aligned to the size of CLPixelPacket,
1080 then use the host buffer directly from the GPU; otherwise,
1081 create a buffer on the GPU and copy the data over */
1082 if (ALIGNED(inputPixels,CLPixelPacket))
1084 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1088 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1090 /* create a CL buffer from image pixel buffer */
1091 length = inputImage->columns * inputImage->rows;
1092 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1093 if (clStatus != CL_SUCCESS)
1095 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1102 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
1103 assert(filteredImage != NULL);
1104 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
1106 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1109 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
1110 if (filteredPixels == (void *) NULL)
1112 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1116 if (ALIGNED(filteredPixels,CLPixelPacket))
1118 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1119 hostPtr = filteredPixels;
1123 mem_flags = CL_MEM_WRITE_ONLY;
1126 /* create a CL buffer from image pixel buffer */
1127 length = inputImage->columns * inputImage->rows;
1128 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1129 if (clStatus != CL_SUCCESS)
1131 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1136 /* create processing kernel */
1138 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
1139 kernel=AcquireKernelInfo(geometry);
1140 if (kernel == (KernelInfo *) NULL)
1142 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
1146 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus);
1147 if (clStatus != CL_SUCCESS)
1149 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1152 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
1153 if (clStatus != CL_SUCCESS)
1155 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
1159 for (i = 0; i < kernel->width; i++)
1161 kernelBufferPtr[i] = (float) kernel->values[i];
1164 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
1165 if (clStatus != CL_SUCCESS)
1167 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1173 unsigned int offsetRows;
1176 /* create temp buffer */
1178 length = inputImage->columns * (inputImage->rows / 2 + 1 + (kernel->width-1) / 2);
1179 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
1180 if (clStatus != CL_SUCCESS)
1182 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1187 /* get the OpenCL kernels */
1189 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection");
1190 if (blurRowKernel == NULL)
1192 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1196 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumnSection");
1197 if (blurColumnKernel == NULL)
1199 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1204 for (sec = 0; sec < 2; sec++)
1207 /* need logic to decide this value */
1208 int chunkSize = 256;
1211 imageColumns = inputImage->columns;
1213 imageRows = inputImage->rows / 2 + (kernel->width-1) / 2;
1215 imageRows = (inputImage->rows - inputImage->rows / 2) + (kernel->width-1) / 2;
1217 offsetRows = sec * inputImage->rows / 2;
1219 kernelWidth = kernel->width;
1221 /* set the kernel arguments */
1223 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1224 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1225 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
1226 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1227 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1228 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1229 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1230 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
1231 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
1232 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec);
1233 if (clStatus != CL_SUCCESS)
1235 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1240 /* launch the kernel */
1245 gsize[0] = chunkSize*((imageColumns+chunkSize-1)/chunkSize);
1246 gsize[1] = imageRows;
1247 wsize[0] = chunkSize;
1250 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1251 if (clStatus != CL_SUCCESS)
1253 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1256 clEnv->library->clFlush(queue);
1261 /* need logic to decide this value */
1262 int chunkSize = 256;
1265 imageColumns = inputImage->columns;
1267 imageRows = inputImage->rows / 2;
1269 imageRows = (inputImage->rows - inputImage->rows / 2);
1271 offsetRows = sec * inputImage->rows / 2;
1273 kernelWidth = kernel->width;
1275 /* set the kernel arguments */
1277 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1278 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1279 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
1280 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1281 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1282 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1283 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1284 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *)NULL);
1285 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
1286 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&sec);
1287 if (clStatus != CL_SUCCESS)
1289 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1294 /* launch the kernel */
1299 gsize[0] = imageColumns;
1300 gsize[1] = chunkSize*((imageRows+chunkSize-1)/chunkSize);
1302 wsize[1] = chunkSize;
1304 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1305 if (clStatus != CL_SUCCESS)
1307 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1310 clEnv->library->clFlush(queue);
1318 if (ALIGNED(filteredPixels,CLPixelPacket))
1320 length = inputImage->columns * inputImage->rows;
1321 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1325 length = inputImage->columns * inputImage->rows;
1326 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1328 if (clStatus != CL_SUCCESS)
1330 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1334 outputReady = MagickTrue;
1337 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1339 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
1340 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
1341 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
1342 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
1343 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
1344 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
1345 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1346 if (kernel!=NULL) DestroyKernelInfo(kernel);
1347 if (outputReady == MagickFalse)
1349 if (filteredImage != NULL)
1351 DestroyImage(filteredImage);
1352 filteredImage = NULL;
1355 return filteredImage;
1359 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1363 % B l u r I m a g e w i t h O p e n C L %
1367 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1369 % BlurImage() blurs an image. We convolve the image with a Gaussian operator
1370 % of the given radius and standard deviation (sigma). For reasonable results,
1371 % the radius should be larger than sigma. Use a radius of 0 and BlurImage()
1372 % selects a suitable radius for you.
1374 % The format of the BlurImage method is:
1376 % Image *BlurImage(const Image *image,const double radius,
1377 % const double sigma,ExceptionInfo *exception)
1378 % Image *BlurImageChannel(const Image *image,const ChannelType channel,
1379 % const double radius,const double sigma,ExceptionInfo *exception)
1381 % A description of each parameter follows:
1383 % o image: the image.
1385 % o channel: the channel type.
1387 % o radius: the radius of the Gaussian, in pixels, not counting the center
1390 % o sigma: the standard deviation of the Gaussian, in pixels.
1392 % o exception: return any errors or warnings in this structure.
1397 Image* AccelerateBlurImage(const Image *image, const ChannelType channel, const double radius, const double sigma,ExceptionInfo *exception)
1399 MagickBooleanType status;
1400 Image* filteredImage = NULL;
1402 assert(image != NULL);
1403 assert(exception != (ExceptionInfo *) NULL);
1405 status = checkOpenCLEnvironment(exception);
1406 if (status == MagickFalse)
1409 status = checkAccelerateCondition(image, channel);
1410 if (status == MagickFalse)
1413 if (splitImage(image) && (image->rows / 2 > radius))
1414 filteredImage = ComputeBlurImageSection(image, channel, radius, sigma, exception);
1416 filteredImage = ComputeBlurImage(image, channel, radius, sigma, exception);
1418 return filteredImage;
1422 static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType channel, const double angle, ExceptionInfo *exception)
1425 MagickBooleanType outputReady;
1426 Image* filteredImage;
1430 size_t global_work_size[2];
1433 cl_mem_flags mem_flags;
1434 cl_mem inputImageBuffer, filteredImageBuffer, sinThetaBuffer, cosThetaBuffer;
1435 cl_kernel radialBlurKernel;
1436 cl_command_queue queue;
1438 const void *inputPixels;
1439 void *filteredPixels;
1443 MagickSizeType length;
1445 MagickPixelPacket bias;
1446 cl_float4 biasPixel;
1447 cl_float2 blurCenter;
1449 unsigned int cossin_theta_size;
1450 float offset, theta;
1454 outputReady = MagickFalse;
1456 filteredImage = NULL;
1457 inputImageBuffer = NULL;
1458 filteredImageBuffer = NULL;
1459 sinThetaBuffer = NULL;
1460 cosThetaBuffer = NULL;
1462 radialBlurKernel = NULL;
1465 clEnv = GetDefaultOpenCLEnv();
1466 context = GetOpenCLContext(clEnv);
1469 /* Create and initialize OpenCL buffers. */
1472 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
1473 if (inputPixels == (const void *) NULL)
1475 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
1479 /* If the host pointer is aligned to the size of CLPixelPacket,
1480 then use the host buffer directly from the GPU; otherwise,
1481 create a buffer on the GPU and copy the data over */
1482 if (ALIGNED(inputPixels,CLPixelPacket))
1484 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1488 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1490 /* create a CL buffer from image pixel buffer */
1491 length = inputImage->columns * inputImage->rows;
1492 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1493 if (clStatus != CL_SUCCESS)
1495 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1500 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
1501 assert(filteredImage != NULL);
1502 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
1504 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1507 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
1508 if (filteredPixels == (void *) NULL)
1510 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1514 if (ALIGNED(filteredPixels,CLPixelPacket))
1516 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1517 hostPtr = filteredPixels;
1521 mem_flags = CL_MEM_WRITE_ONLY;
1524 /* create a CL buffer from image pixel buffer */
1525 length = inputImage->columns * inputImage->rows;
1526 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1527 if (clStatus != CL_SUCCESS)
1529 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1533 blurCenter.s[0] = (float) (inputImage->columns-1)/2.0;
1534 blurCenter.s[1] = (float) (inputImage->rows-1)/2.0;
1535 blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]);
1536 cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL);
1538 /* create a buffer for sin_theta and cos_theta */
1539 sinThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
1540 if (clStatus != CL_SUCCESS)
1542 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1545 cosThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
1546 if (clStatus != CL_SUCCESS)
1548 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1553 queue = AcquireOpenCLCommandQueue(clEnv);
1554 sinThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
1555 if (clStatus != CL_SUCCESS)
1557 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
1561 cosThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
1562 if (clStatus != CL_SUCCESS)
1564 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
1568 theta=DegreesToRadians(angle)/(MagickRealType) (cossin_theta_size-1);
1569 offset=theta*(MagickRealType) (cossin_theta_size-1)/2.0;
1570 for (i=0; i < (ssize_t) cossin_theta_size; i++)
1572 cosThetaPtr[i]=(float)cos((double) (theta*i-offset));
1573 sinThetaPtr[i]=(float)sin((double) (theta*i-offset));
1576 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL);
1577 clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL);
1578 if (clStatus != CL_SUCCESS)
1580 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1584 /* get the OpenCL kernel */
1585 radialBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RadialBlur");
1586 if (radialBlurKernel == NULL)
1588 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1593 /* set the kernel arguments */
1595 clStatus=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1596 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1598 GetMagickPixelPacket(inputImage,&bias);
1599 biasPixel.s[0] = bias.red;
1600 biasPixel.s[1] = bias.green;
1601 biasPixel.s[2] = bias.blue;
1602 biasPixel.s[3] = bias.opacity;
1603 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_float4), &biasPixel);
1604 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(ChannelType), &channel);
1606 matte = (inputImage->matte != MagickFalse)?1:0;
1607 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &matte);
1609 clStatus=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_float2), &blurCenter);
1611 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer);
1612 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer);
1613 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size);
1614 if (clStatus != CL_SUCCESS)
1616 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1621 global_work_size[0] = inputImage->columns;
1622 global_work_size[1] = inputImage->rows;
1623 /* launch the kernel */
1624 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, radialBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
1625 if (clStatus != CL_SUCCESS)
1627 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1630 clEnv->library->clFlush(queue);
1632 if (ALIGNED(filteredPixels,CLPixelPacket))
1634 length = inputImage->columns * inputImage->rows;
1635 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1639 length = inputImage->columns * inputImage->rows;
1640 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1642 if (clStatus != CL_SUCCESS)
1644 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1647 outputReady = MagickTrue;
1650 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1652 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
1653 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
1654 if (sinThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(sinThetaBuffer);
1655 if (cosThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(cosThetaBuffer);
1656 if (radialBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, radialBlurKernel);
1657 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1658 if (outputReady == MagickFalse)
1660 if (filteredImage != NULL)
1662 DestroyImage(filteredImage);
1663 filteredImage = NULL;
1666 return filteredImage;
1670 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1674 % 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 %
1678 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1680 % RadialBlurImage() applies a radial blur to the image.
1682 % Andrew Protano contributed this effect.
1684 % The format of the RadialBlurImage method is:
1686 % Image *RadialBlurImage(const Image *image,const double angle,
1687 % ExceptionInfo *exception)
1688 % Image *RadialBlurImageChannel(const Image *image,const ChannelType channel,
1689 % const double angle,ExceptionInfo *exception)
1691 % A description of each parameter follows:
1693 % o image: the image.
1695 % o channel: the channel type.
1697 % o angle: the angle of the radial blur.
1699 % o exception: return any errors or warnings in this structure.
1704 Image* AccelerateRadialBlurImage(const Image *image, const ChannelType channel, const double angle, ExceptionInfo *exception)
1706 MagickBooleanType status;
1707 Image* filteredImage;
1710 assert(image != NULL);
1711 assert(exception != NULL);
1713 status = checkOpenCLEnvironment(exception);
1714 if (status == MagickFalse)
1717 status = checkAccelerateCondition(image, channel);
1718 if (status == MagickFalse)
1721 filteredImage = ComputeRadialBlurImage(image, channel, angle, exception);
1722 return filteredImage;
1727 static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType channel,const double radius,const double sigma,
1728 const double gain,const double threshold,ExceptionInfo *exception)
1730 MagickBooleanType outputReady = MagickFalse;
1731 Image* filteredImage = NULL;
1732 MagickCLEnv clEnv = NULL;
1736 const void *inputPixels;
1737 void *filteredPixels;
1738 cl_mem_flags mem_flags;
1740 KernelInfo *kernel = NULL;
1741 char geometry[MaxTextExtent];
1743 cl_context context = NULL;
1744 cl_mem inputImageBuffer = NULL;
1745 cl_mem filteredImageBuffer = NULL;
1746 cl_mem tempImageBuffer = NULL;
1747 cl_mem imageKernelBuffer = NULL;
1748 cl_kernel blurRowKernel = NULL;
1749 cl_kernel unsharpMaskBlurColumnKernel = NULL;
1750 cl_command_queue queue = NULL;
1753 float* kernelBufferPtr;
1754 MagickSizeType length;
1755 unsigned int kernelWidth;
1758 unsigned int imageColumns, imageRows;
1762 clEnv = GetDefaultOpenCLEnv();
1763 context = GetOpenCLContext(clEnv);
1764 queue = AcquireOpenCLCommandQueue(clEnv);
1766 /* Create and initialize OpenCL buffers. */
1769 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
1770 if (inputPixels == (const void *) NULL)
1772 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
1776 /* If the host pointer is aligned to the size of CLPixelPacket,
1777 then use the host buffer directly from the GPU; otherwise,
1778 create a buffer on the GPU and copy the data over */
1779 if (ALIGNED(inputPixels,CLPixelPacket))
1781 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1785 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1787 /* create a CL buffer from image pixel buffer */
1788 length = inputImage->columns * inputImage->rows;
1789 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
1790 if (clStatus != CL_SUCCESS)
1792 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1799 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
1800 assert(filteredImage != NULL);
1801 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
1803 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
1806 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
1807 if (filteredPixels == (void *) NULL)
1809 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
1813 if (ALIGNED(filteredPixels,CLPixelPacket))
1815 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1816 hostPtr = filteredPixels;
1820 mem_flags = CL_MEM_WRITE_ONLY;
1824 /* create a CL buffer from image pixel buffer */
1825 length = inputImage->columns * inputImage->rows;
1826 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
1827 if (clStatus != CL_SUCCESS)
1829 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1834 /* create the blur kernel */
1836 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
1837 kernel=AcquireKernelInfo(geometry);
1838 if (kernel == (KernelInfo *) NULL)
1840 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
1844 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
1845 if (clStatus != CL_SUCCESS)
1847 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1852 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
1853 if (clStatus != CL_SUCCESS)
1855 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
1858 for (i = 0; i < kernel->width; i++)
1860 kernelBufferPtr[i] = (float) kernel->values[i];
1862 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
1863 if (clStatus != CL_SUCCESS)
1865 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1871 /* create temp buffer */
1873 length = inputImage->columns * inputImage->rows;
1874 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
1875 if (clStatus != CL_SUCCESS)
1877 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1882 /* get the opencl kernel */
1884 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
1885 if (blurRowKernel == NULL)
1887 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1891 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn");
1892 if (unsharpMaskBlurColumnKernel == NULL)
1894 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1902 imageColumns = inputImage->columns;
1903 imageRows = inputImage->rows;
1905 kernelWidth = kernel->width;
1907 /* set the kernel arguments */
1909 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1910 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1911 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
1912 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1913 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1914 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1915 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1916 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
1917 if (clStatus != CL_SUCCESS)
1919 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1924 /* launch the kernel */
1929 gsize[0] = chunkSize*((inputImage->columns+chunkSize-1)/chunkSize);
1930 gsize[1] = inputImage->rows;
1931 wsize[0] = chunkSize;
1934 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1935 if (clStatus != CL_SUCCESS)
1937 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1940 clEnv->library->clFlush(queue);
1946 imageColumns = inputImage->columns;
1947 imageRows = inputImage->rows;
1948 kernelWidth = kernel->width;
1949 fGain = (float)gain;
1950 fThreshold = (float)threshold;
1953 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
1954 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
1955 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1956 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
1957 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
1958 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
1959 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
1960 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
1961 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
1962 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
1963 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
1964 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
1966 if (clStatus != CL_SUCCESS)
1968 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1973 /* launch the kernel */
1978 gsize[0] = inputImage->columns;
1979 gsize[1] = chunkSize*((inputImage->rows+chunkSize-1)/chunkSize);
1981 wsize[1] = chunkSize;
1983 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
1984 if (clStatus != CL_SUCCESS)
1986 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1989 clEnv->library->clFlush(queue);
1995 if (ALIGNED(filteredPixels,CLPixelPacket))
1997 length = inputImage->columns * inputImage->rows;
1998 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2002 length = inputImage->columns * inputImage->rows;
2003 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2005 if (clStatus != CL_SUCCESS)
2007 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2011 outputReady = MagickTrue;
2014 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2016 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
2017 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
2018 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
2019 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
2020 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
2021 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
2022 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
2023 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2024 if (outputReady == MagickFalse)
2026 if (filteredImage != NULL)
2028 DestroyImage(filteredImage);
2029 filteredImage = NULL;
2032 return filteredImage;
2036 static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const ChannelType channel,const double radius,const double sigma,
2037 const double gain,const double threshold,ExceptionInfo *exception)
2039 MagickBooleanType outputReady = MagickFalse;
2040 Image* filteredImage = NULL;
2041 MagickCLEnv clEnv = NULL;
2045 const void *inputPixels;
2046 void *filteredPixels;
2047 cl_mem_flags mem_flags;
2049 KernelInfo *kernel = NULL;
2050 char geometry[MaxTextExtent];
2052 cl_context context = NULL;
2053 cl_mem inputImageBuffer = NULL;
2054 cl_mem filteredImageBuffer = NULL;
2055 cl_mem tempImageBuffer = NULL;
2056 cl_mem imageKernelBuffer = NULL;
2057 cl_kernel blurRowKernel = NULL;
2058 cl_kernel unsharpMaskBlurColumnKernel = NULL;
2059 cl_command_queue queue = NULL;
2062 float* kernelBufferPtr;
2063 MagickSizeType length;
2064 unsigned int kernelWidth;
2067 unsigned int imageColumns, imageRows;
2071 clEnv = GetDefaultOpenCLEnv();
2072 context = GetOpenCLContext(clEnv);
2073 queue = AcquireOpenCLCommandQueue(clEnv);
2075 /* Create and initialize OpenCL buffers. */
2078 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
2079 if (inputPixels == (const void *) NULL)
2081 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
2085 /* If the host pointer is aligned to the size of CLPixelPacket,
2086 then use the host buffer directly from the GPU; otherwise,
2087 create a buffer on the GPU and copy the data over */
2088 if (ALIGNED(inputPixels,CLPixelPacket))
2090 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2094 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2096 /* create a CL buffer from image pixel buffer */
2097 length = inputImage->columns * inputImage->rows;
2098 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2099 if (clStatus != CL_SUCCESS)
2101 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2108 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
2109 assert(filteredImage != NULL);
2110 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
2112 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2115 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
2116 if (filteredPixels == (void *) NULL)
2118 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2122 if (ALIGNED(filteredPixels,CLPixelPacket))
2124 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2125 hostPtr = filteredPixels;
2129 mem_flags = CL_MEM_WRITE_ONLY;
2133 /* create a CL buffer from image pixel buffer */
2134 length = inputImage->columns * inputImage->rows;
2135 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2136 if (clStatus != CL_SUCCESS)
2138 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2143 /* create the blur kernel */
2145 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
2146 kernel=AcquireKernelInfo(geometry);
2147 if (kernel == (KernelInfo *) NULL)
2149 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
2153 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
2154 if (clStatus != CL_SUCCESS)
2156 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2161 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
2162 if (clStatus != CL_SUCCESS)
2164 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
2167 for (i = 0; i < kernel->width; i++)
2169 kernelBufferPtr[i] = (float) kernel->values[i];
2171 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
2172 if (clStatus != CL_SUCCESS)
2174 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
2180 unsigned int offsetRows;
2183 /* create temp buffer */
2185 length = inputImage->columns * (inputImage->rows / 2 + 1 + (kernel->width-1) / 2);
2186 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
2187 if (clStatus != CL_SUCCESS)
2189 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2194 /* get the opencl kernel */
2196 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection");
2197 if (blurRowKernel == NULL)
2199 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2203 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumnSection");
2204 if (unsharpMaskBlurColumnKernel == NULL)
2206 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2211 for (sec = 0; sec < 2; sec++)
2216 imageColumns = inputImage->columns;
2218 imageRows = inputImage->rows / 2 + (kernel->width-1) / 2;
2220 imageRows = (inputImage->rows - inputImage->rows / 2) + (kernel->width-1) / 2;
2222 offsetRows = sec * inputImage->rows / 2;
2224 kernelWidth = kernel->width;
2226 /* set the kernel arguments */
2228 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
2229 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2230 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
2231 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2232 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2233 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2234 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2235 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL);
2236 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
2237 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec);
2238 if (clStatus != CL_SUCCESS)
2240 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2244 /* launch the kernel */
2249 gsize[0] = chunkSize*((imageColumns+chunkSize-1)/chunkSize);
2250 gsize[1] = imageRows;
2251 wsize[0] = chunkSize;
2254 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2255 if (clStatus != CL_SUCCESS)
2257 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2260 clEnv->library->clFlush(queue);
2267 imageColumns = inputImage->columns;
2269 imageRows = inputImage->rows / 2;
2271 imageRows = (inputImage->rows - inputImage->rows / 2);
2273 offsetRows = sec * inputImage->rows / 2;
2275 kernelWidth = kernel->width;
2277 fGain = (float)gain;
2278 fThreshold = (float)threshold;
2281 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
2282 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
2283 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
2284 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
2285 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
2286 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
2287 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
2288 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
2289 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
2290 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
2291 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
2292 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
2293 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&offsetRows);
2294 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&sec);
2296 if (clStatus != CL_SUCCESS)
2298 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2303 /* launch the kernel */
2308 gsize[0] = imageColumns;
2309 gsize[1] = chunkSize*((imageRows+chunkSize-1)/chunkSize);
2311 wsize[1] = chunkSize;
2313 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
2314 if (clStatus != CL_SUCCESS)
2316 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2319 clEnv->library->clFlush(queue);
2325 if (ALIGNED(filteredPixels,CLPixelPacket))
2327 length = inputImage->columns * inputImage->rows;
2328 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2332 length = inputImage->columns * inputImage->rows;
2333 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2335 if (clStatus != CL_SUCCESS)
2337 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2341 outputReady = MagickTrue;
2344 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2346 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
2347 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
2348 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
2349 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
2350 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
2351 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
2352 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
2353 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2354 if (outputReady == MagickFalse)
2356 if (filteredImage != NULL)
2358 DestroyImage(filteredImage);
2359 filteredImage = NULL;
2362 return filteredImage;
2367 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2371 % 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 %
2375 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2377 % UnsharpMaskImage() sharpens one or more image channels. We convolve the
2378 % image with a Gaussian operator of the given radius and standard deviation
2379 % (sigma). For reasonable results, radius should be larger than sigma. Use a
2380 % radius of 0 and UnsharpMaskImage() selects a suitable radius for you.
2382 % The format of the UnsharpMaskImage method is:
2384 % Image *UnsharpMaskImage(const Image *image,const double radius,
2385 % const double sigma,const double amount,const double threshold,
2386 % ExceptionInfo *exception)
2387 % Image *UnsharpMaskImageChannel(const Image *image,
2388 % const ChannelType channel,const double radius,const double sigma,
2389 % const double gain,const double threshold,ExceptionInfo *exception)
2391 % A description of each parameter follows:
2393 % o image: the image.
2395 % o channel: the channel type.
2397 % o radius: the radius of the Gaussian, in pixels, not counting the center
2400 % o sigma: the standard deviation of the Gaussian, in pixels.
2402 % o gain: the percentage of the difference between the original and the
2403 % blur image that is added back into the original.
2405 % o threshold: the threshold in pixels needed to apply the diffence gain.
2407 % o exception: return any errors or warnings in this structure.
2413 Image* AccelerateUnsharpMaskImage(const Image *image, const ChannelType channel,const double radius,const double sigma,
2414 const double gain,const double threshold,ExceptionInfo *exception)
2416 MagickBooleanType status;
2417 Image* filteredImage;
2420 assert(image != NULL);
2421 assert(exception != NULL);
2423 status = checkOpenCLEnvironment(exception);
2424 if (status == MagickFalse)
2427 status = checkAccelerateCondition(image, channel);
2428 if (status == MagickFalse)
2431 if (splitImage(image) && (image->rows / 2 > radius))
2432 filteredImage = ComputeUnsharpMaskImageSection(image,channel,radius,sigma,gain,threshold,exception);
2434 filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
2435 return filteredImage;
2439 static MagickBooleanType resizeHorizontalFilter(cl_mem inputImage
2440 , const unsigned int inputImageColumns, const unsigned int inputImageRows, const unsigned int matte
2441 , cl_mem resizedImage, const unsigned int resizedColumns, const unsigned int resizedRows
2442 , const ResizeFilter* resizeFilter, cl_mem resizeFilterCubicCoefficients, const float xFactor
2443 , MagickCLEnv clEnv, cl_command_queue queue, ExceptionInfo *exception)
2445 MagickBooleanType status = MagickFalse;
2447 float scale, support;
2449 cl_kernel horizontalKernel = NULL;
2451 size_t global_work_size[2];
2452 size_t local_work_size[2];
2453 int resizeFilterType, resizeWindowType;
2454 float resizeFilterScale, resizeFilterSupport, resizeFilterWindowSupport, resizeFilterBlur;
2455 size_t totalLocalMemorySize;
2456 size_t imageCacheLocalMemorySize, pixelAccumulatorLocalMemorySize
2457 , weightAccumulatorLocalMemorySize, gammaAccumulatorLocalMemorySize;
2458 size_t deviceLocalMemorySize;
2459 int cacheRangeStart, cacheRangeEnd, numCachedPixels;
2461 const unsigned int workgroupSize = 256;
2462 unsigned int pixelPerWorkgroup;
2463 unsigned int chunkSize;
2466 Apply filter to resize vertically from image to resize image.
2468 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
2469 support=scale*GetResizeFilterSupport(resizeFilter);
2473 Support too small even for nearest neighbour: Reduce to point
2476 support=(MagickRealType) 0.5;
2479 scale=PerceptibleReciprocal(scale);
2481 if (resizedColumns < workgroupSize)
2484 pixelPerWorkgroup = 32;
2488 chunkSize = workgroupSize;
2489 pixelPerWorkgroup = workgroupSize;
2492 /* get the local memory size supported by the device */
2493 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
2495 DisableMSCWarning(4127)
2499 /* calculate the local memory size needed per workgroup */
2500 cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
2501 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
2502 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
2503 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
2504 totalLocalMemorySize = imageCacheLocalMemorySize;
2506 /* local size for the pixel accumulator */
2507 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
2508 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
2510 /* local memory size for the weight accumulator */
2511 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2512 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
2514 /* local memory size for the gamma accumulator */
2516 gammaAccumulatorLocalMemorySize = sizeof(float);
2518 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2519 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
2521 if (totalLocalMemorySize <= deviceLocalMemorySize)
2525 pixelPerWorkgroup = pixelPerWorkgroup/2;
2526 chunkSize = chunkSize/2;
2527 if (pixelPerWorkgroup == 0
2530 /* quit, fallback to CPU */
2536 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
2537 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
2540 if (resizeFilterType == SincFastWeightingFunction
2541 && resizeWindowType == SincFastWeightingFunction)
2543 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilterSinc");
2547 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter");
2549 if (horizontalKernel == NULL)
2551 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2556 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&inputImage);
2557 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageColumns);
2558 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageRows);
2559 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
2560 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
2561 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
2563 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
2564 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
2566 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
2567 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
2568 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
2570 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
2571 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
2573 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
2574 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
2576 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
2577 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
2579 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
2580 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
2583 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
2584 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
2585 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
2586 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
2589 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
2590 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
2591 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
2593 if (clStatus != CL_SUCCESS)
2595 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2599 global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
2600 global_work_size[1] = resizedRows;
2602 local_work_size[0] = workgroupSize;
2603 local_work_size[1] = 1;
2604 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
2605 if (clStatus != CL_SUCCESS)
2607 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2610 clEnv->library->clFlush(queue);
2611 status = MagickTrue;
2615 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2617 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
2623 static MagickBooleanType resizeVerticalFilter(cl_mem inputImage
2624 , const unsigned int inputImageColumns, const unsigned int inputImageRows, const unsigned int matte
2625 , cl_mem resizedImage, const unsigned int resizedColumns, const unsigned int resizedRows
2626 , const ResizeFilter* resizeFilter, cl_mem resizeFilterCubicCoefficients, const float yFactor
2627 , MagickCLEnv clEnv, cl_command_queue queue, ExceptionInfo *exception)
2629 MagickBooleanType status = MagickFalse;
2631 float scale, support;
2633 cl_kernel horizontalKernel = NULL;
2635 size_t global_work_size[2];
2636 size_t local_work_size[2];
2637 int resizeFilterType, resizeWindowType;
2638 float resizeFilterScale, resizeFilterSupport, resizeFilterWindowSupport, resizeFilterBlur;
2639 size_t totalLocalMemorySize;
2640 size_t imageCacheLocalMemorySize, pixelAccumulatorLocalMemorySize
2641 , weightAccumulatorLocalMemorySize, gammaAccumulatorLocalMemorySize;
2642 size_t deviceLocalMemorySize;
2643 int cacheRangeStart, cacheRangeEnd, numCachedPixels;
2645 const unsigned int workgroupSize = 256;
2646 unsigned int pixelPerWorkgroup;
2647 unsigned int chunkSize;
2650 Apply filter to resize vertically from image to resize image.
2652 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
2653 support=scale*GetResizeFilterSupport(resizeFilter);
2657 Support too small even for nearest neighbour: Reduce to point
2660 support=(MagickRealType) 0.5;
2663 scale=PerceptibleReciprocal(scale);
2665 if (resizedRows < workgroupSize)
2668 pixelPerWorkgroup = 32;
2672 chunkSize = workgroupSize;
2673 pixelPerWorkgroup = workgroupSize;
2676 /* get the local memory size supported by the device */
2677 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
2679 DisableMSCWarning(4127)
2683 /* calculate the local memory size needed per workgroup */
2684 cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
2685 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
2686 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
2687 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
2688 totalLocalMemorySize = imageCacheLocalMemorySize;
2690 /* local size for the pixel accumulator */
2691 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
2692 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
2694 /* local memory size for the weight accumulator */
2695 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2696 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
2698 /* local memory size for the gamma accumulator */
2700 gammaAccumulatorLocalMemorySize = sizeof(float);
2702 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
2703 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
2705 if (totalLocalMemorySize <= deviceLocalMemorySize)
2709 pixelPerWorkgroup = pixelPerWorkgroup/2;
2710 chunkSize = chunkSize/2;
2711 if (pixelPerWorkgroup == 0
2714 /* quit, fallback to CPU */
2720 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
2721 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
2723 if (resizeFilterType == SincFastWeightingFunction
2724 && resizeWindowType == SincFastWeightingFunction)
2725 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilterSinc");
2727 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter");
2729 if (horizontalKernel == NULL)
2731 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2736 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&inputImage);
2737 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageColumns);
2738 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&inputImageRows);
2739 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
2740 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&yFactor);
2741 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
2743 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
2744 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
2746 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
2747 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
2748 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
2750 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
2751 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
2753 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
2754 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
2756 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
2757 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
2759 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
2760 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
2763 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
2764 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
2765 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
2766 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
2769 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
2770 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
2771 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
2773 if (clStatus != CL_SUCCESS)
2775 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2779 global_work_size[0] = resizedColumns;
2780 global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
2782 local_work_size[0] = 1;
2783 local_work_size[1] = workgroupSize;
2784 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
2785 if (clStatus != CL_SUCCESS)
2787 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2790 clEnv->library->clFlush(queue);
2791 status = MagickTrue;
2795 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2797 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
2804 static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedColumns, const size_t resizedRows
2805 , const ResizeFilter* resizeFilter, ExceptionInfo *exception)
2808 MagickBooleanType outputReady = MagickFalse;
2809 Image* filteredImage = NULL;
2810 MagickCLEnv clEnv = NULL;
2813 MagickBooleanType status;
2814 const void *inputPixels;
2815 void* filteredPixels;
2817 const MagickRealType* resizeFilterCoefficient;
2818 float* mappedCoefficientBuffer;
2819 float xFactor, yFactor;
2820 MagickSizeType length;
2822 cl_mem_flags mem_flags;
2823 cl_context context = NULL;
2824 cl_mem inputImageBuffer = NULL;
2825 cl_mem tempImageBuffer = NULL;
2826 cl_mem filteredImageBuffer = NULL;
2827 cl_mem cubicCoefficientsBuffer = NULL;
2828 cl_command_queue queue = NULL;
2832 clEnv = GetDefaultOpenCLEnv();
2833 context = GetOpenCLContext(clEnv);
2835 /* Create and initialize OpenCL buffers. */
2837 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
2838 if (inputPixels == (const void *) NULL)
2840 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
2844 /* If the host pointer is aligned to the size of CLPixelPacket,
2845 then use the host buffer directly from the GPU; otherwise,
2846 create a buffer on the GPU and copy the data over */
2847 if (ALIGNED(inputPixels,CLPixelPacket))
2849 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2853 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2855 /* create a CL buffer from image pixel buffer */
2856 length = inputImage->columns * inputImage->rows;
2857 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
2858 if (clStatus != CL_SUCCESS)
2860 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2864 cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus);
2865 if (clStatus != CL_SUCCESS)
2867 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2870 queue = AcquireOpenCLCommandQueue(clEnv);
2871 mappedCoefficientBuffer = (float*)clEnv->library->clEnqueueMapBuffer(queue, cubicCoefficientsBuffer, CL_TRUE, CL_MAP_WRITE, 0, 7 * sizeof(float)
2872 , 0, NULL, NULL, &clStatus);
2873 if (clStatus != CL_SUCCESS)
2875 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
2878 resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter);
2879 for (i = 0; i < 7; i++)
2881 mappedCoefficientBuffer[i] = (float) resizeFilterCoefficient[i];
2883 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL);
2884 if (clStatus != CL_SUCCESS)
2886 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
2890 filteredImage = CloneImage(inputImage,resizedColumns,resizedRows,MagickTrue,exception);
2891 if (filteredImage == NULL)
2894 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
2896 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
2899 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
2900 if (filteredPixels == (void *) NULL)
2902 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
2906 if (ALIGNED(filteredPixels,CLPixelPacket))
2908 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2909 hostPtr = filteredPixels;
2913 mem_flags = CL_MEM_WRITE_ONLY;
2917 /* create a CL buffer from image pixel buffer */
2918 length = filteredImage->columns * filteredImage->rows;
2919 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
2920 if (clStatus != CL_SUCCESS)
2922 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2926 xFactor=(float) resizedColumns/(float) inputImage->columns;
2927 yFactor=(float) resizedRows/(float) inputImage->rows;
2928 if (xFactor > yFactor)
2931 length = resizedColumns*inputImage->rows;
2932 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
2933 if (clStatus != CL_SUCCESS)
2935 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2939 status = resizeHorizontalFilter(inputImageBuffer, inputImage->columns, inputImage->rows, (inputImage->matte != MagickFalse)?1:0
2940 , tempImageBuffer, resizedColumns, inputImage->rows
2941 , resizeFilter, cubicCoefficientsBuffer
2942 , xFactor, clEnv, queue, exception);
2943 if (status != MagickTrue)
2946 status = resizeVerticalFilter(tempImageBuffer, resizedColumns, inputImage->rows, (inputImage->matte != MagickFalse)?1:0
2947 , filteredImageBuffer, resizedColumns, resizedRows
2948 , resizeFilter, cubicCoefficientsBuffer
2949 , yFactor, clEnv, queue, exception);
2950 if (status != MagickTrue)
2955 length = inputImage->columns*resizedRows;
2956 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
2957 if (clStatus != CL_SUCCESS)
2959 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2963 status = resizeVerticalFilter(inputImageBuffer, inputImage->columns, inputImage->rows, (inputImage->matte != MagickFalse)?1:0
2964 , tempImageBuffer, inputImage->columns, resizedRows
2965 , resizeFilter, cubicCoefficientsBuffer
2966 , yFactor, clEnv, queue, exception);
2967 if (status != MagickTrue)
2970 status = resizeHorizontalFilter(tempImageBuffer, inputImage->columns, resizedRows, (inputImage->matte != MagickFalse)?1:0
2971 , filteredImageBuffer, resizedColumns, resizedRows
2972 , resizeFilter, cubicCoefficientsBuffer
2973 , xFactor, clEnv, queue, exception);
2974 if (status != MagickTrue)
2977 length = resizedColumns*resizedRows;
2978 if (ALIGNED(filteredPixels,CLPixelPacket))
2980 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2984 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2986 if (clStatus != CL_SUCCESS)
2988 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2991 outputReady = MagickTrue;
2994 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2996 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
2997 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
2998 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
2999 if (cubicCoefficientsBuffer!=NULL) clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer);
3000 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3001 if (outputReady == MagickFalse)
3003 if (filteredImage != NULL)
3005 DestroyImage(filteredImage);
3006 filteredImage = NULL;
3010 return filteredImage;
3013 const ResizeWeightingFunctionType supportedResizeWeighting[] =
3015 BoxWeightingFunction
3016 ,TriangleWeightingFunction
3017 ,HanningWeightingFunction
3018 ,HammingWeightingFunction
3019 ,BlackmanWeightingFunction
3020 ,CubicBCWeightingFunction
3021 ,SincWeightingFunction
3022 ,SincFastWeightingFunction
3023 ,LastWeightingFunction
3026 static MagickBooleanType gpuSupportedResizeWeighting(ResizeWeightingFunctionType f)
3028 MagickBooleanType supported = MagickFalse;
3032 if (supportedResizeWeighting[i] == LastWeightingFunction)
3034 if (supportedResizeWeighting[i] == f)
3036 supported = MagickTrue;
3045 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3049 % A c c e l e r a t e R e s i z e I m a g e %
3053 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3055 % AccelerateResizeImage() is an OpenCL implementation of ResizeImage()
3057 % AccelerateResizeImage() scales an image to the desired dimensions, using the given
3058 % filter (see AcquireFilterInfo()).
3060 % If an undefined filter is given the filter defaults to Mitchell for a
3061 % colormapped image, a image with a matte channel, or if the image is
3062 % enlarged. Otherwise the filter defaults to a Lanczos.
3064 % AccelerateResizeImage() was inspired by Paul Heckbert's "zoom" program.
3066 % The format of the AccelerateResizeImage method is:
3068 % Image *ResizeImage(Image *image,const size_t columns,
3069 % const size_t rows, const ResizeFilter* filter,
3070 % ExceptionInfo *exception)
3072 % A description of each parameter follows:
3074 % o image: the image.
3076 % o columns: the number of columns in the scaled image.
3078 % o rows: the number of rows in the scaled image.
3080 % o filter: Image filter to use.
3082 % o exception: return any errors or warnings in this structure.
3087 Image* AccelerateResizeImage(const Image* image, const size_t resizedColumns, const size_t resizedRows
3088 , const ResizeFilter* resizeFilter, ExceptionInfo *exception)
3090 MagickBooleanType status;
3091 Image* filteredImage;
3093 assert(image != NULL);
3094 assert(resizeFilter != NULL);
3096 status = checkOpenCLEnvironment(exception);
3097 if (status == MagickFalse)
3100 status = checkAccelerateCondition(image, AllChannels);
3101 if (status == MagickFalse)
3104 if (gpuSupportedResizeWeighting(GetResizeFilterWeightingType(resizeFilter)) == MagickFalse
3105 || gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(resizeFilter)) == MagickFalse)
3108 filteredImage = ComputeResizeImage(image,resizedColumns,resizedRows,resizeFilter,exception);
3109 return filteredImage;
3114 static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBooleanType sharpen, ExceptionInfo *exception)
3116 MagickBooleanType outputReady = MagickFalse;
3117 MagickCLEnv clEnv = NULL;
3120 size_t global_work_size[2];
3122 void *inputPixels = NULL;
3123 MagickSizeType length;
3124 unsigned int uSharpen;
3127 cl_mem_flags mem_flags;
3128 cl_context context = NULL;
3129 cl_mem inputImageBuffer = NULL;
3130 cl_kernel filterKernel = NULL;
3131 cl_command_queue queue = NULL;
3133 clEnv = GetDefaultOpenCLEnv();
3134 context = GetOpenCLContext(clEnv);
3136 /* Create and initialize OpenCL buffers. */
3137 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3138 if (inputPixels == (void *) NULL)
3140 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3144 /* If the host pointer is aligned to the size of CLPixelPacket,
3145 then use the host buffer directly from the GPU; otherwise,
3146 create a buffer on the GPU and copy the data over */
3147 if (ALIGNED(inputPixels,CLPixelPacket))
3149 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3153 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3155 /* create a CL buffer from image pixel buffer */
3156 length = inputImage->columns * inputImage->rows;
3157 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3158 if (clStatus != CL_SUCCESS)
3160 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3164 filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast");
3165 if (filterKernel == NULL)
3167 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3172 clStatus=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3174 uSharpen = (sharpen == MagickFalse)?0:1;
3175 clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen);
3176 if (clStatus != CL_SUCCESS)
3178 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3182 global_work_size[0] = inputImage->columns;
3183 global_work_size[1] = inputImage->rows;
3184 /* launch the kernel */
3185 queue = AcquireOpenCLCommandQueue(clEnv);
3186 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3187 if (clStatus != CL_SUCCESS)
3189 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3192 clEnv->library->clFlush(queue);
3194 if (ALIGNED(inputPixels,CLPixelPacket))
3196 length = inputImage->columns * inputImage->rows;
3197 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3201 length = inputImage->columns * inputImage->rows;
3202 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3204 if (clStatus != CL_SUCCESS)
3206 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3209 outputReady = MagickTrue;
3212 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3214 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
3215 if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel);
3216 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3221 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3225 % C o n t r a s t I m a g e w i t h O p e n C L %
3229 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3231 % ContrastImage() enhances the intensity differences between the lighter and
3232 % darker elements of the image. Set sharpen to a MagickTrue to increase the
3233 % image contrast otherwise the contrast is reduced.
3235 % The format of the ContrastImage method is:
3237 % MagickBooleanType ContrastImage(Image *image,
3238 % const MagickBooleanType sharpen)
3240 % A description of each parameter follows:
3242 % o image: the image.
3244 % o sharpen: Increase or decrease image contrast.
3249 MagickBooleanType AccelerateContrastImage(Image* image, const MagickBooleanType sharpen, ExceptionInfo* exception)
3251 MagickBooleanType status;
3253 assert(image != NULL);
3254 assert(exception != NULL);
3256 status = checkOpenCLEnvironment(exception);
3257 if (status == MagickFalse)
3260 status = checkAccelerateCondition(image, AllChannels);
3261 if (status == MagickFalse)
3264 status = ComputeContrastImage(image,sharpen,exception);
3270 MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, double percent_hue, double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
3282 MagickBooleanType outputReady;
3288 MagickSizeType length;
3291 cl_command_queue queue;
3292 cl_kernel modulateKernel;
3294 cl_mem inputImageBuffer;
3295 cl_mem_flags mem_flags;
3299 Image * inputImage = image;
3302 inputImageBuffer = NULL;
3303 modulateKernel = NULL;
3305 assert(inputImage != (Image *) NULL);
3306 assert(inputImage->signature == MagickSignature);
3307 if (inputImage->debug != MagickFalse)
3308 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",inputImage->filename);
3311 * initialize opencl env
3313 clEnv = GetDefaultOpenCLEnv();
3314 context = GetOpenCLContext(clEnv);
3315 queue = AcquireOpenCLCommandQueue(clEnv);
3317 outputReady = MagickFalse;
3319 /* Create and initialize OpenCL buffers.
3320 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
3321 assume this will get a writable image
3323 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3324 if (inputPixels == (void *) NULL)
3326 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3330 /* If the host pointer is aligned to the size of CLPixelPacket,
3331 then use the host buffer directly from the GPU; otherwise,
3332 create a buffer on the GPU and copy the data over
3334 if (ALIGNED(inputPixels,CLPixelPacket))
3336 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3340 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3342 /* create a CL buffer from image pixel buffer */
3343 length = inputImage->columns * inputImage->rows;
3344 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3345 if (clStatus != CL_SUCCESS)
3347 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3351 modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate");
3352 if (modulateKernel == NULL)
3354 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3358 bright=percent_brightness;
3360 saturation=percent_saturation;
3364 clStatus=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3365 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&bright);
3366 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&hue);
3367 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&saturation);
3368 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color);
3369 if (clStatus != CL_SUCCESS)
3371 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3372 printf("no kernel\n");
3377 size_t global_work_size[2];
3378 global_work_size[0] = inputImage->columns;
3379 global_work_size[1] = inputImage->rows;
3380 /* launch the kernel */
3381 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3382 if (clStatus != CL_SUCCESS)
3384 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3387 clEnv->library->clFlush(queue);
3390 if (ALIGNED(inputPixels,CLPixelPacket))
3392 length = inputImage->columns * inputImage->rows;
3393 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3397 length = inputImage->columns * inputImage->rows;
3398 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3400 if (clStatus != CL_SUCCESS)
3402 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3406 outputReady = MagickTrue;
3409 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3412 //ReleasePixelCachePixels();
3416 if (inputImageBuffer!=NULL)
3417 clEnv->library->clReleaseMemObject(inputImageBuffer);
3418 if (modulateKernel!=NULL)
3419 RelinquishOpenCLKernel(clEnv, modulateKernel);
3421 RelinquishOpenCLCommandQueue(clEnv, queue);
3428 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3432 % M o d u l a t e I m a g e w i t h O p e n C L %
3436 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3438 % ModulateImage() lets you control the brightness, saturation, and hue
3439 % of an image. Modulate represents the brightness, saturation, and hue
3440 % as one parameter (e.g. 90,150,100). If the image colorspace is HSL, the
3441 % modulation is lightness, saturation, and hue. For HWB, use blackness,
3442 % whiteness, and hue. And for HCL, use chrome, luma, and hue.
3444 % The format of the ModulateImage method is:
3446 % MagickBooleanType ModulateImage(Image *image,const char *modulate)
3448 % A description of each parameter follows:
3450 % o image: the image.
3452 % o percent_*: Define the percent change in brightness, saturation, and
3458 MagickBooleanType AccelerateModulateImage(Image* image, double percent_brightness, double percent_hue, double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
3460 MagickBooleanType status;
3462 assert(image != NULL);
3463 assert(exception != NULL);
3465 status = checkOpenCLEnvironment(exception);
3466 if (status == MagickFalse)
3469 status = checkAccelerateCondition(image, AllChannels);
3470 if (status == MagickFalse)
3473 if ((colorspace != HSLColorspace && colorspace != UndefinedColorspace))
3477 status = ComputeModulateImage(image,percent_brightness, percent_hue, percent_saturation, colorspace, exception);
3481 MagickBooleanType ComputeNegateImageChannel(Image* image, const ChannelType channel, const MagickBooleanType magick_unused(grayscale), ExceptionInfo* exception)
3486 MagickBooleanType outputReady;
3492 MagickSizeType length;
3495 cl_command_queue queue;
3496 cl_kernel negateKernel;
3498 cl_mem inputImageBuffer;
3499 cl_mem_flags mem_flags;
3503 Image * inputImage = image;
3505 magick_unreferenced(grayscale);
3508 inputImageBuffer = NULL;
3509 negateKernel = NULL;
3511 assert(inputImage != (Image *) NULL);
3512 assert(inputImage->signature == MagickSignature);
3513 if (inputImage->debug != MagickFalse)
3514 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",inputImage->filename);
3517 * initialize opencl env
3519 clEnv = GetDefaultOpenCLEnv();
3520 context = GetOpenCLContext(clEnv);
3521 queue = AcquireOpenCLCommandQueue(clEnv);
3523 outputReady = MagickFalse;
3525 /* Create and initialize OpenCL buffers.
3526 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
3527 assume this will get a writable image
3529 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3530 if (inputPixels == (void *) NULL)
3532 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3536 /* If the host pointer is aligned to the size of CLPixelPacket,
3537 then use the host buffer directly from the GPU; otherwise,
3538 create a buffer on the GPU and copy the data over
3540 if (ALIGNED(inputPixels,CLPixelPacket))
3542 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3546 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3548 /* create a CL buffer from image pixel buffer */
3549 length = inputImage->columns * inputImage->rows;
3550 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3551 if (clStatus != CL_SUCCESS)
3553 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3557 negateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Negate");
3558 if (negateKernel == NULL)
3560 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3565 clStatus=clEnv->library->clSetKernelArg(negateKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3566 clStatus=clEnv->library->clSetKernelArg(negateKernel,i++,sizeof(ChannelType),(void *)&channel);
3567 if (clStatus != CL_SUCCESS)
3569 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3570 printf("no kernel\n");
3575 size_t global_work_size[2];
3576 global_work_size[0] = inputImage->columns;
3577 global_work_size[1] = inputImage->rows;
3578 /* launch the kernel */
3579 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, negateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3580 if (clStatus != CL_SUCCESS)
3582 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3585 clEnv->library->clFlush(queue);
3588 if (ALIGNED(inputPixels,CLPixelPacket))
3590 length = inputImage->columns * inputImage->rows;
3591 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3595 length = inputImage->columns * inputImage->rows;
3596 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3598 if (clStatus != CL_SUCCESS)
3600 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3604 outputReady = MagickTrue;
3607 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3610 //ReleasePixelCachePixels();
3614 if (inputImageBuffer!=NULL)
3615 clEnv->library->clReleaseMemObject(inputImageBuffer);
3616 if (negateKernel!=NULL)
3617 RelinquishOpenCLKernel(clEnv, negateKernel);
3619 RelinquishOpenCLCommandQueue(clEnv, queue);
3627 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3631 % N e g a t e I m a g e w i t h O p e n C L %
3635 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3638 % A description of each parameter follows:
3640 % o image: the image.
3642 % o channel: the channel.
3644 % o grayscale: If MagickTrue, only negate grayscale pixels within the image.
3649 MagickBooleanType AccelerateNegateImageChannel(Image* image, const ChannelType channel, const MagickBooleanType grayscale, ExceptionInfo* exception)
3651 MagickBooleanType status;
3653 assert(image != NULL);
3654 assert(exception != NULL);
3656 status = checkOpenCLEnvironment(exception);
3657 if (status == MagickFalse)
3660 status = checkAccelerateCondition(image, AllChannels);
3661 if (status == MagickFalse)
3664 status = ComputeNegateImageChannel(image,channel,grayscale,exception);
3670 MagickBooleanType ComputeGrayscaleImage(Image* image, const PixelIntensityMethod method, ExceptionInfo* exception)
3675 cl_int intensityMethod;
3678 MagickBooleanType outputReady;
3684 MagickSizeType length;
3687 cl_command_queue queue;
3688 cl_kernel grayscaleKernel;
3690 cl_mem inputImageBuffer;
3691 cl_mem_flags mem_flags;
3695 Image * inputImage = image;
3698 inputImageBuffer = NULL;
3699 grayscaleKernel = NULL;
3701 assert(inputImage != (Image *) NULL);
3702 assert(inputImage->signature == MagickSignature);
3703 if (inputImage->debug != MagickFalse)
3704 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",inputImage->filename);
3707 * initialize opencl env
3709 clEnv = GetDefaultOpenCLEnv();
3710 context = GetOpenCLContext(clEnv);
3711 queue = AcquireOpenCLCommandQueue(clEnv);
3713 outputReady = MagickFalse;
3715 /* Create and initialize OpenCL buffers.
3716 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
3717 assume this will get a writable image
3719 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
3720 if (inputPixels == (void *) NULL)
3722 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
3726 /* If the host pointer is aligned to the size of CLPixelPacket,
3727 then use the host buffer directly from the GPU; otherwise,
3728 create a buffer on the GPU and copy the data over
3730 if (ALIGNED(inputPixels,CLPixelPacket))
3732 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3736 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3738 /* create a CL buffer from image pixel buffer */
3739 length = inputImage->columns * inputImage->rows;
3740 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
3741 if (clStatus != CL_SUCCESS)
3743 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3747 intensityMethod = method;
3748 colorspace = image->colorspace;
3750 grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Grayscale");
3751 if (grayscaleKernel == NULL)
3753 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3758 clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3759 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&intensityMethod);
3760 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&colorspace);
3761 if (clStatus != CL_SUCCESS)
3763 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3764 printf("no kernel\n");
3769 size_t global_work_size[2];
3770 global_work_size[0] = inputImage->columns;
3771 global_work_size[1] = inputImage->rows;
3772 /* launch the kernel */
3773 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, grayscaleKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3774 if (clStatus != CL_SUCCESS)
3776 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3779 clEnv->library->clFlush(queue);
3782 if (ALIGNED(inputPixels,CLPixelPacket))
3784 length = inputImage->columns * inputImage->rows;
3785 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3789 length = inputImage->columns * inputImage->rows;
3790 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3792 if (clStatus != CL_SUCCESS)
3794 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
3798 outputReady = MagickTrue;
3801 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3804 //ReleasePixelCachePixels();
3808 if (inputImageBuffer!=NULL)
3809 clEnv->library->clReleaseMemObject(inputImageBuffer);
3810 if (grayscaleKernel!=NULL)
3811 RelinquishOpenCLKernel(clEnv, grayscaleKernel);
3813 RelinquishOpenCLCommandQueue(clEnv, queue);
3819 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3823 % G r a y s c a l e I m a g e w i t h O p e n C L %
3827 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3829 % GrayscaleImage() converts the colors in the reference image to gray.
3831 % The format of the GrayscaleImageChannel method is:
3833 % MagickBooleanType GrayscaleImage(Image *image,
3834 % const PixelIntensityMethod method)
3836 % A description of each parameter follows:
3838 % o image: the image.
3840 % o channel: the channel.
3845 MagickBooleanType AccelerateGrayscaleImage(Image* image, const PixelIntensityMethod method, ExceptionInfo* exception)
3847 MagickBooleanType status;
3849 assert(image != NULL);
3850 assert(exception != NULL);
3852 status = checkOpenCLEnvironment(exception);
3853 if (status == MagickFalse)
3856 status = checkAccelerateCondition(image, AllChannels);
3857 if (status == MagickFalse)
3860 if (method == Rec601LuminancePixelIntensityMethod || method == Rec709LuminancePixelIntensityMethod)
3863 if (image->colorspace != sRGBColorspace)
3866 status = ComputeGrayscaleImage(image,method,exception);
3871 static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
3872 cl_command_queue queue,
3873 cl_mem inputImageBuffer,
3874 cl_mem histogramBuffer,
3876 const ChannelType channel,
3877 ExceptionInfo * _exception)
3880 *exception=_exception;
3885 MagickBooleanType outputReady;
3889 size_t global_work_size[2];
3891 cl_kernel histogramKernel;
3896 histogramKernel = NULL;
3898 outputReady = MagickFalse;
3899 method = inputImage->intensity;
3900 colorspace = inputImage->colorspace;
3902 /* get the OpenCL kernel */
3903 histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram");
3904 if (histogramKernel == NULL)
3906 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3910 /* set the kernel arguments */
3912 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
3913 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel);
3914 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&method);
3915 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&colorspace);
3916 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer);
3917 if (clStatus != CL_SUCCESS)
3919 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3923 /* launch the kernel */
3924 global_work_size[0] = inputImage->columns;
3925 global_work_size[1] = inputImage->rows;
3927 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
3929 if (clStatus != CL_SUCCESS)
3931 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3934 clEnv->library->clFlush(queue);
3936 outputReady = MagickTrue;
3939 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3941 if (histogramKernel!=NULL)
3942 RelinquishOpenCLKernel(clEnv, histogramKernel);
3948 MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const ChannelType channel, ExceptionInfo * _exception)
3950 #define EqualizeImageTag "Equalize/Image"
3953 *exception=_exception;
3970 Image * image = inputImage;
3972 MagickBooleanType outputReady;
3977 MagickBooleanType status;
3979 size_t global_work_size[2];
3982 cl_mem_flags mem_flags;
3985 cl_mem inputImageBuffer;
3986 cl_mem histogramBuffer;
3987 cl_mem equalizeMapBuffer;
3988 cl_kernel histogramKernel;
3989 cl_kernel equalizeKernel;
3990 cl_command_queue queue;
3994 MagickSizeType length;
3997 inputImageBuffer = NULL;
3998 histogramBuffer = NULL;
3999 equalizeMapBuffer = NULL;
4000 histogramKernel = NULL;
4001 equalizeKernel = NULL;
4004 outputReady = MagickFalse;
4006 assert(inputImage != (Image *) NULL);
4007 assert(inputImage->signature == MagickSignature);
4008 if (inputImage->debug != MagickFalse)
4009 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",inputImage->filename);
4012 * initialize opencl env
4014 clEnv = GetDefaultOpenCLEnv();
4015 context = GetOpenCLContext(clEnv);
4016 queue = AcquireOpenCLCommandQueue(clEnv);
4019 Allocate and initialize histogram arrays.
4021 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram));
4022 if (histogram == (cl_uint4 *) NULL)
4023 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
4025 /* reset histogram */
4026 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram));
4028 /* Create and initialize OpenCL buffers. */
4029 /* inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); */
4030 /* assume this will get a writable image */
4031 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
4033 if (inputPixels == (void *) NULL)
4035 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
4038 /* If the host pointer is aligned to the size of CLPixelPacket,
4039 then use the host buffer directly from the GPU; otherwise,
4040 create a buffer on the GPU and copy the data over */
4041 if (ALIGNED(inputPixels,CLPixelPacket))
4043 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4047 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
4049 /* create a CL buffer from image pixel buffer */
4050 length = inputImage->columns * inputImage->rows;
4051 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4052 if (clStatus != CL_SUCCESS)
4054 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4058 /* If the host pointer is aligned to the size of cl_uint,
4059 then use the host buffer directly from the GPU; otherwise,
4060 create a buffer on the GPU and copy the data over */
4061 if (ALIGNED(histogram,cl_uint4))
4063 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4064 hostPtr = histogram;
4068 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4069 hostPtr = histogram;
4071 /* create a CL buffer for histogram */
4072 length = (MaxMap+1);
4073 histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
4074 if (clStatus != CL_SUCCESS)
4076 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4080 status = LaunchHistogramKernel(clEnv, queue, inputImageBuffer, histogramBuffer, image, channel, exception);
4081 if (status == MagickFalse)
4084 /* read from the kenel output */
4085 if (ALIGNED(histogram,cl_uint4))
4087 length = (MaxMap+1);
4088 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
4092 length = (MaxMap+1);
4093 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
4095 if (clStatus != CL_SUCCESS)
4097 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4101 /* unmap, don't block gpu to use this buffer again. */
4102 if (ALIGNED(histogram,cl_uint4))
4104 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
4105 if (clStatus != CL_SUCCESS)
4107 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
4112 /* recreate input buffer later, in case image updated */
4113 #ifdef RECREATEBUFFER
4114 if (inputImageBuffer!=NULL)
4115 clEnv->library->clReleaseMemObject(inputImageBuffer);
4119 equalize_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*equalize_map));
4120 if (equalize_map == (PixelPacket *) NULL)
4121 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
4123 map=(FloatPixelPacket *) AcquireQuantumMemory(MaxMap+1UL,sizeof(*map));
4124 if (map == (FloatPixelPacket *) NULL)
4125 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
4128 Integrate the histogram to get the equalization map.
4130 (void) ResetMagickMemory(&intensity,0,sizeof(intensity));
4131 for (i=0; i <= (ssize_t) MaxMap; i++)
4133 if ((channel & SyncChannels) != 0)
4135 intensity.red+=histogram[i].s[2];
4139 if ((channel & RedChannel) != 0)
4140 intensity.red+=histogram[i].s[2];
4141 if ((channel & GreenChannel) != 0)
4142 intensity.green+=histogram[i].s[1];
4143 if ((channel & BlueChannel) != 0)
4144 intensity.blue+=histogram[i].s[0];
4145 if ((channel & OpacityChannel) != 0)
4146 intensity.opacity+=histogram[i].s[3];
4148 if (((channel & IndexChannel) != 0) &&
4149 (image->colorspace == CMYKColorspace))
4151 intensity.index+=histogram[i].index;
4157 white=map[(int) MaxMap];
4158 (void) ResetMagickMemory(equalize_map,0,(MaxMap+1)*sizeof(*equalize_map));
4159 for (i=0; i <= (ssize_t) MaxMap; i++)
4161 if ((channel & SyncChannels) != 0)
4163 if (white.red != black.red)
4164 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4165 (map[i].red-black.red))/(white.red-black.red)));
4168 if (((channel & RedChannel) != 0) && (white.red != black.red))
4169 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4170 (map[i].red-black.red))/(white.red-black.red)));
4171 if (((channel & GreenChannel) != 0) && (white.green != black.green))
4172 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4173 (map[i].green-black.green))/(white.green-black.green)));
4174 if (((channel & BlueChannel) != 0) && (white.blue != black.blue))
4175 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4176 (map[i].blue-black.blue))/(white.blue-black.blue)));
4177 if (((channel & OpacityChannel) != 0) && (white.opacity != black.opacity))
4178 equalize_map[i].opacity=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4179 (map[i].opacity-black.opacity))/(white.opacity-black.opacity)));
4181 if ((((channel & IndexChannel) != 0) &&
4182 (image->colorspace == CMYKColorspace)) &&
4183 (white.index != black.index))
4184 equalize_map[i].index=ScaleMapToQuantum((MagickRealType) ((MaxMap*
4185 (map[i].index-black.index))/(white.index-black.index)));
4189 if (image->storage_class == PseudoClass)
4194 for (i=0; i < (ssize_t) image->colors; i++)
4196 if ((channel & SyncChannels) != 0)
4198 if (white.red != black.red)
4200 image->colormap[i].red=equalize_map[
4201 ScaleQuantumToMap(image->colormap[i].red)].red;
4202 image->colormap[i].green=equalize_map[
4203 ScaleQuantumToMap(image->colormap[i].green)].red;
4204 image->colormap[i].blue=equalize_map[
4205 ScaleQuantumToMap(image->colormap[i].blue)].red;
4206 image->colormap[i].opacity=equalize_map[
4207 ScaleQuantumToMap(image->colormap[i].opacity)].red;
4211 if (((channel & RedChannel) != 0) && (white.red != black.red))
4212 image->colormap[i].red=equalize_map[
4213 ScaleQuantumToMap(image->colormap[i].red)].red;
4214 if (((channel & GreenChannel) != 0) && (white.green != black.green))
4215 image->colormap[i].green=equalize_map[
4216 ScaleQuantumToMap(image->colormap[i].green)].green;
4217 if (((channel & BlueChannel) != 0) && (white.blue != black.blue))
4218 image->colormap[i].blue=equalize_map[
4219 ScaleQuantumToMap(image->colormap[i].blue)].blue;
4220 if (((channel & OpacityChannel) != 0) &&
4221 (white.opacity != black.opacity))
4222 image->colormap[i].opacity=equalize_map[
4223 ScaleQuantumToMap(image->colormap[i].opacity)].opacity;
4231 /* GPU can work on this again, image and equalize map as input
4232 image: uchar4 (CLPixelPacket)
4233 equalize_map: uchar4 (PixelPacket)
4234 black, white: float4 (FloatPixelPacket) */
4236 #ifdef RECREATEBUFFER
4237 /* If the host pointer is aligned to the size of CLPixelPacket,
4238 then use the host buffer directly from the GPU; otherwise,
4239 create a buffer on the GPU and copy the data over */
4240 if (ALIGNED(inputPixels,CLPixelPacket))
4242 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4246 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4248 /* create a CL buffer from image pixel buffer */
4249 length = inputImage->columns * inputImage->rows;
4250 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4251 if (clStatus != CL_SUCCESS)
4253 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4258 /* Create and initialize OpenCL buffers. */
4259 if (ALIGNED(equalize_map, PixelPacket))
4261 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4262 hostPtr = equalize_map;
4266 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4267 hostPtr = equalize_map;
4269 /* create a CL buffer for eqaulize_map */
4270 length = (MaxMap+1);
4271 equalizeMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
4272 if (clStatus != CL_SUCCESS)
4274 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4278 /* get the OpenCL kernel */
4279 equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize");
4280 if (equalizeKernel == NULL)
4282 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4286 /* set the kernel arguments */
4288 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
4289 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(ChannelType),&channel);
4290 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&equalizeMapBuffer);
4291 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&white);
4292 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black);
4293 if (clStatus != CL_SUCCESS)
4295 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4299 /* launch the kernel */
4300 global_work_size[0] = inputImage->columns;
4301 global_work_size[1] = inputImage->rows;
4303 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4305 if (clStatus != CL_SUCCESS)
4307 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4310 clEnv->library->clFlush(queue);
4312 /* read the data back */
4313 if (ALIGNED(inputPixels,CLPixelPacket))
4315 length = inputImage->columns * inputImage->rows;
4316 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4320 length = inputImage->columns * inputImage->rows;
4321 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
4323 if (clStatus != CL_SUCCESS)
4325 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4329 outputReady = MagickTrue;
4332 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4335 /*ReleasePixelCachePixels();*/
4339 if (inputImageBuffer!=NULL)
4340 clEnv->library->clReleaseMemObject(inputImageBuffer);
4343 map=(FloatPixelPacket *) RelinquishMagickMemory(map);
4345 if (equalizeMapBuffer!=NULL)
4346 clEnv->library->clReleaseMemObject(equalizeMapBuffer);
4347 if (equalize_map!=NULL)
4348 equalize_map=(PixelPacket *) RelinquishMagickMemory(equalize_map);
4350 if (histogramBuffer!=NULL)
4351 clEnv->library->clReleaseMemObject(histogramBuffer);
4352 if (histogram!=NULL)
4353 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
4355 if (histogramKernel!=NULL)
4356 RelinquishOpenCLKernel(clEnv, histogramKernel);
4357 if (equalizeKernel!=NULL)
4358 RelinquishOpenCLKernel(clEnv, equalizeKernel);
4361 RelinquishOpenCLCommandQueue(clEnv, queue);
4367 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4371 % E q u a l i z e I m a g e w i t h O p e n C L %
4375 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4377 % EqualizeImage() applies a histogram equalization to the image.
4379 % The format of the EqualizeImage method is:
4381 % MagickBooleanType EqualizeImage(Image *image)
4382 % MagickBooleanType EqualizeImageChannel(Image *image,
4383 % const ChannelType channel)
4385 % A description of each parameter follows:
4387 % o image: the image.
4389 % o channel: the channel.
4395 MagickBooleanType AccelerateEqualizeImage(Image* image, const ChannelType channel, ExceptionInfo* exception)
4397 MagickBooleanType status;
4399 assert(image != NULL);
4400 assert(exception != NULL);
4402 status = checkOpenCLEnvironment(exception);
4403 if (status == MagickFalse)
4406 status = checkAccelerateCondition(image, channel);
4407 if (status == MagickFalse)
4410 status = checkHistogramCondition(image, channel);
4411 if (status == MagickFalse)
4414 status = ComputeEqualizeImage(image,channel,exception);
4420 MagickExport MagickBooleanType ComputeContrastStretchImageChannel(Image *image,
4421 const ChannelType channel,const double black_point,const double white_point,
4422 ExceptionInfo * _exception)
4424 #define MaxRange(color) ((MagickRealType) ScaleQuantumToMap((Quantum) (color)))
4425 #define ContrastStretchImageTag "ContrastStretch/Image"
4428 *exception=_exception;
4448 MagickBooleanType outputReady;
4453 MagickBooleanType status;
4455 size_t global_work_size[2];
4458 cl_mem_flags mem_flags;
4461 cl_mem inputImageBuffer;
4462 cl_mem histogramBuffer;
4463 cl_mem stretchMapBuffer;
4464 cl_kernel histogramKernel;
4465 cl_kernel stretchKernel;
4466 cl_command_queue queue;
4470 MagickSizeType length;
4474 inputImageBuffer = NULL;
4475 histogramBuffer = NULL;
4476 stretchMapBuffer = NULL;
4477 histogramKernel = NULL;
4478 stretchKernel = NULL;
4481 outputReady = MagickFalse;
4484 assert(image != (Image *) NULL);
4485 assert(image->signature == MagickSignature);
4486 if (image->debug != MagickFalse)
4487 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
4489 //exception=(&image->exception);
4492 * initialize opencl env
4494 clEnv = GetDefaultOpenCLEnv();
4495 context = GetOpenCLContext(clEnv);
4496 queue = AcquireOpenCLCommandQueue(clEnv);
4499 Allocate and initialize histogram arrays.
4501 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL, sizeof(*histogram));
4503 if (histogram == (cl_uint4 *) NULL)
4504 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed", image->filename);
4506 /* reset histogram */
4507 (void) ResetMagickMemory(histogram,0,(MaxMap+1)*sizeof(*histogram));
4510 if (IsGrayImage(image,exception) != MagickFalse)
4511 (void) SetImageColorspace(image,GRAYColorspace);
4520 /* Create and initialize OpenCL buffers. */
4521 /* inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); */
4522 /* assume this will get a writable image */
4523 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
4525 if (inputPixels == (void *) NULL)
4527 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
4530 /* If the host pointer is aligned to the size of CLPixelPacket,
4531 then use the host buffer directly from the GPU; otherwise,
4532 create a buffer on the GPU and copy the data over */
4533 if (ALIGNED(inputPixels,CLPixelPacket))
4535 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4539 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
4541 /* create a CL buffer from image pixel buffer */
4542 length = inputImage->columns * inputImage->rows;
4543 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4544 if (clStatus != CL_SUCCESS)
4546 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4550 /* If the host pointer is aligned to the size of cl_uint,
4551 then use the host buffer directly from the GPU; otherwise,
4552 create a buffer on the GPU and copy the data over */
4553 if (ALIGNED(histogram,cl_uint4))
4555 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4556 hostPtr = histogram;
4560 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4561 hostPtr = histogram;
4563 /* create a CL buffer for histogram */
4564 length = (MaxMap+1);
4565 histogramBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus);
4566 if (clStatus != CL_SUCCESS)
4568 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4572 status = LaunchHistogramKernel(clEnv, queue, inputImageBuffer, histogramBuffer, image, channel, exception);
4573 if (status == MagickFalse)
4576 /* read from the kenel output */
4577 if (ALIGNED(histogram,cl_uint4))
4579 length = (MaxMap+1);
4580 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
4584 length = (MaxMap+1);
4585 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length * sizeof(cl_uint4), histogram, 0, NULL, NULL);
4587 if (clStatus != CL_SUCCESS)
4589 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4593 /* unmap, don't block gpu to use this buffer again. */
4594 if (ALIGNED(histogram,cl_uint4))
4596 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
4597 if (clStatus != CL_SUCCESS)
4599 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
4604 /* recreate input buffer later, in case image updated */
4605 #ifdef RECREATEBUFFER
4606 if (inputImageBuffer!=NULL)
4607 clEnv->library->clReleaseMemObject(inputImageBuffer);
4612 Find the histogram boundaries by locating the black/white levels.
4615 white.red=MaxRange(QuantumRange);
4616 if ((channel & RedChannel) != 0)
4619 for (i=0; i <= (ssize_t) MaxMap; i++)
4621 intensity+=histogram[i].s[2];
4622 if (intensity > black_point)
4625 black.red=(MagickRealType) i;
4627 for (i=(ssize_t) MaxMap; i != 0; i--)
4629 intensity+=histogram[i].s[2];
4630 if (intensity > ((double) image->columns*image->rows-white_point))
4633 white.red=(MagickRealType) i;
4636 white.green=MaxRange(QuantumRange);
4637 if ((channel & GreenChannel) != 0)
4640 for (i=0; i <= (ssize_t) MaxMap; i++)
4642 intensity+=histogram[i].s[2];
4643 if (intensity > black_point)
4646 black.green=(MagickRealType) i;
4648 for (i=(ssize_t) MaxMap; i != 0; i--)
4650 intensity+=histogram[i].s[2];
4651 if (intensity > ((double) image->columns*image->rows-white_point))
4654 white.green=(MagickRealType) i;
4657 white.blue=MaxRange(QuantumRange);
4658 if ((channel & BlueChannel) != 0)
4661 for (i=0; i <= (ssize_t) MaxMap; i++)
4663 intensity+=histogram[i].s[2];
4664 if (intensity > black_point)
4667 black.blue=(MagickRealType) i;
4669 for (i=(ssize_t) MaxMap; i != 0; i--)
4671 intensity+=histogram[i].s[2];
4672 if (intensity > ((double) image->columns*image->rows-white_point))
4675 white.blue=(MagickRealType) i;
4678 white.opacity=MaxRange(QuantumRange);
4679 if ((channel & OpacityChannel) != 0)
4682 for (i=0; i <= (ssize_t) MaxMap; i++)
4684 intensity+=histogram[i].s[2];
4685 if (intensity > black_point)
4688 black.opacity=(MagickRealType) i;
4690 for (i=(ssize_t) MaxMap; i != 0; i--)
4692 intensity+=histogram[i].s[2];
4693 if (intensity > ((double) image->columns*image->rows-white_point))
4696 white.opacity=(MagickRealType) i;
4700 white.index=MaxRange(QuantumRange);
4701 if (((channel & IndexChannel) != 0) && (image->colorspace == CMYKColorspace))
4704 for (i=0; i <= (ssize_t) MaxMap; i++)
4706 intensity+=histogram[i].index;
4707 if (intensity > black_point)
4710 black.index=(MagickRealType) i;
4712 for (i=(ssize_t) MaxMap; i != 0; i--)
4714 intensity+=histogram[i].index;
4715 if (intensity > ((double) image->columns*image->rows-white_point))
4718 white.index=(MagickRealType) i;
4723 stretch_map=(PixelPacket *) AcquireQuantumMemory(MaxMap+1UL,
4724 sizeof(*stretch_map));
4726 if (stretch_map == (PixelPacket *) NULL)
4727 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed",
4731 Stretch the histogram to create the stretched image mapping.
4733 (void) ResetMagickMemory(stretch_map,0,(MaxMap+1)*sizeof(*stretch_map));
4734 for (i=0; i <= (ssize_t) MaxMap; i++)
4736 if ((channel & RedChannel) != 0)
4738 if (i < (ssize_t) black.red)
4739 stretch_map[i].red=(Quantum) 0;
4741 if (i > (ssize_t) white.red)
4742 stretch_map[i].red=QuantumRange;
4744 if (black.red != white.red)
4745 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap*
4746 (i-black.red)/(white.red-black.red)));
4748 if ((channel & GreenChannel) != 0)
4750 if (i < (ssize_t) black.green)
4751 stretch_map[i].green=0;
4753 if (i > (ssize_t) white.green)
4754 stretch_map[i].green=QuantumRange;
4756 if (black.green != white.green)
4757 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap*
4758 (i-black.green)/(white.green-black.green)));
4760 if ((channel & BlueChannel) != 0)
4762 if (i < (ssize_t) black.blue)
4763 stretch_map[i].blue=0;
4765 if (i > (ssize_t) white.blue)
4766 stretch_map[i].blue= QuantumRange;
4768 if (black.blue != white.blue)
4769 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap*
4770 (i-black.blue)/(white.blue-black.blue)));
4772 if ((channel & OpacityChannel) != 0)
4774 if (i < (ssize_t) black.opacity)
4775 stretch_map[i].opacity=0;
4777 if (i > (ssize_t) white.opacity)
4778 stretch_map[i].opacity=QuantumRange;
4780 if (black.opacity != white.opacity)
4781 stretch_map[i].opacity=ScaleMapToQuantum((MagickRealType) (MaxMap*
4782 (i-black.opacity)/(white.opacity-black.opacity)));
4785 if (((channel & IndexChannel) != 0) &&
4786 (image->colorspace == CMYKColorspace))
4788 if (i < (ssize_t) black.index)
4789 stretch_map[i].index=0;
4791 if (i > (ssize_t) white.index)
4792 stretch_map[i].index=QuantumRange;
4794 if (black.index != white.index)
4795 stretch_map[i].index=ScaleMapToQuantum((MagickRealType) (MaxMap*
4796 (i-black.index)/(white.index-black.index)));
4804 if (((channel & OpacityChannel) != 0) || (((channel & IndexChannel) != 0) &&
4805 (image->colorspace == CMYKColorspace)))
4806 image->storage_class=DirectClass;
4807 if (image->storage_class == PseudoClass)
4812 for (i=0; i < (ssize_t) image->colors; i++)
4814 if ((channel & RedChannel) != 0)
4816 if (black.red != white.red)
4817 image->colormap[i].red=stretch_map[
4818 ScaleQuantumToMap(image->colormap[i].red)].red;
4820 if ((channel & GreenChannel) != 0)
4822 if (black.green != white.green)
4823 image->colormap[i].green=stretch_map[
4824 ScaleQuantumToMap(image->colormap[i].green)].green;
4826 if ((channel & BlueChannel) != 0)
4828 if (black.blue != white.blue)
4829 image->colormap[i].blue=stretch_map[
4830 ScaleQuantumToMap(image->colormap[i].blue)].blue;
4832 if ((channel & OpacityChannel) != 0)
4834 if (black.opacity != white.opacity)
4835 image->colormap[i].opacity=stretch_map[
4836 ScaleQuantumToMap(image->colormap[i].opacity)].opacity;
4846 /* GPU can work on this again, image and equalize map as input
4847 image: uchar4 (CLPixelPacket)
4848 stretch_map: uchar4 (PixelPacket)
4849 black, white: float4 (FloatPixelPacket) */
4851 #ifdef RECREATEBUFFER
4852 /* If the host pointer is aligned to the size of CLPixelPacket,
4853 then use the host buffer directly from the GPU; otherwise,
4854 create a buffer on the GPU and copy the data over */
4855 if (ALIGNED(inputPixels,CLPixelPacket))
4857 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
4861 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4863 /* create a CL buffer from image pixel buffer */
4864 length = inputImage->columns * inputImage->rows;
4865 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
4866 if (clStatus != CL_SUCCESS)
4868 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4873 /* Create and initialize OpenCL buffers. */
4874 if (ALIGNED(stretch_map, PixelPacket))
4876 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
4877 hostPtr = stretch_map;
4881 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
4882 hostPtr = stretch_map;
4884 /* create a CL buffer for stretch_map */
4885 length = (MaxMap+1);
4886 stretchMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
4887 if (clStatus != CL_SUCCESS)
4889 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4893 /* get the OpenCL kernel */
4894 stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Stretch");
4895 if (stretchKernel == NULL)
4897 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4901 /* set the kernel arguments */
4903 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&inputImageBuffer);
4904 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(ChannelType),&channel);
4905 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&stretchMapBuffer);
4906 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&white);
4907 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black);
4908 if (clStatus != CL_SUCCESS)
4910 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4914 /* launch the kernel */
4915 global_work_size[0] = inputImage->columns;
4916 global_work_size[1] = inputImage->rows;
4918 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
4920 if (clStatus != CL_SUCCESS)
4922 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4925 clEnv->library->clFlush(queue);
4927 /* read the data back */
4928 if (ALIGNED(inputPixels,CLPixelPacket))
4930 length = inputImage->columns * inputImage->rows;
4931 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
4935 length = inputImage->columns * inputImage->rows;
4936 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
4938 if (clStatus != CL_SUCCESS)
4940 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
4944 outputReady = MagickTrue;
4947 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4950 /*ReleasePixelCachePixels();*/
4954 if (inputImageBuffer!=NULL)
4955 clEnv->library->clReleaseMemObject(inputImageBuffer);
4957 if (stretchMapBuffer!=NULL)
4958 clEnv->library->clReleaseMemObject(stretchMapBuffer);
4959 if (stretch_map!=NULL)
4960 stretch_map=(PixelPacket *) RelinquishMagickMemory(stretch_map);
4963 if (histogramBuffer!=NULL)
4964 clEnv->library->clReleaseMemObject(histogramBuffer);
4965 if (histogram!=NULL)
4966 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
4969 if (histogramKernel!=NULL)
4970 RelinquishOpenCLKernel(clEnv, histogramKernel);
4971 if (stretchKernel!=NULL)
4972 RelinquishOpenCLKernel(clEnv, stretchKernel);
4975 RelinquishOpenCLCommandQueue(clEnv, queue);
4982 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4986 % C o n t r a s t S t r e t c h I m a g e w i t h O p e n C L %
4990 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4992 % ContrastStretchImage() is a simple image enhancement technique that attempts
4993 % to improve the contrast in an image by `stretching' the range of intensity
4994 % values it contains to span a desired range of values. It differs from the
4995 % more sophisticated histogram equalization in that it can only apply a
4996 % linear scaling function to the image pixel values. As a result the
4997 % `enhancement' is less harsh.
4999 % The format of the ContrastStretchImage method is:
5001 % MagickBooleanType ContrastStretchImage(Image *image,
5002 % const char *levels)
5003 % MagickBooleanType ContrastStretchImageChannel(Image *image,
5004 % const size_t channel,const double black_point,
5005 % const double white_point)
5007 % A description of each parameter follows:
5009 % o image: the image.
5011 % o channel: the channel.
5013 % o black_point: the black point.
5015 % o white_point: the white point.
5017 % o levels: Specify the levels where the black and white points have the
5018 % range of 0 to number-of-pixels (e.g. 1%, 10x90%, etc.).
5022 MagickExport MagickBooleanType AccelerateContrastStretchImageChannel(
5023 Image * image, const ChannelType channel, const double black_point, const double white_point,
5024 ExceptionInfo* exception)
5026 MagickBooleanType status;
5028 assert(image != NULL);
5029 assert(exception != NULL);
5031 status = checkOpenCLEnvironment(exception);
5032 if (status == MagickFalse)
5035 status = checkAccelerateCondition(image, channel);
5036 if (status == MagickFalse)
5039 status = checkHistogramCondition(image, channel);
5040 if (status == MagickFalse)
5043 status = ComputeContrastStretchImageChannel(image,channel, black_point, white_point, exception);
5049 static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exception)
5052 MagickBooleanType outputReady = MagickFalse;
5053 MagickCLEnv clEnv = NULL;
5056 size_t global_work_size[2];
5058 const void *inputPixels = NULL;
5059 Image* filteredImage = NULL;
5060 void *filteredPixels = NULL;
5062 MagickSizeType length;
5064 cl_mem_flags mem_flags;
5065 cl_context context = NULL;
5066 cl_mem inputImageBuffer = NULL;
5067 cl_mem tempImageBuffer[2];
5068 cl_mem filteredImageBuffer = NULL;
5069 cl_command_queue queue = NULL;
5070 cl_kernel hullPass1 = NULL;
5071 cl_kernel hullPass2 = NULL;
5073 unsigned int imageWidth, imageHeight;
5078 X[4] = {0, 1, 1,-1},
5079 Y[4] = {1, 0, 1, 1};
5081 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
5082 clEnv = GetDefaultOpenCLEnv();
5083 context = GetOpenCLContext(clEnv);
5084 queue = AcquireOpenCLCommandQueue(clEnv);
5086 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
5087 if (inputPixels == (void *) NULL)
5089 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
5093 if (ALIGNED(inputPixels,CLPixelPacket))
5095 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
5099 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
5101 /* create a CL buffer from image pixel buffer */
5102 length = inputImage->columns * inputImage->rows;
5103 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
5104 if (clStatus != CL_SUCCESS)
5106 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5110 mem_flags = CL_MEM_READ_WRITE;
5111 length = inputImage->columns * inputImage->rows;
5112 for (k = 0; k < 2; k++)
5114 tempImageBuffer[k] = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus);
5115 if (clStatus != CL_SUCCESS)
5117 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5122 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
5123 assert(filteredImage != NULL);
5124 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
5126 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
5129 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
5130 if (filteredPixels == (void *) NULL)
5132 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
5136 if (ALIGNED(filteredPixels,CLPixelPacket))
5138 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
5139 hostPtr = filteredPixels;
5143 mem_flags = CL_MEM_WRITE_ONLY;
5146 /* create a CL buffer from image pixel buffer */
5147 length = inputImage->columns * inputImage->rows;
5148 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
5149 if (clStatus != CL_SUCCESS)
5151 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5155 hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1");
5156 hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2");
5158 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&inputImageBuffer);
5159 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1));
5160 imageWidth = inputImage->columns;
5161 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth);
5162 imageHeight = inputImage->rows;
5163 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight);
5164 matte = (inputImage->matte==MagickFalse)?0:1;
5165 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
5166 if (clStatus != CL_SUCCESS)
5168 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5172 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1));
5173 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer);
5174 imageWidth = inputImage->columns;
5175 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth);
5176 imageHeight = inputImage->rows;
5177 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight);
5178 matte = (inputImage->matte==MagickFalse)?0:1;
5179 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
5180 if (clStatus != CL_SUCCESS)
5182 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5187 global_work_size[0] = inputImage->columns;
5188 global_work_size[1] = inputImage->rows;
5191 for (k = 0; k < 4; k++)
5200 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5201 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5202 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5203 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5204 if (clStatus != CL_SUCCESS)
5206 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5209 /* launch the kernel */
5210 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5211 if (clStatus != CL_SUCCESS)
5213 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5216 /* launch the kernel */
5217 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5218 if (clStatus != CL_SUCCESS)
5220 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5226 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
5227 offset.s[0] = -X[k];
5228 offset.s[1] = -Y[k];
5230 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5231 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5232 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5233 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5234 if (clStatus != CL_SUCCESS)
5236 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5239 /* launch the kernel */
5240 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5241 if (clStatus != CL_SUCCESS)
5243 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5246 /* launch the kernel */
5247 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5248 if (clStatus != CL_SUCCESS)
5250 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5254 offset.s[0] = -X[k];
5255 offset.s[1] = -Y[k];
5257 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5258 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5259 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5260 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5261 if (clStatus != CL_SUCCESS)
5263 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5266 /* launch the kernel */
5267 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5268 if (clStatus != CL_SUCCESS)
5270 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5273 /* launch the kernel */
5274 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5275 if (clStatus != CL_SUCCESS)
5277 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5284 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
5285 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
5286 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
5287 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
5290 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer);
5292 if (clStatus != CL_SUCCESS)
5294 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5297 /* launch the kernel */
5298 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5299 if (clStatus != CL_SUCCESS)
5301 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5304 /* launch the kernel */
5305 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
5306 if (clStatus != CL_SUCCESS)
5308 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5313 if (ALIGNED(filteredPixels,CLPixelPacket))
5315 length = inputImage->columns * inputImage->rows;
5316 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
5320 length = inputImage->columns * inputImage->rows;
5321 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
5323 if (clStatus != CL_SUCCESS)
5325 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
5329 outputReady = MagickTrue;
5332 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5334 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
5335 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
5336 for (k = 0; k < 2; k++)
5338 if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
5340 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
5341 if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
5342 if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
5343 if (outputReady == MagickFalse)
5345 if (filteredImage != NULL)
5347 DestroyImage(filteredImage);
5348 filteredImage = NULL;
5351 return filteredImage;
5355 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
5359 % D e s p e c k l e I m a g e w i t h O p e n C L %
5363 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
5365 % DespeckleImage() reduces the speckle noise in an image while perserving the
5366 % edges of the original image. A speckle removing filter uses a complementary
5367 % hulling technique (raising pixels that are darker than their surrounding
5368 % neighbors, then complementarily lowering pixels that are brighter than their
5369 % surrounding neighbors) to reduce the speckle index of that image (reference
5370 % Crimmins speckle removal).
5372 % The format of the DespeckleImage method is:
5374 % Image *DespeckleImage(const Image *image,ExceptionInfo *exception)
5376 % A description of each parameter follows:
5378 % o image: the image.
5380 % o exception: return any errors or warnings in this structure.
5385 Image* AccelerateDespeckleImage(const Image* image, ExceptionInfo* exception)
5387 MagickBooleanType status;
5388 Image* newImage = NULL;
5390 assert(image != NULL);
5391 assert(exception != NULL);
5393 status = checkOpenCLEnvironment(exception);
5394 if (status == MagickFalse)
5397 status = checkAccelerateCondition(image, AllChannels);
5398 if (status == MagickFalse)
5401 newImage = ComputeDespeckleImage(image,exception);
5405 static Image* ComputeAddNoiseImage(const Image* inputImage,
5406 const ChannelType channel, const NoiseType noise_type,
5407 ExceptionInfo *exception)
5409 MagickBooleanType outputReady = MagickFalse;
5410 MagickCLEnv clEnv = NULL;
5413 size_t global_work_size[2];
5415 const void *inputPixels = NULL;
5416 Image* filteredImage = NULL;
5417 void *filteredPixels = NULL;
5419 unsigned int inputColumns, inputRows;
5421 float *randomNumberBufferPtr = NULL;
5422 MagickSizeType length;
5423 unsigned int numRandomNumberPerPixel;
5424 unsigned int numRowsPerKernelLaunch;
5425 unsigned int numRandomNumberPerBuffer;
5430 RandomInfo **restrict random_info;
5432 #if defined(MAGICKCORE_OPENMP_SUPPORT)
5436 cl_mem_flags mem_flags;
5437 cl_context context = NULL;
5438 cl_mem inputImageBuffer = NULL;
5439 cl_mem randomNumberBuffer = NULL;
5440 cl_mem filteredImageBuffer = NULL;
5441 cl_command_queue queue = NULL;
5442 cl_kernel addNoiseKernel = NULL;
5445 clEnv = GetDefaultOpenCLEnv();
5446 context = GetOpenCLContext(clEnv);
5447 queue = AcquireOpenCLCommandQueue(clEnv);
5449 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
5450 if (inputPixels == (void *) NULL)
5452 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
5456 if (ALIGNED(inputPixels,CLPixelPacket))
5458 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
5462 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
5464 /* create a CL buffer from image pixel buffer */
5465 length = inputImage->columns * inputImage->rows;
5466 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
5467 if (clStatus != CL_SUCCESS)
5469 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5474 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
5475 assert(filteredImage != NULL);
5476 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
5478 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
5481 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
5482 if (filteredPixels == (void *) NULL)
5484 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
5488 if (ALIGNED(filteredPixels,CLPixelPacket))
5490 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
5491 hostPtr = filteredPixels;
5495 mem_flags = CL_MEM_WRITE_ONLY;
5498 /* create a CL buffer from image pixel buffer */
5499 length = inputImage->columns * inputImage->rows;
5500 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
5501 if (clStatus != CL_SUCCESS)
5503 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5507 /* find out how many random numbers needed by pixel */
5508 numRandomNumberPerPixel = 0;
5510 unsigned int numRandPerChannel = 0;
5515 case LaplacianNoise:
5518 numRandPerChannel = 1;
5521 case MultiplicativeGaussianNoise:
5523 numRandPerChannel = 2;
5527 if ((channel & RedChannel) != 0)
5528 numRandomNumberPerPixel+=numRandPerChannel;
5529 if ((channel & GreenChannel) != 0)
5530 numRandomNumberPerPixel+=numRandPerChannel;
5531 if ((channel & BlueChannel) != 0)
5532 numRandomNumberPerPixel+=numRandPerChannel;
5533 if ((channel & OpacityChannel) != 0)
5534 numRandomNumberPerPixel+=numRandPerChannel;
5537 numRowsPerKernelLaunch = 512;
5538 /* create a buffer for random numbers */
5539 numRandomNumberPerBuffer = (inputImage->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
5540 randomNumberBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, numRandomNumberPerBuffer*sizeof(float)
5544 /* set up the random number generators */
5546 option=GetImageArtifact(inputImage,"attenuate");
5547 if (option != (char *) NULL)
5548 attenuate=StringToDouble(option,(char **) NULL);
5549 random_info=AcquireRandomInfoThreadSet();
5550 #if defined(MAGICKCORE_OPENMP_SUPPORT)
5551 key=GetRandomSecretKey(random_info[0]);
5554 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
5557 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&inputImageBuffer);
5558 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
5559 inputColumns = inputImage->columns;
5560 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
5561 inputRows = inputImage->rows;
5562 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
5563 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
5564 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
5566 option=GetImageArtifact(inputImage,"attenuate");
5567 if (option != (char *) NULL)
5568 attenuate=(float)StringToDouble(option,(char **) NULL);
5569 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
5570 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
5571 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
5573 global_work_size[0] = inputColumns;
5574 for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
5576 /* Generate random numbers in the buffer */
5577 randomNumberBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, randomNumberBuffer, CL_TRUE, CL_MAP_WRITE, 0
5578 , numRandomNumberPerBuffer*sizeof(float), 0, NULL, NULL, &clStatus);
5579 if (clStatus != CL_SUCCESS)
5581 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
5585 #if defined(MAGICKCORE_OPENMP_SUPPORT)
5586 #pragma omp parallel for schedule(static,4) \
5587 num_threads((key == ~0UL) == 0 ? 1 : (size_t) GetMagickResourceLimit(ThreadResource))
5589 for (i = 0; i < numRandomNumberPerBuffer; i++)
5591 const int id = GetOpenMPThreadId();
5592 randomNumberBufferPtr[i] = (float)GetPseudoRandomValue(random_info[id]);
5595 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, randomNumberBuffer, randomNumberBufferPtr, 0, NULL, NULL);
5596 if (clStatus != CL_SUCCESS)
5598 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.",".");
5602 /* set the row offset */
5603 clEnv->library->clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
5604 global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
5605 clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
5608 if (ALIGNED(filteredPixels,CLPixelPacket))
5610 length = inputImage->columns * inputImage->rows;
5611 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
5615 length = inputImage->columns * inputImage->rows;
5616 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
5618 if (clStatus != CL_SUCCESS)
5620 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
5624 outputReady = MagickTrue;
5627 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5629 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
5630 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
5631 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
5632 if (randomNumberBuffer!=NULL) clEnv->library->clReleaseMemObject(randomNumberBuffer);
5633 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
5634 if (outputReady == MagickFalse
5635 && filteredImage != NULL)
5637 DestroyImage(filteredImage);
5638 filteredImage = NULL;
5640 return filteredImage;
5644 static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage,
5645 const ChannelType channel, const NoiseType noise_type,
5646 ExceptionInfo *exception)
5648 MagickBooleanType outputReady = MagickFalse;
5649 MagickCLEnv clEnv = NULL;
5652 size_t global_work_size[2];
5653 size_t random_work_size;
5655 const void *inputPixels = NULL;
5656 Image* filteredImage = NULL;
5657 void *filteredPixels = NULL;
5659 unsigned int inputColumns, inputRows;
5661 MagickSizeType length;
5662 unsigned int numRandomNumberPerPixel;
5663 unsigned int numRowsPerKernelLaunch;
5664 unsigned int numRandomNumberPerBuffer;
5665 unsigned int numRandomNumberGenerators;
5666 unsigned int initRandom;
5673 cl_mem_flags mem_flags;
5674 cl_context context = NULL;
5675 cl_mem inputImageBuffer = NULL;
5676 cl_mem randomNumberBuffer = NULL;
5677 cl_mem filteredImageBuffer = NULL;
5678 cl_mem randomNumberSeedsBuffer = NULL;
5679 cl_command_queue queue = NULL;
5680 cl_kernel addNoiseKernel = NULL;
5681 cl_kernel randomNumberGeneratorKernel = NULL;
5684 clEnv = GetDefaultOpenCLEnv();
5685 context = GetOpenCLContext(clEnv);
5686 queue = AcquireOpenCLCommandQueue(clEnv);
5688 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
5689 if (inputPixels == (void *) NULL)
5691 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
5695 if (ALIGNED(inputPixels,CLPixelPacket))
5697 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
5701 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
5703 /* create a CL buffer from image pixel buffer */
5704 length = inputImage->columns * inputImage->rows;
5705 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
5706 if (clStatus != CL_SUCCESS)
5708 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5713 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,MagickTrue,exception);
5714 assert(filteredImage != NULL);
5715 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
5717 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
5720 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
5721 if (filteredPixels == (void *) NULL)
5723 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
5727 if (ALIGNED(filteredPixels,CLPixelPacket))
5729 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
5730 hostPtr = filteredPixels;
5734 mem_flags = CL_MEM_WRITE_ONLY;
5737 /* create a CL buffer from image pixel buffer */
5738 length = inputImage->columns * inputImage->rows;
5739 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
5740 if (clStatus != CL_SUCCESS)
5742 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5746 /* find out how many random numbers needed by pixel */
5747 numRandomNumberPerPixel = 0;
5749 unsigned int numRandPerChannel = 0;
5754 case LaplacianNoise:
5757 numRandPerChannel = 1;
5760 case MultiplicativeGaussianNoise:
5762 numRandPerChannel = 2;
5766 if ((channel & RedChannel) != 0)
5767 numRandomNumberPerPixel+=numRandPerChannel;
5768 if ((channel & GreenChannel) != 0)
5769 numRandomNumberPerPixel+=numRandPerChannel;
5770 if ((channel & BlueChannel) != 0)
5771 numRandomNumberPerPixel+=numRandPerChannel;
5772 if ((channel & OpacityChannel) != 0)
5773 numRandomNumberPerPixel+=numRandPerChannel;
5776 numRowsPerKernelLaunch = 512;
5778 /* create a buffer for random numbers */
5779 numRandomNumberPerBuffer = (inputImage->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
5780 randomNumberBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, numRandomNumberPerBuffer*sizeof(float)
5784 /* setup the random number generators */
5785 unsigned long* seeds;
5786 numRandomNumberGenerators = 512;
5787 randomNumberSeedsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR|CL_MEM_READ_WRITE
5788 , numRandomNumberGenerators * 4 * sizeof(unsigned long), NULL, &clStatus);
5789 if (clStatus != CL_SUCCESS)
5791 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
5794 seeds = (unsigned long*) clEnv->library->clEnqueueMapBuffer(queue, randomNumberSeedsBuffer, CL_TRUE, CL_MAP_WRITE, 0
5795 , numRandomNumberGenerators*4*sizeof(unsigned long), 0, NULL, NULL, &clStatus);
5796 if (clStatus != CL_SUCCESS)
5798 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
5802 for (i = 0; i < numRandomNumberGenerators; i++) {
5803 RandomInfo* randomInfo = AcquireRandomInfo();
5804 const unsigned long* s = GetRandomInfoSeed(randomInfo);
5807 fNormalize = GetRandomInfoNormalize(randomInfo);
5810 randomInfo = DestroyRandomInfo(randomInfo);
5813 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, randomNumberSeedsBuffer, seeds, 0, NULL, NULL);
5814 if (clStatus != CL_SUCCESS)
5816 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.",".");
5820 randomNumberGeneratorKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE
5821 ,"randomNumberGeneratorKernel");
5824 clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberSeedsBuffer);
5825 clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(float),(void *)&fNormalize);
5826 clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
5828 clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&initRandom);
5829 clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerBuffer);
5831 random_work_size = numRandomNumberGenerators;
5834 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
5836 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&inputImageBuffer);
5837 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
5838 inputColumns = inputImage->columns;
5839 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
5840 inputRows = inputImage->rows;
5841 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
5842 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
5843 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
5845 option=GetImageArtifact(inputImage,"attenuate");
5846 if (option != (char *) NULL)
5847 attenuate=(float)StringToDouble(option,(char **) NULL);
5848 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
5849 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
5850 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
5852 global_work_size[0] = inputColumns;
5853 for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
5855 size_t generator_local_size = 64;
5856 /* Generate random numbers in the buffer */
5857 clEnv->library->clEnqueueNDRangeKernel(queue,randomNumberGeneratorKernel,1,NULL
5858 ,&random_work_size,&generator_local_size,0,NULL,NULL);
5859 if (initRandom != 0)
5861 /* make sure we only do init once */
5863 clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,3,sizeof(unsigned int),(void *)&initRandom);
5866 /* set the row offset */
5867 clEnv->library->clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
5868 global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
5869 clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
5872 if (ALIGNED(filteredPixels,CLPixelPacket))
5874 length = inputImage->columns * inputImage->rows;
5875 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
5879 length = inputImage->columns * inputImage->rows;
5880 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
5882 if (clStatus != CL_SUCCESS)
5884 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
5888 outputReady = MagickTrue;
5891 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5893 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
5894 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
5895 if (randomNumberGeneratorKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomNumberGeneratorKernel);
5896 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
5897 if (randomNumberBuffer!=NULL) clEnv->library->clReleaseMemObject(randomNumberBuffer);
5898 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
5899 if (randomNumberSeedsBuffer!=NULL) clEnv->library->clReleaseMemObject(randomNumberSeedsBuffer);
5900 if (outputReady == MagickFalse
5901 && filteredImage != NULL)
5903 DestroyImage(filteredImage);
5904 filteredImage = NULL;
5906 return filteredImage;
5912 Image* AccelerateAddNoiseImage(const Image *image, const ChannelType channel,
5913 const NoiseType noise_type,ExceptionInfo *exception)
5915 MagickBooleanType status;
5916 Image* filteredImage = NULL;
5918 assert(image != NULL);
5919 assert(exception != NULL);
5921 status = checkOpenCLEnvironment(exception);
5922 if (status == MagickFalse)
5925 status = checkAccelerateCondition(image, channel);
5926 if (status == MagickFalse)
5929 DisableMSCWarning(4127)
5930 if (sizeof(unsigned long) == 4)
5932 filteredImage = ComputeAddNoiseImageOptRandomNum(image,channel,noise_type,exception);
5934 filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
5936 return filteredImage;
5939 static MagickBooleanType LaunchRandomImageKernel(MagickCLEnv clEnv,
5940 cl_command_queue queue,
5941 cl_mem inputImageBuffer,
5942 const unsigned int imageColumns,
5943 const unsigned int imageRows,
5945 const unsigned int numGenerators,
5946 ExceptionInfo *exception)
5948 MagickBooleanType status = MagickFalse;
5949 size_t global_work_size;
5950 size_t local_work_size;
5954 cl_kernel randomImageKernel = NULL;
5956 randomImageKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RandomImage");
5959 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&inputImageBuffer);
5960 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageColumns);
5961 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_uint),(void*)&imageRows);
5962 clEnv->library->clSetKernelArg(randomImageKernel,k++,sizeof(cl_mem),(void*)&seedBuffer);
5964 const float randNormNumerator = 1.0f;
5965 const unsigned int randNormDenominator = (unsigned int)(~0UL);
5966 clEnv->library->clSetKernelArg(randomImageKernel,k++,
5967 sizeof(float),(void*)&randNormNumerator);
5968 clEnv->library->clSetKernelArg(randomImageKernel,k++,
5969 sizeof(cl_uint),(void*)&randNormDenominator);
5973 global_work_size = numGenerators;
5974 local_work_size = 64;
5976 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue,randomImageKernel,1,NULL,&global_work_size,
5977 &local_work_size,0,NULL,NULL);
5979 if (clStatus != CL_SUCCESS)
5981 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
5982 "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5985 status = MagickTrue;
5988 if (randomImageKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomImageKernel);
5992 static MagickBooleanType ComputeRandomImage(Image* inputImage,
5993 ExceptionInfo* exception)
5995 MagickBooleanType status = MagickFalse;
5997 MagickBooleanType outputReady = MagickFalse;
5998 MagickCLEnv clEnv = NULL;
6002 void *inputPixels = NULL;
6003 MagickSizeType length;
6005 cl_mem_flags mem_flags;
6006 cl_context context = NULL;
6007 cl_mem inputImageBuffer = NULL;
6008 cl_command_queue queue = NULL;
6010 /* Don't release this buffer in this function !!! */
6011 cl_mem randomNumberSeedsBuffer;
6013 clEnv = GetDefaultOpenCLEnv();
6014 context = GetOpenCLContext(clEnv);
6016 /* Create and initialize OpenCL buffers. */
6017 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
6018 if (inputPixels == (void *) NULL)
6020 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename);
6024 /* If the host pointer is aligned to the size of CLPixelPacket,
6025 then use the host buffer directly from the GPU; otherwise,
6026 create a buffer on the GPU and copy the data over */
6027 if (ALIGNED(inputPixels,CLPixelPacket))
6029 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
6033 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
6035 /* create a CL buffer from image pixel buffer */
6036 length = inputImage->columns * inputImage->rows;
6037 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6038 if (clStatus != CL_SUCCESS)
6040 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6044 queue = AcquireOpenCLCommandQueue(clEnv);
6046 randomNumberSeedsBuffer = GetAndLockRandSeedBuffer(clEnv);
6047 if (randomNumberSeedsBuffer==NULL)
6049 (void) OpenCLThrowMagickException(exception, GetMagickModule(),
6050 ResourceLimitWarning, "Failed to get GPU random number generators.",
6055 status = LaunchRandomImageKernel(clEnv,queue,
6057 inputImage->columns,
6059 randomNumberSeedsBuffer,
6060 GetNumRandGenerators(clEnv),
6062 if (status==MagickFalse)
6067 if (ALIGNED(inputPixels,CLPixelPacket))
6069 length = inputImage->columns * inputImage->rows;
6070 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
6074 length = inputImage->columns * inputImage->rows;
6075 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
6077 if (clStatus != CL_SUCCESS)
6079 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
6082 outputReady = MagickTrue;
6085 OpenCLLogException(__FUNCTION__,__LINE__,exception);
6087 UnlockRandSeedBuffer(clEnv);
6088 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
6089 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
6093 MagickExport MagickBooleanType AccelerateRandomImage(Image* image, ExceptionInfo* exception)
6095 MagickBooleanType status = MagickFalse;
6097 status = checkOpenCLEnvironment(exception);
6098 if (status==MagickFalse)
6101 status = checkAccelerateCondition(image, AllChannels);
6102 if (status==MagickFalse)
6105 status = ComputeRandomImage(image,exception);
6109 static Image* ComputeMotionBlurImage(const Image *inputImage,
6110 const ChannelType channel, const double *kernel, const size_t width,
6111 const OffsetInfo *offset, ExceptionInfo *exception)
6113 MagickBooleanType outputReady;
6114 Image* filteredImage;
6118 size_t global_work_size[2];
6119 size_t local_work_size[2];
6122 cl_mem_flags mem_flags;
6123 cl_mem inputImageBuffer, filteredImageBuffer, imageKernelBuffer,
6125 cl_kernel motionBlurKernel;
6126 cl_command_queue queue;
6128 const void *inputPixels;
6129 void *filteredPixels;
6131 float* kernelBufferPtr;
6132 int* offsetBufferPtr;
6133 MagickSizeType length;
6135 MagickPixelPacket bias;
6136 cl_float4 biasPixel;
6137 unsigned int imageWidth, imageHeight;
6141 outputReady = MagickFalse;
6143 filteredImage = NULL;
6144 inputImageBuffer = NULL;
6145 filteredImageBuffer = NULL;
6146 imageKernelBuffer = NULL;
6147 motionBlurKernel = NULL;
6151 clEnv = GetDefaultOpenCLEnv();
6152 context = GetOpenCLContext(clEnv);
6154 /* Create and initialize OpenCL buffers. */
6157 inputPixels = AcquirePixelCachePixels(inputImage, &length, exception);
6158 if (inputPixels == (const void *) NULL)
6160 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
6161 "UnableToReadPixelCache.","`%s'",inputImage->filename);
6165 // If the host pointer is aligned to the size of CLPixelPacket,
6166 // then use the host buffer directly from the GPU; otherwise,
6167 // create a buffer on the GPU and copy the data over
6168 if (ALIGNED(inputPixels,CLPixelPacket))
6170 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
6174 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
6176 // create a CL buffer from image pixel buffer
6177 length = inputImage->columns * inputImage->rows;
6178 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6179 length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6180 if (clStatus != CL_SUCCESS)
6182 (void) ThrowMagickException(exception, GetMagickModule(),
6183 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6188 filteredImage = CloneImage(inputImage,inputImage->columns,inputImage->rows,
6189 MagickTrue,exception);
6190 assert(filteredImage != NULL);
6191 if (SetImageStorageClass(filteredImage,DirectClass) != MagickTrue)
6193 (void) ThrowMagickException(exception, GetMagickModule(),
6194 ResourceLimitError, "CloneImage failed.", "'%s'", ".");
6197 filteredPixels = GetPixelCachePixels(filteredImage, &length, exception);
6198 if (filteredPixels == (void *) NULL)
6200 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
6201 "UnableToReadPixelCache.","`%s'",filteredImage->filename);
6205 if (ALIGNED(filteredPixels,CLPixelPacket))
6207 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
6208 hostPtr = filteredPixels;
6212 mem_flags = CL_MEM_WRITE_ONLY;
6215 // create a CL buffer from image pixel buffer
6216 length = inputImage->columns * inputImage->rows;
6217 filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6218 length * sizeof(CLPixelPacket), hostPtr, &clStatus);
6219 if (clStatus != CL_SUCCESS)
6221 (void) ThrowMagickException(exception, GetMagickModule(),
6222 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6227 imageKernelBuffer = clEnv->library->clCreateBuffer(context,
6228 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(float), NULL,
6230 if (clStatus != CL_SUCCESS)
6232 (void) ThrowMagickException(exception, GetMagickModule(),
6233 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6237 queue = AcquireOpenCLCommandQueue(clEnv);
6238 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
6239 CL_TRUE, CL_MAP_WRITE, 0, width * sizeof(float), 0, NULL, NULL, &clStatus);
6240 if (clStatus != CL_SUCCESS)
6242 (void) ThrowMagickException(exception, GetMagickModule(),
6243 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
6246 for (i = 0; i < width; i++)
6248 kernelBufferPtr[i] = (float) kernel[i];
6250 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr,
6252 if (clStatus != CL_SUCCESS)
6254 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6255 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
6259 offsetBuffer = clEnv->library->clCreateBuffer(context,
6260 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(cl_int2), NULL,
6262 if (clStatus != CL_SUCCESS)
6264 (void) ThrowMagickException(exception, GetMagickModule(),
6265 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
6269 offsetBufferPtr = (int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
6270 CL_MAP_WRITE, 0, width * sizeof(cl_int2), 0, NULL, NULL, &clStatus);
6271 if (clStatus != CL_SUCCESS)
6273 (void) ThrowMagickException(exception, GetMagickModule(),
6274 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
6277 for (i = 0; i < width; i++)
6279 offsetBufferPtr[2*i] = (int)offset[i].x;
6280 offsetBufferPtr[2*i+1] = (int)offset[i].y;
6282 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
6284 if (clStatus != CL_SUCCESS)
6286 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6287 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
6292 // get the OpenCL kernel
6293 motionBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
6295 if (motionBlurKernel == NULL)
6297 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6298 "AcquireOpenCLKernel failed.", "'%s'", ".");
6302 // set the kernel arguments
6304 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6305 (void *)&inputImageBuffer);
6306 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6307 (void *)&filteredImageBuffer);
6308 imageWidth = inputImage->columns;
6309 imageHeight = inputImage->rows;
6310 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
6312 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
6314 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6315 (void *)&imageKernelBuffer);
6316 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
6318 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
6319 (void *)&offsetBuffer);
6321 GetMagickPixelPacket(inputImage,&bias);
6322 biasPixel.s[0] = bias.red;
6323 biasPixel.s[1] = bias.green;
6324 biasPixel.s[2] = bias.blue;
6325 biasPixel.s[3] = bias.opacity;
6326 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_float4), &biasPixel);
6328 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(ChannelType), &channel);
6329 matte = (inputImage->matte == MagickTrue)?1:0;
6330 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), &matte);
6331 if (clStatus != CL_SUCCESS)
6333 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6334 "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
6338 // launch the kernel
6339 local_work_size[0] = 16;
6340 local_work_size[1] = 16;
6341 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
6342 inputImage->columns,local_work_size[0]);
6343 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
6344 inputImage->rows,local_work_size[1]);
6345 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
6346 global_work_size, local_work_size, 0, NULL, NULL);
6348 if (clStatus != CL_SUCCESS)
6350 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6351 "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
6354 clEnv->library->clFlush(queue);
6356 if (ALIGNED(filteredPixels,CLPixelPacket))
6358 length = inputImage->columns * inputImage->rows;
6359 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE,
6360 CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL,
6365 length = inputImage->columns * inputImage->rows;
6366 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0,
6367 length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
6369 if (clStatus != CL_SUCCESS)
6371 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
6372 "Reading output image from CL buffer failed.", "'%s'", ".");
6375 outputReady = MagickTrue;
6379 if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
6380 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
6381 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
6382 if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, motionBlurKernel);
6383 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
6384 if (outputReady == MagickFalse)
6386 if (filteredImage != NULL)
6388 DestroyImage(filteredImage);
6389 filteredImage = NULL;
6393 return filteredImage;
6398 Image* AccelerateMotionBlurImage(const Image *image, const ChannelType channel,
6399 const double* kernel, const size_t width, const OffsetInfo *offset,
6400 ExceptionInfo *exception)
6402 MagickBooleanType status;
6403 Image* filteredImage = NULL;
6405 assert(image != NULL);
6406 assert(kernel != (double *) NULL);
6407 assert(offset != (OffsetInfo *) NULL);
6408 assert(exception != (ExceptionInfo *) NULL);
6410 status = checkOpenCLEnvironment(exception);
6411 if (status == MagickFalse)
6414 status = checkAccelerateCondition(image, channel);
6415 if (status == MagickFalse)
6418 filteredImage = ComputeMotionBlurImage(image, channel, kernel, width,
6420 return filteredImage;
6425 static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv,
6426 cl_command_queue queue,
6427 cl_mem inputImageBuffer,
6428 const unsigned int inputWidth, const unsigned int inputHeight,
6429 const unsigned int matte,
6430 const ChannelType channel,const CompositeOperator compose,
6431 const cl_mem compositeImageBuffer,
6432 const unsigned int compositeWidth,
6433 const unsigned int compositeHeight,
6434 const float destination_dissolve,const float source_dissolve,
6435 ExceptionInfo *magick_unused(exception))
6437 size_t global_work_size[2];
6438 size_t local_work_size[2];
6439 unsigned int composeOp;
6443 cl_kernel compositeKernel = NULL;
6445 magick_unreferenced(exception);
6447 compositeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
6451 clStatus=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&inputImageBuffer);
6452 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputWidth);
6453 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&inputHeight);
6454 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(cl_mem),(void*)&compositeImageBuffer);
6455 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeWidth);
6456 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&compositeHeight);
6457 composeOp = (unsigned int)compose;
6458 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&composeOp);
6459 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(ChannelType),(void*)&channel);
6460 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(unsigned int),(void*)&matte);
6461 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&destination_dissolve);
6462 clStatus|=clEnv->library->clSetKernelArg(compositeKernel,k++,sizeof(float),(void*)&source_dissolve);
6464 if (clStatus!=CL_SUCCESS)
6467 local_work_size[0] = 64;
6468 local_work_size[1] = 1;
6470 global_work_size[0] = padGlobalWorkgroupSizeToLocalWorkgroupSize(inputWidth,
6471 local_work_size[0]);
6472 global_work_size[1] = inputHeight;
6473 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL,
6474 global_work_size, local_work_size, 0, NULL, NULL);
6477 RelinquishOpenCLKernel(clEnv, compositeKernel);
6479 return (clStatus==CL_SUCCESS)?MagickTrue:MagickFalse;
6483 static MagickBooleanType ComputeCompositeImage(Image *inputImage,
6484 const ChannelType channel,const CompositeOperator compose,
6485 const Image *compositeImage,const ssize_t magick_unused(x_offset),const ssize_t magick_unused(y_offset),
6486 const float destination_dissolve,const float source_dissolve,
6487 ExceptionInfo *exception)
6489 MagickBooleanType status = MagickFalse;
6491 MagickBooleanType outputReady = MagickFalse;
6492 MagickCLEnv clEnv = NULL;
6496 void *inputPixels = NULL;
6497 const void *composePixels = NULL;
6498 MagickSizeType length;
6500 cl_mem_flags mem_flags;
6501 cl_context context = NULL;
6502 cl_mem inputImageBuffer = NULL;
6503 cl_mem compositeImageBuffer = NULL;
6504 cl_command_queue queue = NULL;
6506 magick_unreferenced(x_offset);
6507 magick_unreferenced(y_offset);
6509 clEnv = GetDefaultOpenCLEnv();
6510 context = GetOpenCLContext(clEnv);
6511 queue = AcquireOpenCLCommandQueue(clEnv);
6513 /* Create and initialize OpenCL buffers. */
6514 inputPixels = GetPixelCachePixels(inputImage, &length, exception);
6515 if (inputPixels == (void *) NULL)
6517 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,
6518 "UnableToReadPixelCache.","`%s'",inputImage->filename);
6522 /* If the host pointer is aligned to the size of CLPixelPacket,
6523 then use the host buffer directly from the GPU; otherwise,
6524 create a buffer on the GPU and copy the data over */
6525 if (ALIGNED(inputPixels,CLPixelPacket))
6527 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
6531 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
6533 /* create a CL buffer from image pixel buffer */
6534 length = inputImage->columns * inputImage->rows;
6535 inputImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6536 length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
6537 if (clStatus != CL_SUCCESS)
6539 (void) OpenCLThrowMagickException(exception, GetMagickModule(),
6540 ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6545 /* Create and initialize OpenCL buffers. */
6546 composePixels = AcquirePixelCachePixels(compositeImage, &length, exception);
6547 if (composePixels == (void *) NULL)
6549 (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,
6550 "UnableToReadPixelCache.","`%s'",compositeImage->filename);
6554 /* If the host pointer is aligned to the size of CLPixelPacket,
6555 then use the host buffer directly from the GPU; otherwise,
6556 create a buffer on the GPU and copy the data over */
6557 if (ALIGNED(composePixels,CLPixelPacket))
6559 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
6563 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
6565 /* create a CL buffer from image pixel buffer */
6566 length = compositeImage->columns * compositeImage->rows;
6567 compositeImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags,
6568 length * sizeof(CLPixelPacket), (void*)composePixels, &clStatus);
6569 if (clStatus != CL_SUCCESS)
6571 (void) OpenCLThrowMagickException(exception, GetMagickModule(),
6572 ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
6576 status = LaunchCompositeKernel(clEnv,queue,inputImageBuffer,
6577 (unsigned int) inputImage->columns,
6578 (unsigned int) inputImage->rows,
6579 (unsigned int) inputImage->matte,
6580 channel, compose, compositeImageBuffer,
6581 (unsigned int) compositeImage->columns,
6582 (unsigned int) compositeImage->rows,
6583 destination_dissolve,source_dissolve,
6586 if (status==MagickFalse)
6589 length = inputImage->columns * inputImage->rows;
6590 if (ALIGNED(inputPixels,CLPixelPacket))
6592 clEnv->library->clEnqueueMapBuffer(queue, inputImageBuffer, CL_TRUE,
6593 CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL,
6598 clStatus = clEnv->library->clEnqueueReadBuffer(queue, inputImageBuffer, CL_TRUE, 0,
6599 length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
6601 if (clStatus==CL_SUCCESS)
6602 outputReady = MagickTrue;
6605 if (inputImageBuffer!=NULL) clEnv->library->clReleaseMemObject(inputImageBuffer);
6606 if (compositeImageBuffer!=NULL) clEnv->library->clReleaseMemObject(compositeImageBuffer);
6607 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
6614 MagickBooleanType AccelerateCompositeImage(Image *image,
6615 const ChannelType channel,const CompositeOperator compose,
6616 const Image *composite,const ssize_t x_offset,const ssize_t y_offset,
6617 const float destination_dissolve,const float source_dissolve,
6618 ExceptionInfo *exception)
6620 MagickBooleanType status;
6622 assert(image != NULL);
6623 assert(composite != NULL);
6624 assert(exception != (ExceptionInfo *) NULL);
6626 status = checkOpenCLEnvironment(exception);
6627 if (status == MagickFalse)
6630 status = checkAccelerateCondition(image, channel);
6631 if (status == MagickFalse)
6634 /* only support zero offset and
6635 images with the size for now */
6638 || image->columns!=composite->columns
6639 || image->rows!=composite->rows)
6643 case ColorDodgeCompositeOp:
6644 case BlendCompositeOp:
6647 // unsupported compose operator, quit
6651 status = ComputeCompositeImage(image,channel,compose,composite,
6652 x_offset,y_offset,destination_dissolve,source_dissolve,exception);
6659 #else /* MAGICKCORE_OPENCL_SUPPORT */
6661 MagickExport Image *AccelerateConvolveImageChannel(
6662 const Image *magick_unused(image),const ChannelType magick_unused(channel),
6663 const KernelInfo *magick_unused(kernel),
6664 ExceptionInfo *magick_unused(exception))
6666 magick_unreferenced(image);
6667 magick_unreferenced(channel);
6668 magick_unreferenced(kernel);
6669 magick_unreferenced(exception);
6674 MagickExport MagickBooleanType AccelerateFunctionImage(
6675 Image *magick_unused(image),const ChannelType magick_unused(channel),
6676 const MagickFunction magick_unused(function),
6677 const size_t magick_unused(number_parameters),
6678 const double *magick_unused(parameters),
6679 ExceptionInfo *magick_unused(exception))
6681 magick_unreferenced(image);
6682 magick_unreferenced(channel);
6683 magick_unreferenced(function);
6684 magick_unreferenced(number_parameters);
6685 magick_unreferenced(parameters);
6686 magick_unreferenced(exception);
6691 MagickExport Image *AccelerateBlurImage(const Image *magick_unused(image),
6692 const ChannelType magick_unused(channel),const double magick_unused(radius),
6693 const double magick_unused(sigma),ExceptionInfo *magick_unused(exception))
6695 magick_unreferenced(image);
6696 magick_unreferenced(channel);
6697 magick_unreferenced(radius);
6698 magick_unreferenced(sigma);
6699 magick_unreferenced(exception);
6704 MagickExport Image *AccelerateRadialBlurImage(
6705 const Image *magick_unused(image),const ChannelType magick_unused(channel),
6706 const double magick_unused(angle),ExceptionInfo *magick_unused(exception))
6708 magick_unreferenced(image);
6709 magick_unreferenced(channel);
6710 magick_unreferenced(angle);
6711 magick_unreferenced(exception);
6717 MagickExport Image *AccelerateUnsharpMaskImage(
6718 const Image *magick_unused(image),const ChannelType magick_unused(channel),
6719 const double magick_unused(radius),const double magick_unused(sigma),
6720 const double magick_unused(gain),const double magick_unused(threshold),
6721 ExceptionInfo *magick_unused(exception))
6723 magick_unreferenced(image);
6724 magick_unreferenced(channel);
6725 magick_unreferenced(radius);
6726 magick_unreferenced(sigma);
6727 magick_unreferenced(gain);
6728 magick_unreferenced(threshold);
6729 magick_unreferenced(exception);
6735 MagickBooleanType AccelerateCompositeImage(Image *image,
6736 const ChannelType channel,const CompositeOperator compose,
6737 const Image *composite,const ssize_t x_offset,const ssize_t y_offset,
6738 const float destination_dissolve,const float source_dissolve,
6739 ExceptionInfo *exception)
6741 magick_unreferenced(image);
6742 magick_unreferenced(channel);
6743 magick_unreferenced(compose);
6744 magick_unreferenced(composite);
6745 magick_unreferenced(x_offset);
6746 magick_unreferenced(y_offset);
6747 magick_unreferenced(destination_dissolve);
6748 magick_unreferenced(source_dissolve);
6749 magick_unreferenced(exception);
6755 MagickExport MagickBooleanType AccelerateContrastImage(
6756 Image* magick_unused(image),const MagickBooleanType magick_unused(sharpen),
6757 ExceptionInfo* magick_unused(exception))
6759 magick_unreferenced(image);
6760 magick_unreferenced(sharpen);
6761 magick_unreferenced(exception);
6766 MagickExport MagickBooleanType AccelerateContrastStretchImageChannel(
6767 Image * image, const ChannelType channel, const double black_point, const double white_point,
6768 ExceptionInfo* magick_unused(exception))
6770 magick_unreferenced(image);
6771 magick_unreferenced(channel);
6772 magick_unreferenced(black_point);
6773 magick_unreferenced(white_point);
6774 magick_unreferenced(exception);
6779 MagickExport MagickBooleanType AccelerateEqualizeImage(
6780 Image* magick_unused(image), const ChannelType magick_unused(channel),
6781 ExceptionInfo* magick_unused(exception))
6783 magick_unreferenced(image);
6784 magick_unreferenced(channel);
6785 magick_unreferenced(exception);
6790 MagickExport Image *AccelerateDespeckleImage(const Image* magick_unused(image),
6791 ExceptionInfo* magick_unused(exception))
6793 magick_unreferenced(image);
6794 magick_unreferenced(exception);
6799 MagickExport Image *AccelerateResizeImage(const Image* magick_unused(image),
6800 const size_t magick_unused(resizedColumns),
6801 const size_t magick_unused(resizedRows),
6802 const ResizeFilter* magick_unused(resizeFilter),
6803 ExceptionInfo *magick_unused(exception))
6805 magick_unreferenced(image);
6806 magick_unreferenced(resizedColumns);
6807 magick_unreferenced(resizedRows);
6808 magick_unreferenced(resizeFilter);
6809 magick_unreferenced(exception);
6815 MagickBooleanType AccelerateModulateImage(
6816 Image* image, double percent_brightness, double percent_hue,
6817 double percent_saturation, ColorspaceType colorspace, ExceptionInfo* exception)
6819 magick_unreferenced(image);
6820 magick_unreferenced(percent_brightness);
6821 magick_unreferenced(percent_hue);
6822 magick_unreferenced(percent_saturation);
6823 magick_unreferenced(colorspace);
6824 magick_unreferenced(exception);
6825 return(MagickFalse);
6829 MagickBooleanType AccelerateNegateImageChannel(
6830 Image* image, const ChannelType channel, const MagickBooleanType grayscale, ExceptionInfo* exception)
6832 magick_unreferenced(image);
6833 magick_unreferenced(channel);
6834 magick_unreferenced(grayscale);
6835 magick_unreferenced(exception);
6836 return(MagickFalse);
6840 MagickBooleanType AccelerateGrayscaleImage(
6841 Image* image, const PixelIntensityMethod method, ExceptionInfo* exception)
6843 magick_unreferenced(image);
6844 magick_unreferenced(method);
6845 magick_unreferenced(exception);
6846 return(MagickFalse);
6849 MagickExport Image *AccelerateAddNoiseImage(const Image *image,
6850 const ChannelType channel, const NoiseType noise_type,ExceptionInfo *exception)
6852 magick_unreferenced(image);
6853 magick_unreferenced(channel);
6854 magick_unreferenced(noise_type);
6855 magick_unreferenced(exception);
6860 MagickExport MagickBooleanType AccelerateRandomImage(Image* image, ExceptionInfo* exception)
6862 magick_unreferenced(image);
6863 magick_unreferenced(exception);
6868 Image* AccelerateMotionBlurImage(const Image *image, const ChannelType channel,
6869 const double* kernel, const size_t width,
6870 const OffsetInfo *offset,
6871 ExceptionInfo *exception)
6873 magick_unreferenced(image);
6874 magick_unreferenced(channel);
6875 magick_unreferenced(kernel);
6876 magick_unreferenced(width);
6877 magick_unreferenced(offset);
6878 magick_unreferenced(exception);
6882 #endif /* MAGICKCORE_OPENCL_SUPPORT */
6884 MagickExport MagickBooleanType AccelerateConvolveImage(
6885 const Image *magick_unused(image),const KernelInfo *magick_unused(kernel),
6886 Image *magick_unused(convolve_image),ExceptionInfo *magick_unused(exception))
6888 magick_unreferenced(image);
6889 magick_unreferenced(kernel);
6890 magick_unreferenced(convolve_image);
6891 magick_unreferenced(exception);
6893 /* legacy, do not use */
6894 return(MagickFalse);