44#include "MagickCore/studio.h"
45#include "MagickCore/accelerate-private.h"
46#include "MagickCore/artifact.h"
47#include "MagickCore/cache.h"
48#include "MagickCore/cache-private.h"
49#include "MagickCore/cache-view.h"
50#include "MagickCore/color-private.h"
51#include "MagickCore/delegate-private.h"
52#include "MagickCore/enhance.h"
53#include "MagickCore/exception.h"
54#include "MagickCore/exception-private.h"
55#include "MagickCore/gem.h"
56#include "MagickCore/image.h"
57#include "MagickCore/image-private.h"
58#include "MagickCore/linked-list.h"
59#include "MagickCore/list.h"
60#include "MagickCore/memory_.h"
61#include "MagickCore/monitor-private.h"
62#include "MagickCore/opencl.h"
63#include "MagickCore/opencl-private.h"
64#include "MagickCore/option.h"
65#include "MagickCore/pixel-accessor.h"
66#include "MagickCore/prepress.h"
67#include "MagickCore/quantize.h"
68#include "MagickCore/quantum-private.h"
69#include "MagickCore/random_.h"
70#include "MagickCore/random-private.h"
71#include "MagickCore/registry.h"
72#include "MagickCore/resize.h"
73#include "MagickCore/resize-private.h"
74#include "MagickCore/semaphore.h"
75#include "MagickCore/splay-tree.h"
76#include "MagickCore/statistic.h"
77#include "MagickCore/string_.h"
78#include "MagickCore/string-private.h"
79#include "MagickCore/token.h"
81#define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
82#define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
84#if defined(MAGICKCORE_OPENCL_SUPPORT)
89#define ALIGNED(pointer,type) ((((size_t)(pointer)) & (sizeof(type)-1)) == 0)
94static const ResizeWeightingFunctionType supportedResizeWeighting[] =
97 TriangleWeightingFunction,
98 HannWeightingFunction,
99 HammingWeightingFunction,
100 BlackmanWeightingFunction,
101 CubicBCWeightingFunction,
102 SincWeightingFunction,
103 SincFastWeightingFunction,
104 LastWeightingFunction
110static MagickBooleanType checkAccelerateCondition(
const Image* image)
113 if (image->storage_class != DirectClass)
117 if (image->colorspace != RGBColorspace &&
118 image->colorspace != sRGBColorspace &&
119 image->colorspace != LinearGRAYColorspace &&
120 image->colorspace != GRAYColorspace)
124 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
125 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
129 if (((image->channels & ReadMaskChannel) != 0) ||
130 ((image->channels & WriteMaskChannel) != 0) ||
131 ((image->channels & CompositeMaskChannel) != 0))
134 if (image->number_channels > 4)
138 if ((image->channel_mask != AllChannels) &&
139 (image->channel_mask > 0x7ffffff))
143 if (GetPixelChannelOffset(image,RedPixelChannel) != 0)
146 if (image->number_channels == 1)
150 if ((image->number_channels == 2) &&
151 (GetPixelChannelOffset(image,AlphaPixelChannel) == 1))
154 if (image->number_channels == 2)
158 if ((GetPixelChannelOffset(image,GreenPixelChannel) != 1) ||
159 (GetPixelChannelOffset(image,BluePixelChannel) != 2))
162 if (image->number_channels == 3)
166 if (GetPixelChannelOffset(image,AlphaPixelChannel) != 3)
172static MagickBooleanType checkAccelerateConditionRGBA(
const Image* image)
174 if (checkAccelerateCondition(image) == MagickFalse)
178 if (image->number_channels != 4)
181 if ((GetPixelRedTraits(image) == UndefinedPixelTrait) ||
182 (GetPixelGreenTraits(image) == UndefinedPixelTrait) ||
183 (GetPixelBlueTraits(image) == UndefinedPixelTrait) ||
184 (GetPixelAlphaTraits(image) == UndefinedPixelTrait))
190static MagickBooleanType checkPixelIntensity(
const Image *image,
191 const PixelIntensityMethod method)
194 if ((method == Rec601LumaPixelIntensityMethod) ||
195 (method == Rec709LumaPixelIntensityMethod))
197 if (image->colorspace == RGBColorspace)
201 if ((method == Rec601LuminancePixelIntensityMethod) ||
202 (method == Rec709LuminancePixelIntensityMethod))
204 if (image->colorspace == sRGBColorspace)
211static MagickBooleanType checkHistogramCondition(
const Image *image,
212 const PixelIntensityMethod method)
215 if ((image->channel_mask & SyncChannels) == 0)
218 return(checkPixelIntensity(image,method));
221static MagickCLEnv getOpenCLEnvironment(
ExceptionInfo* exception)
226 clEnv=GetCurrentOpenCLEnv();
227 if (clEnv == (MagickCLEnv) NULL)
228 return((MagickCLEnv) NULL);
230 if (clEnv->enabled == MagickFalse)
231 return((MagickCLEnv) NULL);
233 if (InitializeOpenCL(clEnv,exception) == MagickFalse)
234 return((MagickCLEnv) NULL);
244 if (((image->channel_mask & RedChannel) != 0) &&
245 ((image->channel_mask & GreenChannel) != 0) &&
246 ((image->channel_mask & BlueChannel) != 0) &&
247 ((image->channel_mask & AlphaChannel) != 0))
248 clone=CloneImage(image,0,0,MagickTrue,exception);
251 clone=CloneImage(image,0,0,MagickTrue,exception);
252 if (clone != (
Image *) NULL)
253 SyncImagePixelCache(clone,exception);
260inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
261 const unsigned int orgGlobalSize,
const unsigned int localGroupSize)
263 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
266static cl_mem createKernelInfo(MagickCLDevice device,
const double radius,
270 geometry[MagickPathExtent];
284 (void) FormatLocaleString(geometry,MagickPathExtent,
285 "blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
286 kernel=AcquireKernelInfo(geometry,exception);
289 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
290 ResourceLimitWarning,
"AcquireKernelInfo failed.",
".");
291 return((cl_mem) NULL);
293 kernelBufferPtr=(
float *) AcquireMagickMemory(kernel->width*
294 sizeof(*kernelBufferPtr));
295 if (kernelBufferPtr == (
float *) NULL)
297 kernel=DestroyKernelInfo(kernel);
298 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
299 ResourceLimitWarning,
"MemoryAllocationFailed.",
".");
300 return((cl_mem) NULL);
302 for (i = 0; i < (ssize_t) kernel->width; i++)
303 kernelBufferPtr[i]=(
float) kernel->values[i];
304 imageKernelBuffer=CreateOpenCLBuffer(device,CL_MEM_COPY_HOST_PTR |
305 CL_MEM_READ_ONLY,kernel->width*
sizeof(*kernelBufferPtr),kernelBufferPtr);
306 *width=(cl_uint) kernel->width;
307 kernelBufferPtr=(
float *) RelinquishMagickMemory(kernelBufferPtr);
308 kernel=DestroyKernelInfo(kernel);
309 if (imageKernelBuffer == (cl_mem) NULL)
310 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
311 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
312 return(imageKernelBuffer);
315static cl_int get32BitChannelValue(
const ChannelType channel)
317#if defined(MAGICKCORE_64BIT_CHANNEL_MASK_SUPPORT)
318 if (channel == AllChannels)
321 return((cl_int) channel);
324static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
325 MagickCLDevice device,cl_command_queue queue,cl_mem imageBuffer,
326 cl_mem histogramBuffer,
Image *image,
const ChannelType channel,
333 channel_mask=get32BitChannelValue(channel),
352 histogramKernel = NULL;
354 outputReady = MagickFalse;
355 colorspace = image->colorspace;
356 method = image->intensity;
359 histogramKernel = AcquireOpenCLKernel(device,
"Histogram");
360 if (histogramKernel == NULL)
362 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
368 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
369 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_int),&channel_mask);
370 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_uint),&colorspace);
371 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_uint),&method);
372 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_mem),(
void *)&histogramBuffer);
373 if (clStatus != CL_SUCCESS)
375 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
380 global_work_size[0] = image->columns;
381 global_work_size[1] = image->rows;
383 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
385 if (clStatus != CL_SUCCESS)
387 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
390 RecordProfileData(device,histogramKernel,event);
392 outputReady = MagickTrue;
396 if (histogramKernel!=NULL)
397 ReleaseOpenCLKernel(histogramKernel);
414static Image *ComputeBlurImage(
const Image* image,MagickCLEnv clEnv,
415 const double radius,
const double sigma,
ExceptionInfo *exception)
421 channel_mask=get32BitChannelValue(image->channel_mask),
460 filteredImageBuffer=NULL;
461 tempImageBuffer=NULL;
462 imageKernelBuffer=NULL;
464 blurColumnKernel=NULL;
465 outputReady=MagickFalse;
467 assert(image != (
Image *) NULL);
468 assert(image->signature == MagickCoreSignature);
469 if (IsEventLogging() != MagickFalse)
470 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
472 device=RequestOpenCLDevice(clEnv);
473 queue=AcquireOpenCLCommandQueue(device);
474 filteredImage=cloneImage(image,exception);
475 if (filteredImage == (
Image *) NULL)
477 if (filteredImage->number_channels != image->number_channels)
479 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
480 if (imageBuffer == (cl_mem) NULL)
482 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
483 if (filteredImageBuffer == (cl_mem) NULL)
486 imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
488 if (imageKernelBuffer == (cl_mem) NULL)
491 length=image->columns*image->rows;
492 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
493 sizeof(cl_float4),(
void *) NULL);
494 if (tempImageBuffer == (cl_mem) NULL)
497 blurRowKernel=AcquireOpenCLKernel(device,
"BlurRow");
498 if (blurRowKernel == (cl_kernel) NULL)
500 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
501 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
505 number_channels=(cl_uint) image->number_channels;
506 imageColumns=(cl_uint) image->columns;
507 imageRows=(cl_uint) image->rows;
510 status =SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
511 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),&number_channels);
512 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_int),&channel_mask);
513 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
514 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
515 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
516 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
517 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_float4)*(chunkSize+kernelWidth),(
void *) NULL);
518 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
519 if (status != CL_SUCCESS)
521 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
522 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
526 gsize[0]=chunkSize*((image->columns+chunkSize-1)/chunkSize);
527 gsize[1]=image->rows;
531 outputReady=EnqueueOpenCLKernel(queue,blurRowKernel,2,(
size_t *) NULL,gsize,
532 lsize,image,filteredImage,MagickFalse,exception);
533 if (outputReady == MagickFalse)
536 blurColumnKernel=AcquireOpenCLKernel(device,
"BlurColumn");
537 if (blurColumnKernel == (cl_kernel) NULL)
539 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
540 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
545 status =SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
546 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),&number_channels);
547 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_int),&channel_mask);
548 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
549 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
550 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
551 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
552 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_float4)*(chunkSize+kernelWidth),(
void *) NULL);
553 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
554 if (status != CL_SUCCESS)
556 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
557 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
561 gsize[0]=image->columns;
562 gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize);
566 outputReady=EnqueueOpenCLKernel(queue,blurColumnKernel,2,(
size_t *) NULL,gsize,
567 lsize,image,filteredImage,MagickFalse,exception);
571 if (imageBuffer != (cl_mem) NULL)
572 ReleaseOpenCLMemObject(imageBuffer);
573 if (filteredImageBuffer != (cl_mem) NULL)
574 ReleaseOpenCLMemObject(filteredImageBuffer);
575 if (tempImageBuffer != (cl_mem) NULL)
576 ReleaseOpenCLMemObject(tempImageBuffer);
577 if (imageKernelBuffer != (cl_mem) NULL)
578 ReleaseOpenCLMemObject(imageKernelBuffer);
579 if (blurRowKernel != (cl_kernel) NULL)
580 ReleaseOpenCLKernel(blurRowKernel);
581 if (blurColumnKernel != (cl_kernel) NULL)
582 ReleaseOpenCLKernel(blurColumnKernel);
583 if (queue != (cl_command_queue) NULL)
584 ReleaseOpenCLCommandQueue(device,queue);
585 if (device != (MagickCLDevice) NULL)
586 ReleaseOpenCLDevice(device);
587 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
588 filteredImage=DestroyImage(filteredImage);
590 return(filteredImage);
593MagickPrivate
Image* AccelerateBlurImage(
const Image *image,
594 const double radius,
const double sigma,
ExceptionInfo *exception)
602 assert(image != NULL);
604 if (IsEventLogging() != MagickFalse)
605 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
607 if (checkAccelerateCondition(image) == MagickFalse)
608 return((
Image *) NULL);
610 clEnv=getOpenCLEnvironment(exception);
611 if (clEnv == (MagickCLEnv) NULL)
612 return((
Image *) NULL);
614 filteredImage=ComputeBlurImage(image,clEnv,radius,sigma,exception);
615 return(filteredImage);
630static MagickBooleanType ComputeContrastImage(
Image *image,MagickCLEnv clEnv,
659 assert(image != (
Image *) NULL);
660 assert(image->signature == MagickCoreSignature);
661 if (IsEventLogging() != MagickFalse)
662 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
666 outputReady=MagickFalse;
668 device=RequestOpenCLDevice(clEnv);
669 queue=AcquireOpenCLCommandQueue(device);
670 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
671 if (imageBuffer == (cl_mem) NULL)
674 contrastKernel=AcquireOpenCLKernel(device,
"Contrast");
675 if (contrastKernel == (cl_kernel) NULL)
677 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
678 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
682 number_channels=(cl_uint) image->number_channels;
683 sign=sharpen != MagickFalse ? 1 : -1;
686 status =SetOpenCLKernelArg(contrastKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
687 status|=SetOpenCLKernelArg(contrastKernel,i++,
sizeof(cl_uint),&number_channels);
688 status|=SetOpenCLKernelArg(contrastKernel,i++,
sizeof(cl_int),&sign);
689 if (status != CL_SUCCESS)
691 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
692 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
696 gsize[0]=image->columns;
697 gsize[1]=image->rows;
699 outputReady=EnqueueOpenCLKernel(queue,contrastKernel,2,(
const size_t *) NULL,
700 gsize,(
const size_t *) NULL,image,(
Image *) NULL,MagickFalse,exception);
704 if (imageBuffer != (cl_mem) NULL)
705 ReleaseOpenCLMemObject(imageBuffer);
706 if (contrastKernel != (cl_kernel) NULL)
707 ReleaseOpenCLKernel(contrastKernel);
708 if (queue != (cl_command_queue) NULL)
709 ReleaseOpenCLCommandQueue(device,queue);
710 if (device != (MagickCLDevice) NULL)
711 ReleaseOpenCLDevice(device);
716MagickPrivate MagickBooleanType AccelerateContrastImage(
Image *image,
725 assert(image != NULL);
727 if (IsEventLogging() != MagickFalse)
728 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
730 if (checkAccelerateCondition(image) == MagickFalse)
733 clEnv=getOpenCLEnvironment(exception);
734 if (clEnv == (MagickCLEnv) NULL)
737 status=ComputeContrastImage(image,clEnv,sharpen,exception);
753static MagickBooleanType ComputeContrastStretchImage(
Image *image,
754 MagickCLEnv clEnv,
const double black_point,
const double white_point,
757#define ContrastStretchImageTag "ContrastStretch/Image"
758#define MaxRange(color) ((cl_float) ScaleQuantumToMap((Quantum) (color)))
767 channel_mask=get32BitChannelValue(image->channel_mask),
818 assert(image != (
Image *) NULL);
819 assert(image->signature == MagickCoreSignature);
820 if (IsEventLogging() != MagickFalse)
821 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
827 histogramBuffer = NULL;
828 stretchMapBuffer = NULL;
829 histogramKernel = NULL;
830 stretchKernel = NULL;
832 outputReady = MagickFalse;
839 device = RequestOpenCLDevice(clEnv);
840 queue = AcquireOpenCLCommandQueue(device);
845 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*histogram));
847 if (histogram == (cl_uint4 *) NULL)
848 ThrowBinaryException(ResourceLimitError,
"MemoryAllocationFailed", image->filename);
851 (void) memset(histogram,0,(MaxMap+1)*
sizeof(*histogram));
867 image_view=AcquireAuthenticCacheView(image,exception);
868 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
870 if (inputPixels == (
void *) NULL)
872 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
878 if (ALIGNED(inputPixels,CLPixelPacket))
880 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
884 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
887 length = image->columns * image->rows;
888 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
889 if (clStatus != CL_SUCCESS)
891 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
898 if (ALIGNED(histogram,cl_uint4))
900 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
905 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
910 histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(cl_uint4), hostPtr, &clStatus);
911 if (clStatus != CL_SUCCESS)
913 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
917 status = LaunchHistogramKernel(clEnv, device, queue, imageBuffer, histogramBuffer, image, image->channel_mask,exception);
918 if (status == MagickFalse)
922 if (ALIGNED(histogram,cl_uint4))
925 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
930 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length *
sizeof(cl_uint4), histogram, 0, NULL, NULL);
932 if (clStatus != CL_SUCCESS)
934 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
939 if (ALIGNED(histogram,cl_uint4))
941 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
942 if (clStatus != CL_SUCCESS)
944 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
".");
951 if (imageBuffer!=NULL)
952 clEnv->library->clReleaseMemObject(imageBuffer);
960 white.x=MaxRange(QuantumRange);
961 if ((image->channel_mask & RedChannel) != 0)
964 for (i=0; i <= (ssize_t) MaxMap; i++)
966 intensity+=histogram[i].s[2];
967 if (intensity > black_point)
970 black.x=(cl_float) i;
972 for (i=(ssize_t) MaxMap; i != 0; i--)
974 intensity+=histogram[i].s[2];
975 if (intensity > ((
double) image->columns*image->rows-white_point))
978 white.x=(cl_float) i;
981 white.y=MaxRange(QuantumRange);
982 if ((image->channel_mask & GreenChannel) != 0)
985 for (i=0; i <= (ssize_t) MaxMap; i++)
987 intensity+=histogram[i].s[2];
988 if (intensity > black_point)
991 black.y=(cl_float) i;
993 for (i=(ssize_t) MaxMap; i != 0; i--)
995 intensity+=histogram[i].s[2];
996 if (intensity > ((
double) image->columns*image->rows-white_point))
999 white.y=(cl_float) i;
1002 white.z=MaxRange(QuantumRange);
1003 if ((image->channel_mask & BlueChannel) != 0)
1006 for (i=0; i <= (ssize_t) MaxMap; i++)
1008 intensity+=histogram[i].s[2];
1009 if (intensity > black_point)
1012 black.z=(cl_float) i;
1014 for (i=(ssize_t) MaxMap; i != 0; i--)
1016 intensity+=histogram[i].s[2];
1017 if (intensity > ((
double) image->columns*image->rows-white_point))
1020 white.z=(cl_float) i;
1023 white.w=MaxRange(QuantumRange);
1024 if ((image->channel_mask & AlphaChannel) != 0)
1027 for (i=0; i <= (ssize_t) MaxMap; i++)
1029 intensity+=histogram[i].s[2];
1030 if (intensity > black_point)
1033 black.w=(cl_float) i;
1035 for (i=(ssize_t) MaxMap; i != 0; i--)
1037 intensity+=histogram[i].s[2];
1038 if (intensity > ((
double) image->columns*image->rows-white_point))
1041 white.w=(cl_float) i;
1044 stretch_map=(
PixelPacket *) AcquireQuantumMemory(MaxMap+1UL,
1045 sizeof(*stretch_map));
1048 ThrowBinaryException(ResourceLimitError,
"MemoryAllocationFailed",
1054 (void) memset(stretch_map,0,(MaxMap+1)*
sizeof(*stretch_map));
1055 for (i=0; i <= (ssize_t) MaxMap; i++)
1057 if ((image->channel_mask & RedChannel) != 0)
1059 if (i < (ssize_t) black.x)
1060 stretch_map[i].red=(Quantum) 0;
1062 if (i > (ssize_t) white.x)
1063 stretch_map[i].red=QuantumRange;
1065 if (black.x != white.x)
1066 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap*
1067 (i-black.x)/(white.x-black.x)));
1069 if ((image->channel_mask & GreenChannel) != 0)
1071 if (i < (ssize_t) black.y)
1072 stretch_map[i].green=0;
1074 if (i > (ssize_t) white.y)
1075 stretch_map[i].green=QuantumRange;
1077 if (black.y != white.y)
1078 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap*
1079 (i-black.y)/(white.y-black.y)));
1081 if ((image->channel_mask & BlueChannel) != 0)
1083 if (i < (ssize_t) black.z)
1084 stretch_map[i].blue=0;
1086 if (i > (ssize_t) white.z)
1087 stretch_map[i].blue= QuantumRange;
1089 if (black.z != white.z)
1090 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap*
1091 (i-black.z)/(white.z-black.z)));
1093 if ((image->channel_mask & AlphaChannel) != 0)
1095 if (i < (ssize_t) black.w)
1096 stretch_map[i].alpha=0;
1098 if (i > (ssize_t) white.w)
1099 stretch_map[i].alpha=QuantumRange;
1101 if (black.w != white.w)
1102 stretch_map[i].alpha=ScaleMapToQuantum((MagickRealType) (MaxMap*
1103 (i-black.w)/(white.w-black.w)));
1110 if (((image->channel_mask & AlphaChannel) != 0) || (((image->channel_mask & IndexChannel) != 0) &&
1111 (image->colorspace == CMYKColorspace)))
1112 image->storage_class=DirectClass;
1113 if (image->storage_class == PseudoClass)
1118 for (i=0; i < (ssize_t) image->colors; i++)
1120 if ((image->channel_mask & RedChannel) != 0)
1122 if (black.x != white.x)
1123 image->colormap[i].red=stretch_map[
1124 ScaleQuantumToMap(image->colormap[i].red)].red;
1126 if ((image->channel_mask & GreenChannel) != 0)
1128 if (black.y != white.y)
1129 image->colormap[i].green=stretch_map[
1130 ScaleQuantumToMap(image->colormap[i].green)].green;
1132 if ((image->channel_mask & BlueChannel) != 0)
1134 if (black.z != white.z)
1135 image->colormap[i].blue=stretch_map[
1136 ScaleQuantumToMap(image->colormap[i].blue)].blue;
1138 if ((image->channel_mask & AlphaChannel) != 0)
1140 if (black.w != white.w)
1141 image->colormap[i].alpha=stretch_map[
1142 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
1157#ifdef RECREATEBUFFER
1161 if (ALIGNED(inputPixels,CLPixelPacket))
1163 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
1167 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
1170 length = image->columns * image->rows;
1171 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1172 if (clStatus != CL_SUCCESS)
1174 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1182 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1183 hostPtr = stretch_map;
1187 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
1188 hostPtr = stretch_map;
1191 length = (MaxMap+1);
1192 stretchMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(
PixelPacket), hostPtr, &clStatus);
1193 if (clStatus != CL_SUCCESS)
1195 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1200 stretchKernel = AcquireOpenCLKernel(device,
"ContrastStretch");
1201 if (stretchKernel == NULL)
1203 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
1209 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1210 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_int),&channel_mask);
1211 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_mem),(
void *)&stretchMapBuffer);
1212 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_float4),&white);
1213 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_float4),&black);
1214 if (clStatus != CL_SUCCESS)
1216 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1221 global_work_size[0] = image->columns;
1222 global_work_size[1] = image->rows;
1224 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1226 if (clStatus != CL_SUCCESS)
1228 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1231 RecordProfileData(device,stretchKernel,event);
1234 if (ALIGNED(inputPixels,CLPixelPacket))
1236 length = image->columns * image->rows;
1237 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1241 length = image->columns * image->rows;
1242 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
1244 if (clStatus != CL_SUCCESS)
1246 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
1250 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
1254 image_view=DestroyCacheView(image_view);
1256 if (imageBuffer!=NULL)
1257 clEnv->library->clReleaseMemObject(imageBuffer);
1259 if (stretchMapBuffer!=NULL)
1260 clEnv->library->clReleaseMemObject(stretchMapBuffer);
1261 if (stretch_map!=NULL)
1262 stretch_map=(
PixelPacket *) RelinquishMagickMemory(stretch_map);
1263 if (histogramBuffer!=NULL)
1264 clEnv->library->clReleaseMemObject(histogramBuffer);
1265 if (histogram!=NULL)
1266 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
1267 if (histogramKernel!=NULL)
1268 ReleaseOpenCLKernel(histogramKernel);
1269 if (stretchKernel!=NULL)
1270 ReleaseOpenCLKernel(stretchKernel);
1272 ReleaseOpenCLCommandQueue(device,queue);
1274 ReleaseOpenCLDevice(device);
1276 return(outputReady);
1279MagickPrivate MagickBooleanType AccelerateContrastStretchImage(
1280 Image *image,
const double black_point,
const double white_point,
1289 assert(image != NULL);
1291 if (IsEventLogging() != MagickFalse)
1292 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
1294 if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
1295 (checkHistogramCondition(image,image->intensity) == MagickFalse))
1296 return(MagickFalse);
1298 clEnv=getOpenCLEnvironment(exception);
1299 if (clEnv == (MagickCLEnv) NULL)
1300 return(MagickFalse);
1302 status=ComputeContrastStretchImage(image,clEnv,black_point,white_point,
1319static Image *ComputeDespeckleImage(
const Image *image,MagickCLEnv clEnv,
1323 X[4] = {0, 1, 1,-1},
1324 Y[4] = {1, 0, 1, 1};
1327 *filteredImage_view,
1347 filteredImageBuffer,
1371 global_work_size[2];
1381 outputReady = MagickFalse;
1383 filteredImage = NULL;
1384 filteredImage_view = NULL;
1385 filteredPixels = NULL;
1387 filteredImageBuffer = NULL;
1391 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
1393 device = RequestOpenCLDevice(clEnv);
1394 queue = AcquireOpenCLCommandQueue(device);
1396 image_view=AcquireAuthenticCacheView(image,exception);
1397 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
1398 if (inputPixels == (
void *) NULL)
1400 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
1404 if (ALIGNED(inputPixels,CLPixelPacket))
1406 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1410 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1413 length = image->columns * image->rows;
1414 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1415 if (clStatus != CL_SUCCESS)
1417 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1421 mem_flags = CL_MEM_READ_WRITE;
1422 length = image->columns * image->rows;
1423 for (k = 0; k < 2; k++)
1425 tempImageBuffer[k] = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), NULL, &clStatus);
1426 if (clStatus != CL_SUCCESS)
1428 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1433 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
1434 assert(filteredImage != NULL);
1435 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1437 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"CloneImage failed.",
".");
1440 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
1441 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
1442 if (filteredPixels == (
void *) NULL)
1444 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
1448 if (ALIGNED(filteredPixels,CLPixelPacket))
1450 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1451 hostPtr = filteredPixels;
1455 mem_flags = CL_MEM_WRITE_ONLY;
1459 length = image->columns * image->rows;
1460 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
1461 if (clStatus != CL_SUCCESS)
1463 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1467 hullPass1 = AcquireOpenCLKernel(device,
"HullPass1");
1468 hullPass2 = AcquireOpenCLKernel(device,
"HullPass2");
1470 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,
sizeof(cl_mem),(
void *)&imageBuffer);
1471 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,
sizeof(cl_mem),(
void *)(tempImageBuffer+1));
1472 imageWidth = (
unsigned int) image->columns;
1473 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,
sizeof(
unsigned int),(
void *)&imageWidth);
1474 imageHeight = (
unsigned int) image->rows;
1475 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,
sizeof(
unsigned int),(
void *)&imageHeight);
1476 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
1477 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,
sizeof(
int),(
void *)&matte);
1478 if (clStatus != CL_SUCCESS)
1480 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1484 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,
sizeof(cl_mem),(
void *)(tempImageBuffer+1));
1485 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,
sizeof(cl_mem),(
void *)tempImageBuffer);
1486 imageWidth = (
unsigned int) image->columns;
1487 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,
sizeof(
unsigned int),(
void *)&imageWidth);
1488 imageHeight = (
unsigned int) image->rows;
1489 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,
sizeof(
unsigned int),(
void *)&imageHeight);
1490 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
1491 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,
sizeof(
int),(
void *)&matte);
1492 if (clStatus != CL_SUCCESS)
1494 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1499 global_work_size[0] = image->columns;
1500 global_work_size[1] = image->rows;
1503 for (k = 0; k < 4; k++)
1512 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
1513 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
1514 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
1515 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
1516 if (clStatus != CL_SUCCESS)
1518 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1522 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1523 if (clStatus != CL_SUCCESS)
1525 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1528 RecordProfileData(device,hullPass1,event);
1531 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1532 if (clStatus != CL_SUCCESS)
1534 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1537 RecordProfileData(device,hullPass2,event);
1540 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,
sizeof(cl_mem),(
void *)(tempImageBuffer));
1541 offset.s[0] = -X[k];
1542 offset.s[1] = -Y[k];
1544 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
1545 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
1546 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
1547 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
1548 if (clStatus != CL_SUCCESS)
1550 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1554 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1555 if (clStatus != CL_SUCCESS)
1557 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1560 RecordProfileData(device,hullPass1,event);
1563 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1564 if (clStatus != CL_SUCCESS)
1566 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1569 RecordProfileData(device,hullPass2,event);
1571 offset.s[0] = -X[k];
1572 offset.s[1] = -Y[k];
1574 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
1575 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
1576 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
1577 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
1578 if (clStatus != CL_SUCCESS)
1580 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1584 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1585 if (clStatus != CL_SUCCESS)
1587 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1590 RecordProfileData(device,hullPass1,event);
1593 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1594 if (clStatus != CL_SUCCESS)
1596 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1599 RecordProfileData(device,hullPass2,event);
1604 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
1605 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
1606 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
1607 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
1610 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
1612 if (clStatus != CL_SUCCESS)
1614 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1618 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1619 if (clStatus != CL_SUCCESS)
1621 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1624 RecordProfileData(device,hullPass1,event);
1627 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1628 if (clStatus != CL_SUCCESS)
1630 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1633 RecordProfileData(device,hullPass2,event);
1636 if (ALIGNED(filteredPixels,CLPixelPacket))
1638 length = image->columns * image->rows;
1639 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1643 length = image->columns * image->rows;
1644 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1646 if (clStatus != CL_SUCCESS)
1648 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
1652 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
1656 image_view=DestroyCacheView(image_view);
1657 if (filteredImage_view != NULL)
1658 filteredImage_view=DestroyCacheView(filteredImage_view);
1661 ReleaseOpenCLCommandQueue(device,queue);
1663 ReleaseOpenCLDevice(device);
1664 if (imageBuffer!=NULL)
1665 clEnv->library->clReleaseMemObject(imageBuffer);
1666 for (k = 0; k < 2; k++)
1668 if (tempImageBuffer[k]!=NULL)
1669 clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
1671 if (filteredImageBuffer!=NULL)
1672 clEnv->library->clReleaseMemObject(filteredImageBuffer);
1673 if (hullPass1!=NULL)
1674 ReleaseOpenCLKernel(hullPass1);
1675 if (hullPass2!=NULL)
1676 ReleaseOpenCLKernel(hullPass2);
1677 if (outputReady == MagickFalse && filteredImage != NULL)
1678 filteredImage=DestroyImage(filteredImage);
1680 return(filteredImage);
1683MagickPrivate
Image *AccelerateDespeckleImage(
const Image* image,
1692 assert(image != NULL);
1695 if (checkAccelerateConditionRGBA(image) == MagickFalse)
1696 return((
Image *) NULL);
1698 clEnv=getOpenCLEnvironment(exception);
1699 if (clEnv == (MagickCLEnv) NULL)
1700 return((
Image *) NULL);
1702 filteredImage=ComputeDespeckleImage(image,clEnv,exception);
1703 return(filteredImage);
1718static MagickBooleanType ComputeEqualizeImage(
Image *image,MagickCLEnv clEnv,
1721#define EqualizeImageTag "Equalize/Image"
1730 channel_mask=get32BitChannelValue(image->channel_mask),
1774 global_work_size[2];
1780 assert(image != (
Image *) NULL);
1781 assert(image->signature == MagickCoreSignature);
1782 if (IsEventLogging() != MagickFalse)
1783 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
1790 histogramBuffer = NULL;
1791 equalizeMapBuffer = NULL;
1792 histogramKernel = NULL;
1793 equalizeKernel = NULL;
1795 outputReady = MagickFalse;
1800 device = RequestOpenCLDevice(clEnv);
1801 queue = AcquireOpenCLCommandQueue(device);
1806 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*histogram));
1807 if (histogram == (cl_uint4 *) NULL)
1808 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
1811 (void) memset(histogram,0,(MaxMap+1)*
sizeof(*histogram));
1816 image_view=AcquireAuthenticCacheView(image,exception);
1817 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
1819 if (inputPixels == (
void *) NULL)
1821 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
1827 if (ALIGNED(inputPixels,CLPixelPacket))
1829 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1833 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1836 length = image->columns * image->rows;
1837 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1838 if (clStatus != CL_SUCCESS)
1840 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1847 if (ALIGNED(histogram,cl_uint4))
1849 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
1850 hostPtr = histogram;
1854 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
1855 hostPtr = histogram;
1858 length = (MaxMap+1);
1859 histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(cl_uint4), hostPtr, &clStatus);
1860 if (clStatus != CL_SUCCESS)
1862 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1866 status = LaunchHistogramKernel(clEnv, device, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception);
1867 if (status == MagickFalse)
1871 if (ALIGNED(histogram,cl_uint4))
1873 length = (MaxMap+1);
1874 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
1878 length = (MaxMap+1);
1879 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length *
sizeof(cl_uint4), histogram, 0, NULL, NULL);
1881 if (clStatus != CL_SUCCESS)
1883 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
1888 if (ALIGNED(histogram,cl_uint4))
1890 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
1891 if (clStatus != CL_SUCCESS)
1893 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
".");
1899#ifdef RECREATEBUFFER
1900 if (imageBuffer!=NULL)
1901 clEnv->library->clReleaseMemObject(imageBuffer);
1905 equalize_map=(
PixelPacket *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*equalize_map));
1907 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
1909 map=(cl_float4 *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*map));
1910 if (map == (cl_float4 *) NULL)
1911 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
1916 (void) memset(&intensity,0,
sizeof(intensity));
1917 for (i=0; i <= (ssize_t) MaxMap; i++)
1919 if ((image->channel_mask & SyncChannels) != 0)
1921 intensity.x+=histogram[i].s[2];
1925 if ((image->channel_mask & RedChannel) != 0)
1926 intensity.x+=histogram[i].s[2];
1927 if ((image->channel_mask & GreenChannel) != 0)
1928 intensity.y+=histogram[i].s[1];
1929 if ((image->channel_mask & BlueChannel) != 0)
1930 intensity.z+=histogram[i].s[0];
1931 if ((image->channel_mask & AlphaChannel) != 0)
1932 intensity.w+=histogram[i].s[3];
1936 white=map[(int) MaxMap];
1937 (void) memset(equalize_map,0,(MaxMap+1)*
sizeof(*equalize_map));
1938 for (i=0; i <= (ssize_t) MaxMap; i++)
1940 if ((image->channel_mask & SyncChannels) != 0)
1942 if (white.x != black.x)
1943 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
1944 (map[i].x-black.x))/(white.x-black.x)));
1947 if (((image->channel_mask & RedChannel) != 0) && (white.x != black.x))
1948 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
1949 (map[i].x-black.x))/(white.x-black.x)));
1950 if (((image->channel_mask & GreenChannel) != 0) && (white.y != black.y))
1951 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
1952 (map[i].y-black.y))/(white.y-black.y)));
1953 if (((image->channel_mask & BlueChannel) != 0) && (white.z != black.z))
1954 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
1955 (map[i].z-black.z))/(white.z-black.z)));
1956 if (((image->channel_mask & AlphaChannel) != 0) && (white.w != black.w))
1957 equalize_map[i].alpha=ScaleMapToQuantum((MagickRealType) ((MaxMap*
1958 (map[i].w-black.w))/(white.w-black.w)));
1961 if (image->storage_class == PseudoClass)
1966 for (i=0; i < (ssize_t) image->colors; i++)
1968 if ((image->channel_mask & SyncChannels) != 0)
1970 if (white.x != black.x)
1972 image->colormap[i].red=equalize_map[
1973 ScaleQuantumToMap(image->colormap[i].red)].red;
1974 image->colormap[i].green=equalize_map[
1975 ScaleQuantumToMap(image->colormap[i].green)].red;
1976 image->colormap[i].blue=equalize_map[
1977 ScaleQuantumToMap(image->colormap[i].blue)].red;
1978 image->colormap[i].alpha=equalize_map[
1979 ScaleQuantumToMap(image->colormap[i].alpha)].red;
1983 if (((image->channel_mask & RedChannel) != 0) && (white.x != black.x))
1984 image->colormap[i].red=equalize_map[
1985 ScaleQuantumToMap(image->colormap[i].red)].red;
1986 if (((image->channel_mask & GreenChannel) != 0) && (white.y != black.y))
1987 image->colormap[i].green=equalize_map[
1988 ScaleQuantumToMap(image->colormap[i].green)].green;
1989 if (((image->channel_mask & BlueChannel) != 0) && (white.z != black.z))
1990 image->colormap[i].blue=equalize_map[
1991 ScaleQuantumToMap(image->colormap[i].blue)].blue;
1992 if (((image->channel_mask & AlphaChannel) != 0) && (white.w != black.w))
1993 image->colormap[i].alpha=equalize_map[
1994 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
2007#ifdef RECREATEBUFFER
2011 if (ALIGNED(inputPixels,CLPixelPacket))
2013 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
2017 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
2020 length = image->columns * image->rows;
2021 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
2022 if (clStatus != CL_SUCCESS)
2024 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2032 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2033 hostPtr = equalize_map;
2037 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
2038 hostPtr = equalize_map;
2041 length = (MaxMap+1);
2042 equalizeMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(
PixelPacket), hostPtr, &clStatus);
2043 if (clStatus != CL_SUCCESS)
2045 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2050 equalizeKernel = AcquireOpenCLKernel(device,
"Equalize");
2051 if (equalizeKernel == NULL)
2053 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2059 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2060 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_int),&channel_mask);
2061 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_mem),(
void *)&equalizeMapBuffer);
2062 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_float4),&white);
2063 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_float4),&black);
2064 if (clStatus != CL_SUCCESS)
2066 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2071 global_work_size[0] = image->columns;
2072 global_work_size[1] = image->rows;
2074 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2076 if (clStatus != CL_SUCCESS)
2078 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2081 RecordProfileData(device,equalizeKernel,event);
2084 if (ALIGNED(inputPixels,CLPixelPacket))
2086 length = image->columns * image->rows;
2087 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2091 length = image->columns * image->rows;
2092 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
2094 if (clStatus != CL_SUCCESS)
2096 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
2100 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
2104 image_view=DestroyCacheView(image_view);
2106 if (imageBuffer!=NULL)
2107 clEnv->library->clReleaseMemObject(imageBuffer);
2109 map=(cl_float4 *) RelinquishMagickMemory(map);
2110 if (equalizeMapBuffer!=NULL)
2111 clEnv->library->clReleaseMemObject(equalizeMapBuffer);
2112 if (equalize_map!=NULL)
2113 equalize_map=(
PixelPacket *) RelinquishMagickMemory(equalize_map);
2114 if (histogramBuffer!=NULL)
2115 clEnv->library->clReleaseMemObject(histogramBuffer);
2116 if (histogram!=NULL)
2117 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
2118 if (histogramKernel!=NULL)
2119 ReleaseOpenCLKernel(histogramKernel);
2120 if (equalizeKernel!=NULL)
2121 ReleaseOpenCLKernel(equalizeKernel);
2123 ReleaseOpenCLCommandQueue(device, queue);
2125 ReleaseOpenCLDevice(device);
2127 return(outputReady);
2130MagickPrivate MagickBooleanType AccelerateEqualizeImage(
Image *image,
2139 assert(image != NULL);
2141 if (IsEventLogging() != MagickFalse)
2142 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2144 if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
2145 (checkHistogramCondition(image,image->intensity) == MagickFalse))
2146 return(MagickFalse);
2148 clEnv=getOpenCLEnvironment(exception);
2149 if (clEnv == (MagickCLEnv) NULL)
2150 return(MagickFalse);
2152 status=ComputeEqualizeImage(image,clEnv,exception);
2168static MagickBooleanType ComputeFunctionImage(
Image *image,MagickCLEnv clEnv,
2169 const MagickFunction function,
const size_t number_parameters,
2176 channel_mask=get32BitChannelValue(image->channel_mask),
2191 *parametersBufferPtr;
2203 assert(image != (
Image *) NULL);
2204 assert(image->signature == MagickCoreSignature);
2205 if (IsEventLogging() != MagickFalse)
2206 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2208 outputReady=MagickFalse;
2210 functionKernel=NULL;
2211 parametersBuffer=NULL;
2213 device=RequestOpenCLDevice(clEnv);
2214 queue=AcquireOpenCLCommandQueue(device);
2215 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
2216 if (imageBuffer == (cl_mem) NULL)
2219 parametersBufferPtr=(
float *) AcquireQuantumMemory(number_parameters,
2221 if (parametersBufferPtr == (
float *) NULL)
2223 for (i=0; i<number_parameters; i++)
2224 parametersBufferPtr[i]=(
float) parameters[i];
2225 parametersBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
2226 CL_MEM_COPY_HOST_PTR,number_parameters*
sizeof(*parametersBufferPtr),
2227 parametersBufferPtr);
2228 parametersBufferPtr=(
float *) RelinquishMagickMemory(parametersBufferPtr);
2229 if (parametersBuffer == (cl_mem) NULL)
2231 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2232 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
2236 functionKernel=AcquireOpenCLKernel(device,
"ComputeFunction");
2237 if (functionKernel == (cl_kernel) NULL)
2239 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2240 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2244 number_channels=(cl_uint) image->number_channels;
2245 number_params=(cl_uint) number_parameters;
2248 status =SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2249 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
2250 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_int),&channel_mask);
2251 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(MagickFunction),(
void *)&function);
2252 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_uint),(
void *)&number_params);
2253 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_mem),(
void *)¶metersBuffer);
2254 if (status != CL_SUCCESS)
2256 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2257 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
2261 gsize[0]=image->columns;
2262 gsize[1]=image->rows;
2263 outputReady=EnqueueOpenCLKernel(queue,functionKernel,2,(
const size_t *) NULL,
2264 gsize,(
const size_t *) NULL,image,(
const Image *) NULL,MagickFalse,
2269 if (imageBuffer != (cl_mem) NULL)
2270 ReleaseOpenCLMemObject(imageBuffer);
2271 if (parametersBuffer != (cl_mem) NULL)
2272 ReleaseOpenCLMemObject(parametersBuffer);
2273 if (functionKernel != (cl_kernel) NULL)
2274 ReleaseOpenCLKernel(functionKernel);
2275 if (queue != (cl_command_queue) NULL)
2276 ReleaseOpenCLCommandQueue(device,queue);
2277 if (device != (MagickCLDevice) NULL)
2278 ReleaseOpenCLDevice(device);
2279 return(outputReady);
2282MagickPrivate MagickBooleanType AccelerateFunctionImage(
Image *image,
2283 const MagickFunction function,
const size_t number_parameters,
2292 assert(image != NULL);
2294 if (IsEventLogging() != MagickFalse)
2295 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2297 if (checkAccelerateCondition(image) == MagickFalse)
2298 return(MagickFalse);
2300 clEnv=getOpenCLEnvironment(exception);
2301 if (clEnv == (MagickCLEnv) NULL)
2302 return(MagickFalse);
2304 status=ComputeFunctionImage(image,clEnv,function,number_parameters,
2305 parameters,exception);
2321static MagickBooleanType ComputeGrayscaleImage(
Image *image,MagickCLEnv clEnv,
2351 assert(image != (
Image *) NULL);
2352 assert(image->signature == MagickCoreSignature);
2353 if (IsEventLogging() != MagickFalse)
2354 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2356 outputReady=MagickFalse;
2358 grayscaleKernel=NULL;
2360 device=RequestOpenCLDevice(clEnv);
2361 queue=AcquireOpenCLCommandQueue(device);
2362 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
2363 if (imageBuffer == (cl_mem) NULL)
2366 grayscaleKernel=AcquireOpenCLKernel(device,
"Grayscale");
2367 if (grayscaleKernel == (cl_kernel) NULL)
2369 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2370 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2374 number_channels=(cl_uint) image->number_channels;
2375 intensityMethod=(cl_uint) method;
2376 colorspace=(cl_uint) image->colorspace;
2379 status =SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2380 status|=SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_uint),&number_channels);
2381 status|=SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_uint),&colorspace);
2382 status|=SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_uint),&intensityMethod);
2383 if (status != CL_SUCCESS)
2385 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2386 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
2390 gsize[0]=image->columns;
2391 gsize[1]=image->rows;
2392 outputReady=EnqueueOpenCLKernel(queue,grayscaleKernel,2,
2393 (
const size_t *) NULL,gsize,(
const size_t *) NULL,image,(
Image *) NULL,
2394 MagickFalse,exception);
2398 if (imageBuffer != (cl_mem) NULL)
2399 ReleaseOpenCLMemObject(imageBuffer);
2400 if (grayscaleKernel != (cl_kernel) NULL)
2401 ReleaseOpenCLKernel(grayscaleKernel);
2402 if (queue != (cl_command_queue) NULL)
2403 ReleaseOpenCLCommandQueue(device,queue);
2404 if (device != (MagickCLDevice) NULL)
2405 ReleaseOpenCLDevice(device);
2407 return(outputReady);
2410MagickPrivate MagickBooleanType AccelerateGrayscaleImage(
Image* image,
2419 assert(image != NULL);
2421 if (IsEventLogging() != MagickFalse)
2422 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2424 if ((checkAccelerateCondition(image) == MagickFalse) ||
2425 (checkPixelIntensity(image,method) == MagickFalse))
2426 return(MagickFalse);
2428 if (image->number_channels < 3)
2429 return(MagickFalse);
2431 if ((GetPixelRedTraits(image) == UndefinedPixelTrait) ||
2432 (GetPixelGreenTraits(image) == UndefinedPixelTrait) ||
2433 (GetPixelBlueTraits(image) == UndefinedPixelTrait))
2434 return(MagickFalse);
2436 clEnv=getOpenCLEnvironment(exception);
2437 if (clEnv == (MagickCLEnv) NULL)
2438 return(MagickFalse);
2440 status=ComputeGrayscaleImage(image,clEnv,method,exception);
2456static Image *ComputeLocalContrastImage(
const Image *image,MagickCLEnv clEnv,
2457 const double radius,
const double strength,
ExceptionInfo *exception)
2460 *filteredImage_view,
2478 filteredImageBuffer,
2511 filteredImage = NULL;
2512 filteredImage_view = NULL;
2514 filteredImageBuffer = NULL;
2515 tempImageBuffer = NULL;
2516 imageKernelBuffer = NULL;
2517 blurRowKernel = NULL;
2518 blurColumnKernel = NULL;
2520 outputReady = MagickFalse;
2522 device = RequestOpenCLDevice(clEnv);
2523 queue = AcquireOpenCLCommandQueue(device);
2527 image_view=AcquireAuthenticCacheView(image,exception);
2528 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
2529 if (inputPixels == (
const void *) NULL)
2531 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
2538 if (ALIGNED(inputPixels,CLPixelPacket))
2540 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2544 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2547 length = image->columns * image->rows;
2548 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
2549 if (clStatus != CL_SUCCESS)
2551 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2558 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
2559 assert(filteredImage != NULL);
2560 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
2562 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"CloneImage failed.",
".");
2565 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
2566 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
2567 if (filteredPixels == (
void *) NULL)
2569 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
2573 if (ALIGNED(filteredPixels,CLPixelPacket))
2575 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2576 hostPtr = filteredPixels;
2580 mem_flags = CL_MEM_WRITE_ONLY;
2585 length = image->columns * image->rows;
2586 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
2587 if (clStatus != CL_SUCCESS)
2589 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2597 length = image->columns * image->rows;
2598 tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length *
sizeof(
float), NULL, &clStatus);
2599 if (clStatus != CL_SUCCESS)
2601 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2608 blurRowKernel = AcquireOpenCLKernel(device,
"LocalContrastBlurRow");
2609 if (blurRowKernel == NULL)
2611 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2615 blurColumnKernel = AcquireOpenCLKernel(device,
"LocalContrastBlurApplyColumn");
2616 if (blurColumnKernel == NULL)
2618 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2624 imageColumns = (
unsigned int) image->columns;
2625 imageRows = (
unsigned int) image->rows;
2626 iRadius = (cl_int) (image->rows > image->columns ? image->rows : image->columns) * 0.002f * fabs(radius);
2628 passes = (((1.0f * imageRows) * imageColumns * iRadius) + 3999999999) / 4000000000.0f;
2629 passes = (passes < 1) ? 1: passes;
2633 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2634 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
2635 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
2636 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_int),(
void *)&iRadius);
2637 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
2638 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
2640 if (clStatus != CL_SUCCESS)
2642 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2650 for (x = 0; x < passes; ++x) {
2656 gsize[1] = (image->rows + passes - 1) / passes;
2660 goffset[1] = x * gsize[1];
2662 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, 0, NULL, &event);
2663 if (clStatus != CL_SUCCESS)
2665 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2668 clEnv->library->clFlush(queue);
2669 RecordProfileData(device,blurRowKernel,event);
2674 cl_float FStrength = strength;
2676 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2677 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
2678 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
2679 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&iRadius);
2680 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_float),(
void *)&FStrength);
2681 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
2682 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
2684 if (clStatus != CL_SUCCESS)
2686 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2694 for (x = 0; x < passes; ++x) {
2699 gsize[0] = ((image->columns + 3) / 4) * 4;
2700 gsize[1] = ((((image->rows + 63) / 64) + (passes + 1)) / passes) * 64;
2704 goffset[1] = x * gsize[1];
2706 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, 0, NULL, &event);
2707 if (clStatus != CL_SUCCESS)
2709 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2712 clEnv->library->clFlush(queue);
2713 RecordProfileData(device,blurColumnKernel,event);
2719 if (ALIGNED(filteredPixels,CLPixelPacket))
2721 length = image->columns * image->rows;
2722 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2726 length = image->columns * image->rows;
2727 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2729 if (clStatus != CL_SUCCESS)
2731 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
2735 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
2739 image_view=DestroyCacheView(image_view);
2740 if (filteredImage_view != NULL)
2741 filteredImage_view=DestroyCacheView(filteredImage_view);
2743 if (imageBuffer!=NULL)
2744 clEnv->library->clReleaseMemObject(imageBuffer);
2745 if (filteredImageBuffer!=NULL)
2746 clEnv->library->clReleaseMemObject(filteredImageBuffer);
2747 if (tempImageBuffer!=NULL)
2748 clEnv->library->clReleaseMemObject(tempImageBuffer);
2749 if (imageKernelBuffer!=NULL)
2750 clEnv->library->clReleaseMemObject(imageKernelBuffer);
2751 if (blurRowKernel!=NULL)
2752 ReleaseOpenCLKernel(blurRowKernel);
2753 if (blurColumnKernel!=NULL)
2754 ReleaseOpenCLKernel(blurColumnKernel);
2756 ReleaseOpenCLCommandQueue(device, queue);
2758 ReleaseOpenCLDevice(device);
2759 if (outputReady == MagickFalse)
2761 if (filteredImage != NULL)
2763 DestroyImage(filteredImage);
2764 filteredImage = NULL;
2768 return(filteredImage);
2771MagickPrivate
Image *AccelerateLocalContrastImage(
const Image *image,
2772 const double radius,
const double strength,
ExceptionInfo *exception)
2780 assert(image != NULL);
2783 if (checkAccelerateConditionRGBA(image) == MagickFalse)
2784 return((
Image *) NULL);
2786 clEnv=getOpenCLEnvironment(exception);
2787 if (clEnv == (MagickCLEnv) NULL)
2788 return((
Image *) NULL);
2790 filteredImage=ComputeLocalContrastImage(image,clEnv,radius,strength,
2792 return(filteredImage);
2807static MagickBooleanType ComputeModulateImage(
Image *image,MagickCLEnv clEnv,
2808 const double percent_brightness,
const double percent_hue,
2809 const double percent_saturation,
const ColorspaceType colorspace,
2854 assert(image != (
Image *) NULL);
2855 assert(image->signature == MagickCoreSignature);
2856 if (IsEventLogging() != MagickFalse)
2857 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2861 modulateKernel = NULL;
2866 device = RequestOpenCLDevice(clEnv);
2867 queue = AcquireOpenCLCommandQueue(device);
2869 outputReady = MagickFalse;
2875 image_view=AcquireAuthenticCacheView(image,exception);
2876 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
2877 if (inputPixels == (
void *) NULL)
2879 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
2887 if (ALIGNED(inputPixels,CLPixelPacket))
2889 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
2893 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
2896 length = image->columns * image->rows;
2897 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
2898 if (clStatus != CL_SUCCESS)
2900 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2904 modulateKernel = AcquireOpenCLKernel(device,
"Modulate");
2905 if (modulateKernel == NULL)
2907 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2911 bright=percent_brightness;
2913 saturation=percent_saturation;
2917 clStatus=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2918 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&bright);
2919 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&hue);
2920 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&saturation);
2921 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&color);
2922 if (clStatus != CL_SUCCESS)
2924 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2929 size_t global_work_size[2];
2930 global_work_size[0] = image->columns;
2931 global_work_size[1] = image->rows;
2933 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2934 if (clStatus != CL_SUCCESS)
2936 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2939 RecordProfileData(device,modulateKernel,event);
2942 if (ALIGNED(inputPixels,CLPixelPacket))
2944 length = image->columns * image->rows;
2945 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2949 length = image->columns * image->rows;
2950 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
2952 if (clStatus != CL_SUCCESS)
2954 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
2958 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
2962 image_view=DestroyCacheView(image_view);
2964 if (imageBuffer!=NULL)
2965 clEnv->library->clReleaseMemObject(imageBuffer);
2966 if (modulateKernel!=NULL)
2967 ReleaseOpenCLKernel(modulateKernel);
2969 ReleaseOpenCLCommandQueue(device,queue);
2971 ReleaseOpenCLDevice(device);
2977MagickPrivate MagickBooleanType AccelerateModulateImage(
Image *image,
2978 const double percent_brightness,
const double percent_hue,
2979 const double percent_saturation,
const ColorspaceType colorspace,
2988 assert(image != NULL);
2990 if (IsEventLogging() != MagickFalse)
2991 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2993 if (checkAccelerateConditionRGBA(image) == MagickFalse)
2994 return(MagickFalse);
2996 if ((colorspace != HSLColorspace) && (colorspace != UndefinedColorspace))
2997 return(MagickFalse);
2999 clEnv=getOpenCLEnvironment(exception);
3000 if (clEnv == (MagickCLEnv) NULL)
3001 return(MagickFalse);
3003 status=ComputeModulateImage(image,clEnv,percent_brightness,percent_hue,
3004 percent_saturation,colorspace,exception);
3020static Image* ComputeMotionBlurImage(
const Image *image,MagickCLEnv clEnv,
3021 const double *kernel,
const size_t width,
const OffsetInfo *offset,
3025 *filteredImage_view,
3035 channel_mask=get32BitChannelValue(image->channel_mask),
3045 filteredImageBuffer,
3078 global_work_size[2],
3091 assert(image != (
Image *) NULL);
3092 assert(image->signature == MagickCoreSignature);
3093 if (IsEventLogging() != MagickFalse)
3094 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
3096 outputReady = MagickFalse;
3097 filteredImage = NULL;
3098 filteredImage_view = NULL;
3100 filteredImageBuffer = NULL;
3101 imageKernelBuffer = NULL;
3102 motionBlurKernel = NULL;
3105 device = RequestOpenCLDevice(clEnv);
3109 image_view=AcquireAuthenticCacheView(image,exception);
3110 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,
3111 image->rows,exception);
3112 if (inputPixels == (
const void *) NULL)
3114 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
3115 "UnableToReadPixelCache.",
"`%s'",image->filename);
3124 if (ALIGNED(inputPixels,CLPixelPacket))
3126 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
3130 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
3135 length = image->columns * image->rows;
3136 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
3137 length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
3138 if (clStatus != CL_SUCCESS)
3140 (void) ThrowMagickException(exception, GetMagickModule(),
3141 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3146 filteredImage = CloneImage(image,image->columns,image->rows,
3147 MagickTrue,exception);
3148 assert(filteredImage != NULL);
3149 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
3151 (void) ThrowMagickException(exception, GetMagickModule(),
3152 ResourceLimitError,
"CloneImage failed.",
".");
3155 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
3156 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
3157 if (filteredPixels == (
void *) NULL)
3159 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
3160 "UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
3164 if (ALIGNED(filteredPixels,CLPixelPacket))
3166 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
3167 hostPtr = filteredPixels;
3171 mem_flags = CL_MEM_WRITE_ONLY;
3177 length = image->columns * image->rows;
3178 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
3179 length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
3180 if (clStatus != CL_SUCCESS)
3182 (void) ThrowMagickException(exception, GetMagickModule(),
3183 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3188 imageKernelBuffer = clEnv->library->clCreateBuffer(device->context,
3189 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width *
sizeof(
float), NULL,
3191 if (clStatus != CL_SUCCESS)
3193 (void) ThrowMagickException(exception, GetMagickModule(),
3194 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3198 queue = AcquireOpenCLCommandQueue(device);
3199 kernelBufferPtr = (
float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
3200 CL_TRUE, CL_MAP_WRITE, 0, width *
sizeof(
float), 0, NULL, NULL, &clStatus);
3201 if (clStatus != CL_SUCCESS)
3203 (void) ThrowMagickException(exception, GetMagickModule(),
3204 ResourceLimitError,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
3207 for (i = 0; i < width; i++)
3209 kernelBufferPtr[i] = (float) kernel[i];
3211 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr,
3213 if (clStatus != CL_SUCCESS)
3215 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3216 "clEnv->library->clEnqueueUnmapMemObject failed.",
".");
3220 offsetBuffer = clEnv->library->clCreateBuffer(device->context,
3221 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width *
sizeof(cl_int2), NULL,
3223 if (clStatus != CL_SUCCESS)
3225 (void) ThrowMagickException(exception, GetMagickModule(),
3226 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3230 offsetBufferPtr = (
int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
3231 CL_MAP_WRITE, 0, width *
sizeof(cl_int2), 0, NULL, NULL, &clStatus);
3232 if (clStatus != CL_SUCCESS)
3234 (void) ThrowMagickException(exception, GetMagickModule(),
3235 ResourceLimitError,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
3238 for (i = 0; i < width; i++)
3240 offsetBufferPtr[2*i] = (int)offset[i].x;
3241 offsetBufferPtr[2*i+1] = (int)offset[i].y;
3243 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
3245 if (clStatus != CL_SUCCESS)
3247 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3248 "clEnv->library->clEnqueueUnmapMemObject failed.",
".");
3256 motionBlurKernel = AcquireOpenCLKernel(device,
"MotionBlur");
3257 if (motionBlurKernel == NULL)
3259 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3260 "AcquireOpenCLKernel failed.",
".");
3268 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3269 (
void *)&imageBuffer);
3270 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3271 (
void *)&filteredImageBuffer);
3272 imageWidth = (
unsigned int) image->columns;
3273 imageHeight = (
unsigned int) image->rows;
3274 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3276 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3278 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3279 (
void *)&imageKernelBuffer);
3280 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3282 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3283 (
void *)&offsetBuffer);
3285 GetPixelInfo(image,&bias);
3286 biasPixel.s[0] = bias.red;
3287 biasPixel.s[1] = bias.green;
3288 biasPixel.s[2] = bias.blue;
3289 biasPixel.s[3] = bias.alpha;
3290 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_float4), &biasPixel);
3292 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_int),&channel_mask);
3293 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
3294 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int), &matte);
3295 if (clStatus != CL_SUCCESS)
3297 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3298 "clEnv->library->clSetKernelArg failed.",
".");
3305 local_work_size[0] = 16;
3306 local_work_size[1] = 16;
3307 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3308 (
unsigned int) image->columns,(
unsigned int) local_work_size[0]);
3309 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3310 (
unsigned int) image->rows,(
unsigned int) local_work_size[1]);
3311 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
3312 global_work_size, local_work_size, 0, NULL, &event);
3314 if (clStatus != CL_SUCCESS)
3316 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3317 "clEnv->library->clEnqueueNDRangeKernel failed.",
".");
3320 RecordProfileData(device,motionBlurKernel,event);
3322 if (ALIGNED(filteredPixels,CLPixelPacket))
3324 length = image->columns * image->rows;
3325 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE,
3326 CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL,
3331 length = image->columns * image->rows;
3332 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0,
3333 length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
3335 if (clStatus != CL_SUCCESS)
3337 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3338 "Reading output image from CL buffer failed.",
".");
3341 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
3345 image_view=DestroyCacheView(image_view);
3346 if (filteredImage_view != NULL)
3347 filteredImage_view=DestroyCacheView(filteredImage_view);
3349 if (filteredImageBuffer!=NULL)
3350 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3351 if (imageBuffer!=NULL)
3352 clEnv->library->clReleaseMemObject(imageBuffer);
3353 if (imageKernelBuffer!=NULL)
3354 clEnv->library->clReleaseMemObject(imageKernelBuffer);
3355 if (motionBlurKernel!=NULL)
3356 ReleaseOpenCLKernel(motionBlurKernel);
3358 ReleaseOpenCLCommandQueue(device,queue);
3360 ReleaseOpenCLDevice(device);
3361 if (outputReady == MagickFalse && filteredImage != NULL)
3362 filteredImage=DestroyImage(filteredImage);
3364 return(filteredImage);
3367MagickPrivate
Image *AccelerateMotionBlurImage(
const Image *image,
3368 const double* kernel,
const size_t width,
const OffsetInfo *offset,
3377 assert(image != NULL);
3378 assert(kernel != (
double *) NULL);
3382 if (checkAccelerateConditionRGBA(image) == MagickFalse)
3383 return((
Image *) NULL);
3385 clEnv=getOpenCLEnvironment(exception);
3386 if (clEnv == (MagickCLEnv) NULL)
3387 return((
Image *) NULL);
3389 filteredImage=ComputeMotionBlurImage(image,clEnv,kernel,width,offset,
3391 return(filteredImage);
3406static MagickBooleanType resizeHorizontalFilter(MagickCLDevice device,
3407 cl_command_queue queue,
const Image *image,
Image *filteredImage,
3408 cl_mem imageBuffer,cl_uint number_channels,cl_uint columns,cl_uint rows,
3409 cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows,
3410 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3420 workgroupSize = 256;
3424 resizeFilterSupport,
3425 resizeFilterWindowSupport,
3439 gammaAccumulatorLocalMemorySize,
3442 imageCacheLocalMemorySize,
3443 pixelAccumulatorLocalMemorySize,
3445 totalLocalMemorySize,
3446 weightAccumulatorLocalMemorySize;
3452 horizontalKernel=NULL;
3453 outputReady=MagickFalse;
3458 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
3459 support=scale*GetResizeFilterSupport(resizeFilter);
3466 support=(float) 0.5;
3469 scale=PerceptibleReciprocal(scale);
3471 if (resizedColumns < workgroupSize)
3474 pixelPerWorkgroup=32;
3478 chunkSize=workgroupSize;
3479 pixelPerWorkgroup=workgroupSize;
3482DisableMSCWarning(4127)
3487 numCachedPixels=(int) ceil((pixelPerWorkgroup-1)/xFactor+2*support);
3488 imageCacheLocalMemorySize=numCachedPixels*
sizeof(CLQuantum)*
3490 totalLocalMemorySize=imageCacheLocalMemorySize;
3493 pixelAccumulatorLocalMemorySize=chunkSize*
sizeof(cl_float4);
3494 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
3497 weightAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
3498 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
3501 if ((number_channels == 4) || (number_channels == 2))
3502 gammaAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
3504 gammaAccumulatorLocalMemorySize=
sizeof(float);
3505 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
3507 if (totalLocalMemorySize <= device->local_memory_size)
3511 pixelPerWorkgroup=pixelPerWorkgroup/2;
3512 chunkSize=chunkSize/2;
3513 if ((pixelPerWorkgroup == 0) || (chunkSize == 0))
3521 resizeFilterType=(int)GetResizeFilterWeightingType(resizeFilter);
3522 resizeWindowType=(int)GetResizeFilterWindowWeightingType(resizeFilter);
3524 horizontalKernel=AcquireOpenCLKernel(device,
"ResizeHorizontalFilter");
3525 if (horizontalKernel == (cl_kernel) NULL)
3527 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3528 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
3532 resizeFilterScale=(float) GetResizeFilterScale(resizeFilter);
3533 resizeFilterSupport=(float) GetResizeFilterSupport(resizeFilter);
3534 resizeFilterWindowSupport=(float) GetResizeFilterWindowSupport(resizeFilter);
3535 resizeFilterBlur=(float) GetResizeFilterBlur(resizeFilter);
3538 status =SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_mem),(
void*)&imageBuffer);
3539 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&number_channels);
3540 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&columns);
3541 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&rows);
3542 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_mem),(
void*)&resizedImageBuffer);
3543 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&resizedColumns);
3544 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&resizedRows);
3545 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&xFactor);
3546 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
int),(
void*)&resizeFilterType);
3547 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
int),(
void*)&resizeWindowType);
3548 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_mem),(
void*)&resizeFilterCubicCoefficients);
3549 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterScale);
3550 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterSupport);
3551 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterWindowSupport);
3552 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterBlur);
3553 status|=SetOpenCLKernelArg(horizontalKernel,i++,imageCacheLocalMemorySize,NULL);
3554 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
int),&numCachedPixels);
3555 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
unsigned int),&pixelPerWorkgroup);
3556 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
unsigned int),&chunkSize);
3557 status|=SetOpenCLKernelArg(horizontalKernel,i++,pixelAccumulatorLocalMemorySize,NULL);
3558 status|=SetOpenCLKernelArg(horizontalKernel,i++,weightAccumulatorLocalMemorySize,NULL);
3559 status|=SetOpenCLKernelArg(horizontalKernel,i++,gammaAccumulatorLocalMemorySize,NULL);
3561 if (status != CL_SUCCESS)
3563 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3564 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
3568 gsize[0]=(resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*
3570 gsize[1]=resizedRows;
3571 lsize[0]=workgroupSize;
3573 outputReady=EnqueueOpenCLKernel(queue,horizontalKernel,2,
3574 (
const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
3579 if (horizontalKernel != (cl_kernel) NULL)
3580 ReleaseOpenCLKernel(horizontalKernel);
3582 return(outputReady);
3585static MagickBooleanType resizeVerticalFilter(MagickCLDevice device,
3586 cl_command_queue queue,
const Image *image,
Image * filteredImage,
3587 cl_mem imageBuffer,cl_uint number_channels,cl_uint columns,cl_uint rows,
3588 cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows,
3589 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3599 workgroupSize = 256;
3603 resizeFilterSupport,
3604 resizeFilterWindowSupport,
3618 gammaAccumulatorLocalMemorySize,
3621 imageCacheLocalMemorySize,
3622 pixelAccumulatorLocalMemorySize,
3624 totalLocalMemorySize,
3625 weightAccumulatorLocalMemorySize;
3631 verticalKernel=NULL;
3632 outputReady=MagickFalse;
3637 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
3638 support=scale*GetResizeFilterSupport(resizeFilter);
3645 support=(float) 0.5;
3648 scale=PerceptibleReciprocal(scale);
3650 if (resizedRows < workgroupSize)
3653 pixelPerWorkgroup=32;
3657 chunkSize=workgroupSize;
3658 pixelPerWorkgroup=workgroupSize;
3661DisableMSCWarning(4127)
3666 numCachedPixels=(int)ceil((pixelPerWorkgroup-1)/yFactor+2*support);
3667 imageCacheLocalMemorySize=numCachedPixels*
sizeof(CLQuantum)*
3669 totalLocalMemorySize=imageCacheLocalMemorySize;
3672 pixelAccumulatorLocalMemorySize=chunkSize*
sizeof(cl_float4);
3673 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
3676 weightAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
3677 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
3680 if ((number_channels == 4) || (number_channels == 2))
3681 gammaAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
3683 gammaAccumulatorLocalMemorySize=
sizeof(float);
3684 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
3686 if (totalLocalMemorySize <= device->local_memory_size)
3690 pixelPerWorkgroup=pixelPerWorkgroup/2;
3691 chunkSize=chunkSize/2;
3692 if ((pixelPerWorkgroup == 0) || (chunkSize == 0))
3700 resizeFilterType=(int)GetResizeFilterWeightingType(resizeFilter);
3701 resizeWindowType=(int)GetResizeFilterWindowWeightingType(resizeFilter);
3703 verticalKernel=AcquireOpenCLKernel(device,
"ResizeVerticalFilter");
3704 if (verticalKernel == (cl_kernel) NULL)
3706 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3707 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
3711 resizeFilterScale=(float) GetResizeFilterScale(resizeFilter);
3712 resizeFilterSupport=(float) GetResizeFilterSupport(resizeFilter);
3713 resizeFilterBlur=(float) GetResizeFilterBlur(resizeFilter);
3714 resizeFilterWindowSupport=(float) GetResizeFilterWindowSupport(resizeFilter);
3717 status =SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_mem),(
void*)&imageBuffer);
3718 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&number_channels);
3719 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&columns);
3720 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&rows);
3721 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_mem),(
void*)&resizedImageBuffer);
3722 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&resizedColumns);
3723 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&resizedRows);
3724 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&yFactor);
3725 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
int),(
void*)&resizeFilterType);
3726 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
int),(
void*)&resizeWindowType);
3727 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_mem),(
void*)&resizeFilterCubicCoefficients);
3728 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterScale);
3729 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterSupport);
3730 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterWindowSupport);
3731 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterBlur);
3732 status|=SetOpenCLKernelArg(verticalKernel,i++,imageCacheLocalMemorySize, NULL);
3733 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
int), &numCachedPixels);
3734 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
unsigned int), &pixelPerWorkgroup);
3735 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
unsigned int), &chunkSize);
3736 status|=SetOpenCLKernelArg(verticalKernel,i++,pixelAccumulatorLocalMemorySize, NULL);
3737 status|=SetOpenCLKernelArg(verticalKernel,i++,weightAccumulatorLocalMemorySize, NULL);
3738 status|=SetOpenCLKernelArg(verticalKernel,i++,gammaAccumulatorLocalMemorySize, NULL);
3740 if (status != CL_SUCCESS)
3742 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3743 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
3747 gsize[0]=resizedColumns;
3748 gsize[1]=(resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*
3751 lsize[1]=workgroupSize;
3752 outputReady=EnqueueOpenCLKernel(queue,verticalKernel,2,(
const size_t *) NULL,
3753 gsize,lsize,image,filteredImage,MagickFalse,exception);
3757 if (verticalKernel != (cl_kernel) NULL)
3758 ReleaseOpenCLKernel(verticalKernel);
3760 return(outputReady);
3763static Image *ComputeResizeImage(
const Image* image,MagickCLEnv clEnv,
3764 const size_t resizedColumns,
const size_t resizedRows,
3771 cubicCoefficientsBuffer,
3772 filteredImageBuffer,
3780 *resizeFilterCoefficient;
3783 coefficientBuffer[7],
3804 filteredImageBuffer=NULL;
3805 tempImageBuffer=NULL;
3806 cubicCoefficientsBuffer=NULL;
3807 outputReady=MagickFalse;
3809 device=RequestOpenCLDevice(clEnv);
3810 queue=AcquireOpenCLCommandQueue(device);
3811 filteredImage=CloneImage(image,resizedColumns,resizedRows,MagickTrue,
3813 if (filteredImage == (
Image *) NULL)
3815 if (filteredImage->number_channels != image->number_channels)
3817 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
3818 if (imageBuffer == (cl_mem) NULL)
3820 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
3821 if (filteredImageBuffer == (cl_mem) NULL)
3824 resizeFilterCoefficient=GetResizeFilterCoefficient(resizeFilter);
3825 for (i = 0; i < 7; i++)
3826 coefficientBuffer[i]=(
float) resizeFilterCoefficient[i];
3827 cubicCoefficientsBuffer=CreateOpenCLBuffer(device,CL_MEM_COPY_HOST_PTR |
3828 CL_MEM_READ_ONLY,
sizeof(coefficientBuffer),&coefficientBuffer);
3829 if (cubicCoefficientsBuffer == (cl_mem) NULL)
3831 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3832 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
3836 number_channels=(cl_uint) image->number_channels;
3837 xFactor=(float) resizedColumns/(
float) image->columns;
3838 yFactor=(float) resizedRows/(
float) image->rows;
3839 if (xFactor > yFactor)
3841 length=resizedColumns*image->rows*number_channels;
3842 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
3843 sizeof(CLQuantum),(
void *) NULL);
3844 if (tempImageBuffer == (cl_mem) NULL)
3846 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3847 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
3851 outputReady=resizeHorizontalFilter(device,queue,image,filteredImage,
3852 imageBuffer,number_channels,(cl_uint) image->columns,
3853 (cl_uint) image->rows,tempImageBuffer,(cl_uint) resizedColumns,
3854 (cl_uint) image->rows,resizeFilter,cubicCoefficientsBuffer,xFactor,
3856 if (outputReady == MagickFalse)
3859 outputReady=resizeVerticalFilter(device,queue,image,filteredImage,
3860 tempImageBuffer,number_channels,(cl_uint) resizedColumns,
3861 (cl_uint) image->rows,filteredImageBuffer,(cl_uint) resizedColumns,
3862 (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,yFactor,
3864 if (outputReady == MagickFalse)
3869 length=image->columns*resizedRows*number_channels;
3870 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
3871 sizeof(CLQuantum),(
void *) NULL);
3872 if (tempImageBuffer == (cl_mem) NULL)
3874 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
3875 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
3879 outputReady=resizeVerticalFilter(device,queue,image,filteredImage,
3880 imageBuffer,number_channels,(cl_uint) image->columns,
3881 (cl_int) image->rows,tempImageBuffer,(cl_uint) image->columns,
3882 (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,yFactor,
3884 if (outputReady == MagickFalse)
3887 outputReady=resizeHorizontalFilter(device,queue,image,filteredImage,
3888 tempImageBuffer,number_channels,(cl_uint) image->columns,
3889 (cl_uint) resizedRows,filteredImageBuffer,(cl_uint) resizedColumns,
3890 (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,xFactor,
3892 if (outputReady == MagickFalse)
3898 if (imageBuffer != (cl_mem) NULL)
3899 ReleaseOpenCLMemObject(imageBuffer);
3900 if (filteredImageBuffer != (cl_mem) NULL)
3901 ReleaseOpenCLMemObject(filteredImageBuffer);
3902 if (tempImageBuffer != (cl_mem) NULL)
3903 ReleaseOpenCLMemObject(tempImageBuffer);
3904 if (cubicCoefficientsBuffer != (cl_mem) NULL)
3905 ReleaseOpenCLMemObject(cubicCoefficientsBuffer);
3906 if (queue != (cl_command_queue) NULL)
3907 ReleaseOpenCLCommandQueue(device,queue);
3908 if (device != (MagickCLDevice) NULL)
3909 ReleaseOpenCLDevice(device);
3910 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
3911 filteredImage=DestroyImage(filteredImage);
3913 return(filteredImage);
3916static MagickBooleanType gpuSupportedResizeWeighting(
3917 ResizeWeightingFunctionType f)
3924 if (supportedResizeWeighting[i] == LastWeightingFunction)
3926 if (supportedResizeWeighting[i] == f)
3929 return(MagickFalse);
3932MagickPrivate
Image *AccelerateResizeImage(
const Image *image,
3933 const size_t resizedColumns,
const size_t resizedRows,
3942 assert(image != NULL);
3945 if (checkAccelerateCondition(image) == MagickFalse)
3946 return((
Image *) NULL);
3948 if ((gpuSupportedResizeWeighting(GetResizeFilterWeightingType(
3949 resizeFilter)) == MagickFalse) ||
3950 (gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(
3951 resizeFilter)) == MagickFalse))
3952 return((
Image *) NULL);
3954 clEnv=getOpenCLEnvironment(exception);
3955 if (clEnv == (MagickCLEnv) NULL)
3956 return((
Image *) NULL);
3958 filteredImage=ComputeResizeImage(image,clEnv,resizedColumns,resizedRows,
3959 resizeFilter,exception);
3960 return(filteredImage);
3975static Image* ComputeRotationalBlurImage(
const Image *image,MagickCLEnv clEnv,
3985 channel_mask=get32BitChannelValue(image->channel_mask),
3990 filteredImageBuffer,
3995 rotationalBlurKernel;
4021 assert(image != (
Image *) NULL);
4022 assert(image->signature == MagickCoreSignature);
4023 if (IsEventLogging() != MagickFalse)
4024 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
4028 filteredImageBuffer=NULL;
4029 sinThetaBuffer=NULL;
4030 cosThetaBuffer=NULL;
4031 rotationalBlurKernel=NULL;
4032 outputReady=MagickFalse;
4034 device=RequestOpenCLDevice(clEnv);
4035 queue=AcquireOpenCLCommandQueue(device);
4036 filteredImage=cloneImage(image,exception);
4037 if (filteredImage == (
Image *) NULL)
4039 if (filteredImage->number_channels != image->number_channels)
4041 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4042 if (imageBuffer == (cl_mem) NULL)
4044 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4045 if (filteredImageBuffer == (cl_mem) NULL)
4048 blurCenter.x=(float) (image->columns-1)/2.0;
4049 blurCenter.y=(float) (image->rows-1)/2.0;
4050 blurRadius=hypot(blurCenter.x,blurCenter.y);
4051 cossin_theta_size=(
unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt(
4052 (
double) blurRadius)+2UL);
4054 cosThetaPtr=(
float *) AcquireQuantumMemory(cossin_theta_size,
sizeof(
float));
4055 if (cosThetaPtr == (
float *) NULL)
4057 sinThetaPtr=(
float *) AcquireQuantumMemory(cossin_theta_size,
sizeof(
float));
4058 if (sinThetaPtr == (
float *) NULL)
4060 cosThetaPtr=(
float *) RelinquishMagickMemory(cosThetaPtr);
4064 theta=DegreesToRadians(angle)/(double) (cossin_theta_size-1);
4065 offset=theta*(float) (cossin_theta_size-1)/2.0;
4066 for (i=0; i < (ssize_t) cossin_theta_size; i++)
4068 cosThetaPtr[i]=(float)cos((
double) (theta*i-offset));
4069 sinThetaPtr[i]=(float)sin((
double) (theta*i-offset));
4072 sinThetaBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
4073 CL_MEM_COPY_HOST_PTR,cossin_theta_size*
sizeof(
float),sinThetaPtr);
4074 sinThetaPtr=(
float *) RelinquishMagickMemory(sinThetaPtr);
4075 cosThetaBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
4076 CL_MEM_COPY_HOST_PTR,cossin_theta_size*
sizeof(
float),cosThetaPtr);
4077 cosThetaPtr=(
float *) RelinquishMagickMemory(cosThetaPtr);
4078 if ((sinThetaBuffer == (cl_mem) NULL) || (cosThetaBuffer == (cl_mem) NULL))
4080 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4081 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4085 rotationalBlurKernel=AcquireOpenCLKernel(device,
"RotationalBlur");
4086 if (rotationalBlurKernel == (cl_kernel) NULL)
4088 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4089 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4093 number_channels=(cl_uint) image->number_channels;
4096 status =SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4097 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_uint),&number_channels);
4098 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_int),&channel_mask);
4099 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_float2), &blurCenter);
4100 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&cosThetaBuffer);
4101 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&sinThetaBuffer);
4102 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_uint), &cossin_theta_size);
4103 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4104 if (status != CL_SUCCESS)
4106 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4107 ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
4111 gsize[0]=image->columns;
4112 gsize[1]=image->rows;
4113 outputReady=EnqueueOpenCLKernel(queue,rotationalBlurKernel,2,
4114 (
const size_t *) NULL,gsize,(
const size_t *) NULL,image,filteredImage,
4115 MagickFalse,exception);
4119 if (imageBuffer != (cl_mem) NULL)
4120 ReleaseOpenCLMemObject(imageBuffer);
4121 if (filteredImageBuffer != (cl_mem) NULL)
4122 ReleaseOpenCLMemObject(filteredImageBuffer);
4123 if (sinThetaBuffer != (cl_mem) NULL)
4124 ReleaseOpenCLMemObject(sinThetaBuffer);
4125 if (cosThetaBuffer != (cl_mem) NULL)
4126 ReleaseOpenCLMemObject(cosThetaBuffer);
4127 if (rotationalBlurKernel != (cl_kernel) NULL)
4128 ReleaseOpenCLKernel(rotationalBlurKernel);
4129 if (queue != (cl_command_queue) NULL)
4130 ReleaseOpenCLCommandQueue(device,queue);
4131 if (device != (MagickCLDevice) NULL)
4132 ReleaseOpenCLDevice(device);
4133 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4134 filteredImage=DestroyImage(filteredImage);
4136 return(filteredImage);
4139MagickPrivate
Image* AccelerateRotationalBlurImage(
const Image *image,
4148 assert(image != NULL);
4150 if (IsEventLogging() != MagickFalse)
4151 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
4153 if (checkAccelerateCondition(image) == MagickFalse)
4154 return((
Image *) NULL);
4156 clEnv=getOpenCLEnvironment(exception);
4157 if (clEnv == (MagickCLEnv) NULL)
4158 return((
Image *) NULL);
4160 filteredImage=ComputeRotationalBlurImage(image,clEnv,angle,exception);
4161 return filteredImage;
4176static Image *ComputeUnsharpMaskImage(
const Image *image,MagickCLEnv clEnv,
4177 const double radius,
const double sigma,
const double gain,
4184 channel_mask=get32BitChannelValue(image->channel_mask),
4189 unsharpMaskBlurColumnKernel;
4192 filteredImageBuffer,
4229 filteredImageBuffer=NULL;
4230 tempImageBuffer=NULL;
4231 imageKernelBuffer=NULL;
4233 unsharpMaskBlurColumnKernel=NULL;
4234 outputReady=MagickFalse;
4236 device=RequestOpenCLDevice(clEnv);
4237 queue=AcquireOpenCLCommandQueue(device);
4238 filteredImage=cloneImage(image,exception);
4239 if (filteredImage == (
Image *) NULL)
4241 if (filteredImage->number_channels != image->number_channels)
4243 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4244 if (imageBuffer == (cl_mem) NULL)
4246 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4247 if (filteredImageBuffer == (cl_mem) NULL)
4250 imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
4253 length=image->columns*image->rows;
4254 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
4255 sizeof(cl_float4),NULL);
4256 if (tempImageBuffer == (cl_mem) NULL)
4258 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4259 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4263 blurRowKernel=AcquireOpenCLKernel(device,
"BlurRow");
4264 if (blurRowKernel == (cl_kernel) NULL)
4266 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4267 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4271 unsharpMaskBlurColumnKernel=AcquireOpenCLKernel(device,
4272 "UnsharpMaskBlurColumn");
4273 if (unsharpMaskBlurColumnKernel == (cl_kernel) NULL)
4275 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4276 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4280 number_channels=(cl_uint) image->number_channels;
4281 imageColumns=(cl_uint) image->columns;
4282 imageRows=(cl_uint) image->rows;
4287 status =SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4288 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),&number_channels);
4289 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_int),&channel_mask);
4290 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4291 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
4292 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
4293 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
4294 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_float4)*(chunkSize+kernelWidth),(
void *) NULL);
4295 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
4296 if (status != CL_SUCCESS)
4298 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4299 ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
4303 gsize[0]=chunkSize*((image->columns+chunkSize-1)/chunkSize);
4304 gsize[1]=image->rows;
4307 outputReady=EnqueueOpenCLKernel(queue,blurRowKernel,2,
4308 (
const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
4313 fThreshold=(float) threshold;
4316 status =SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4317 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
4318 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),&number_channels);
4319 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_int),&channel_mask);
4320 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
4321 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
4322 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,(chunkSize+kernelWidth-1)*
sizeof(cl_float4),NULL);
4323 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,kernelWidth*
sizeof(
float),NULL);
4324 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4325 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
4326 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
float),(
void *)&fGain);
4327 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
float),(
void *)&fThreshold);
4328 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4329 if (status != CL_SUCCESS)
4331 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4332 ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
4336 gsize[0]=image->columns;
4337 gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize);
4340 outputReady=EnqueueOpenCLKernel(queue,unsharpMaskBlurColumnKernel,2,
4341 (
const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
4346 if (imageBuffer != (cl_mem) NULL)
4347 ReleaseOpenCLMemObject(imageBuffer);
4348 if (filteredImageBuffer != (cl_mem) NULL)
4349 ReleaseOpenCLMemObject(filteredImageBuffer);
4350 if (tempImageBuffer != (cl_mem) NULL)
4351 ReleaseOpenCLMemObject(tempImageBuffer);
4352 if (imageKernelBuffer != (cl_mem) NULL)
4353 ReleaseOpenCLMemObject(imageKernelBuffer);
4354 if (blurRowKernel != (cl_kernel) NULL)
4355 ReleaseOpenCLKernel(blurRowKernel);
4356 if (unsharpMaskBlurColumnKernel != (cl_kernel) NULL)
4357 ReleaseOpenCLKernel(unsharpMaskBlurColumnKernel);
4358 if (queue != (cl_command_queue) NULL)
4359 ReleaseOpenCLCommandQueue(device,queue);
4360 if (device != (MagickCLDevice) NULL)
4361 ReleaseOpenCLDevice(device);
4362 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4363 filteredImage=DestroyImage(filteredImage);
4365 return(filteredImage);
4368static Image *ComputeUnsharpMaskImageSingle(
const Image *image,
4369 MagickCLEnv clEnv,
const double radius,
const double sigma,
const double gain,
4376 channel_mask=get32BitChannelValue(image->channel_mask),
4383 filteredImageBuffer,
4413 filteredImageBuffer=NULL;
4414 imageKernelBuffer=NULL;
4415 unsharpMaskKernel=NULL;
4416 outputReady=MagickFalse;
4418 device=RequestOpenCLDevice(clEnv);
4419 queue=AcquireOpenCLCommandQueue(device);
4420 filteredImage=cloneImage(image,exception);
4421 if (filteredImage == (
Image *) NULL)
4423 if (filteredImage->number_channels != image->number_channels)
4425 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4426 if (imageBuffer == (cl_mem) NULL)
4428 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4429 if (filteredImageBuffer == (cl_mem) NULL)
4432 imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
4435 unsharpMaskKernel=AcquireOpenCLKernel(device,
"UnsharpMask");
4436 if (unsharpMaskKernel == NULL)
4438 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4439 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4443 imageColumns=(cl_uint) image->columns;
4444 imageRows=(cl_uint) image->rows;
4445 number_channels=(cl_uint) image->number_channels;
4447 fThreshold=(float) threshold;
4450 status =SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4451 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
4452 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_int),&channel_mask);
4453 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4454 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
4455 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
4456 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
4457 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_float4)*(8 * (32 + kernelWidth)),(
void *) NULL);
4458 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(
float),(
void *)&fGain);
4459 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(
float),(
void *)&fThreshold);
4460 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4461 if (status != CL_SUCCESS)
4463 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4464 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
4468 gsize[0]=((image->columns + 7) / 8)*8;
4469 gsize[1]=((image->rows + 31) / 32)*32;
4472 outputReady=EnqueueOpenCLKernel(queue,unsharpMaskKernel,2,(
const size_t *) NULL,
4473 gsize,lsize,image,filteredImage,MagickFalse,exception);
4477 if (imageBuffer != (cl_mem) NULL)
4478 ReleaseOpenCLMemObject(imageBuffer);
4479 if (filteredImageBuffer != (cl_mem) NULL)
4480 ReleaseOpenCLMemObject(filteredImageBuffer);
4481 if (imageKernelBuffer != (cl_mem) NULL)
4482 ReleaseOpenCLMemObject(imageKernelBuffer);
4483 if (unsharpMaskKernel != (cl_kernel) NULL)
4484 ReleaseOpenCLKernel(unsharpMaskKernel);
4485 if (queue != (cl_command_queue) NULL)
4486 ReleaseOpenCLCommandQueue(device,queue);
4487 if (device != (MagickCLDevice) NULL)
4488 ReleaseOpenCLDevice(device);
4489 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4490 filteredImage=DestroyImage(filteredImage);
4492 return(filteredImage);
4495MagickPrivate
Image *AccelerateUnsharpMaskImage(
const Image *image,
4496 const double radius,
const double sigma,
const double gain,
4505 assert(image != NULL);
4508 if (checkAccelerateCondition(image) == MagickFalse)
4509 return((
Image *) NULL);
4511 clEnv=getOpenCLEnvironment(exception);
4512 if (clEnv == (MagickCLEnv) NULL)
4513 return((
Image *) NULL);
4516 filteredImage=ComputeUnsharpMaskImageSingle(image,clEnv,radius,sigma,gain,
4517 threshold,exception);
4519 filteredImage=ComputeUnsharpMaskImage(image,clEnv,radius,sigma,gain,
4520 threshold,exception);
4521 return(filteredImage);
4524static Image *ComputeWaveletDenoiseImage(
const Image *image,MagickCLEnv clEnv,
4536 SIZE=TILESIZE-2*PAD;
4548 filteredImageBuffer,
4576 filteredImageBuffer=NULL;
4579 outputReady=MagickFalse;
4581 device=RequestOpenCLDevice(clEnv);
4583 if (strcmp(
"Intel(R) HD Graphics",device->name) == 0)
4585 queue=AcquireOpenCLCommandQueue(device);
4586 filteredImage=CloneImage(image,0,0,MagickTrue,
4588 if (filteredImage == (
Image *) NULL)
4590 if (filteredImage->number_channels != image->number_channels)
4592 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4593 if (imageBuffer == (cl_mem) NULL)
4595 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4596 if (filteredImageBuffer == (cl_mem) NULL)
4599 denoiseKernel=AcquireOpenCLKernel(device,
"WaveletDenoise");
4600 if (denoiseKernel == (cl_kernel) NULL)
4602 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4603 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4607 number_channels=(cl_uint)image->number_channels;
4608 width=(cl_uint)image->columns;
4609 height=(cl_uint)image->rows;
4610 max_channels=number_channels;
4611 if ((max_channels == 4) || (max_channels == 2))
4612 max_channels=max_channels-1;
4614 passes=(((1.0f*image->columns)*image->rows)+1999999.0f)/2000000.0f;
4615 passes=(passes < 1) ? 1 : passes;
4618 status =SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4619 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4620 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
4621 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&max_channels);
4622 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_float),(
void *)&thresh);
4623 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_int),(
void *)&PASSES);
4624 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&width);
4625 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&height);
4626 if (status != CL_SUCCESS)
4628 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4629 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
4633 for (x = 0; x < passes; ++x)
4635 gsize[0]=((width+(SIZE-1))/SIZE)*TILESIZE;
4636 gsize[1]=((((height+(SIZE-1))/SIZE)+passes-1)/passes)*4;
4640 goffset[1]=x*gsize[1];
4642 outputReady=EnqueueOpenCLKernel(queue,denoiseKernel,2,goffset,gsize,lsize,
4643 image,filteredImage,MagickTrue,exception);
4644 if (outputReady == MagickFalse)
4650 if (imageBuffer != (cl_mem) NULL)
4651 ReleaseOpenCLMemObject(imageBuffer);
4652 if (filteredImageBuffer != (cl_mem) NULL)
4653 ReleaseOpenCLMemObject(filteredImageBuffer);
4654 if (denoiseKernel != (cl_kernel) NULL)
4655 ReleaseOpenCLKernel(denoiseKernel);
4656 if (queue != (cl_command_queue) NULL)
4657 ReleaseOpenCLCommandQueue(device,queue);
4658 if (device != (MagickCLDevice) NULL)
4659 ReleaseOpenCLDevice(device);
4660 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4661 filteredImage=DestroyImage(filteredImage);
4663 return(filteredImage);
4666MagickPrivate
Image *AccelerateWaveletDenoiseImage(
const Image *image,
4675 assert(image != NULL);
4678 if (checkAccelerateCondition(image) == MagickFalse)
4679 return((
Image *) NULL);
4681 clEnv=getOpenCLEnvironment(exception);
4682 if (clEnv == (MagickCLEnv) NULL)
4683 return((
Image *) NULL);
4685 filteredImage=ComputeWaveletDenoiseImage(image,clEnv,threshold,exception);
4687 return(filteredImage);