44 #include "MagickCore/studio.h"
45 #include "MagickCore/accelerate-private.h"
46 #include "MagickCore/accelerate-kernels-private.h"
47 #include "MagickCore/artifact.h"
48 #include "MagickCore/cache.h"
49 #include "MagickCore/cache-private.h"
50 #include "MagickCore/cache-view.h"
51 #include "MagickCore/color-private.h"
52 #include "MagickCore/delegate-private.h"
53 #include "MagickCore/enhance.h"
54 #include "MagickCore/exception.h"
55 #include "MagickCore/exception-private.h"
56 #include "MagickCore/gem.h"
57 #include "MagickCore/image.h"
58 #include "MagickCore/image-private.h"
59 #include "MagickCore/linked-list.h"
60 #include "MagickCore/list.h"
61 #include "MagickCore/memory_.h"
62 #include "MagickCore/monitor-private.h"
63 #include "MagickCore/opencl.h"
64 #include "MagickCore/opencl-private.h"
65 #include "MagickCore/option.h"
66 #include "MagickCore/pixel-accessor.h"
67 #include "MagickCore/prepress.h"
68 #include "MagickCore/quantize.h"
69 #include "MagickCore/quantum-private.h"
70 #include "MagickCore/random_.h"
71 #include "MagickCore/random-private.h"
72 #include "MagickCore/registry.h"
73 #include "MagickCore/resize.h"
74 #include "MagickCore/resize-private.h"
75 #include "MagickCore/semaphore.h"
76 #include "MagickCore/splay-tree.h"
77 #include "MagickCore/statistic.h"
78 #include "MagickCore/string_.h"
79 #include "MagickCore/string-private.h"
80 #include "MagickCore/token.h"
82 #define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
83 #define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
85 #if defined(MAGICKCORE_OPENCL_SUPPORT)
90 #define ALIGNED(pointer,type) ((((size_t)(pointer)) & (sizeof(type)-1)) == 0)
95 static const ResizeWeightingFunctionType supportedResizeWeighting[] =
98 TriangleWeightingFunction,
99 HannWeightingFunction,
100 HammingWeightingFunction,
101 BlackmanWeightingFunction,
102 CubicBCWeightingFunction,
103 SincWeightingFunction,
104 SincFastWeightingFunction,
105 LastWeightingFunction
111 static MagickBooleanType checkAccelerateCondition(
const Image* image)
114 if (image->storage_class != DirectClass)
118 if (image->colorspace != RGBColorspace &&
119 image->colorspace != sRGBColorspace &&
120 image->colorspace != LinearGRAYColorspace &&
121 image->colorspace != GRAYColorspace)
125 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
126 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
130 if (((image->channels & ReadMaskChannel) != 0) ||
131 ((image->channels & WriteMaskChannel) != 0) ||
132 ((image->channels & CompositeMaskChannel) != 0))
135 if (image->number_channels > 4)
139 if (GetPixelChannelOffset(image,RedPixelChannel) != 0)
142 if (image->number_channels == 1)
146 if ((image->number_channels == 2) &&
147 (GetPixelChannelOffset(image,AlphaPixelChannel) == 1))
150 if (image->number_channels == 2)
154 if ((GetPixelChannelOffset(image,GreenPixelChannel) != 1) ||
155 (GetPixelChannelOffset(image,BluePixelChannel) != 2))
158 if (image->number_channels == 3)
162 if (GetPixelChannelOffset(image,AlphaPixelChannel) != 3)
168 static MagickBooleanType checkAccelerateConditionRGBA(
const Image* image)
170 if (checkAccelerateCondition(image) == MagickFalse)
174 if (image->number_channels != 4)
177 if ((GetPixelRedTraits(image) == UndefinedPixelTrait) ||
178 (GetPixelGreenTraits(image) == UndefinedPixelTrait) ||
179 (GetPixelBlueTraits(image) == UndefinedPixelTrait) ||
180 (GetPixelAlphaTraits(image) == UndefinedPixelTrait))
186 static MagickBooleanType checkPixelIntensity(
const Image *image,
187 const PixelIntensityMethod method)
190 if ((method == Rec601LumaPixelIntensityMethod) ||
191 (method == Rec709LumaPixelIntensityMethod))
193 if (image->colorspace == RGBColorspace)
197 if ((method == Rec601LuminancePixelIntensityMethod) ||
198 (method == Rec709LuminancePixelIntensityMethod))
200 if (image->colorspace == sRGBColorspace)
207 static MagickBooleanType checkHistogramCondition(
const Image *image,
208 const PixelIntensityMethod method)
211 if ((image->channel_mask & SyncChannels) == 0)
214 return(checkPixelIntensity(image,method));
217 static MagickCLEnv getOpenCLEnvironment(
ExceptionInfo* exception)
222 clEnv=GetCurrentOpenCLEnv();
223 if (clEnv == (MagickCLEnv) NULL)
224 return((MagickCLEnv) NULL);
226 if (clEnv->enabled == MagickFalse)
227 return((MagickCLEnv) NULL);
229 if (InitializeOpenCL(clEnv,exception) == MagickFalse)
230 return((MagickCLEnv) NULL);
240 if (((image->channel_mask & RedChannel) != 0) &&
241 ((image->channel_mask & GreenChannel) != 0) &&
242 ((image->channel_mask & BlueChannel) != 0) &&
243 ((image->channel_mask & AlphaChannel) != 0))
244 clone=CloneImage(image,0,0,MagickTrue,exception);
247 clone=CloneImage(image,0,0,MagickTrue,exception);
248 if (clone != (
Image *) NULL)
249 SyncImagePixelCache(clone,exception);
256 inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
257 const unsigned int orgGlobalSize,
const unsigned int localGroupSize)
259 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
262 static cl_mem createKernelInfo(MagickCLDevice device,
const double radius,
266 geometry[MagickPathExtent];
280 (void) FormatLocaleString(geometry,MagickPathExtent,
281 "blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
282 kernel=AcquireKernelInfo(geometry,exception);
285 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
286 ResourceLimitWarning,
"AcquireKernelInfo failed.",
".");
287 return((cl_mem) NULL);
289 kernelBufferPtr=(
float *) AcquireMagickMemory(kernel->width*
290 sizeof(*kernelBufferPtr));
291 if (kernelBufferPtr == (
float *) NULL)
293 kernel=DestroyKernelInfo(kernel);
294 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
295 ResourceLimitWarning,
"MemoryAllocationFailed.",
".");
296 return((cl_mem) NULL);
298 for (i = 0; i < (ssize_t) kernel->width; i++)
299 kernelBufferPtr[i]=(float) kernel->values[i];
300 imageKernelBuffer=CreateOpenCLBuffer(device,CL_MEM_COPY_HOST_PTR |
301 CL_MEM_READ_ONLY,kernel->width*
sizeof(*kernelBufferPtr),kernelBufferPtr);
302 *width=(cl_uint) kernel->width;
303 kernelBufferPtr=(
float *) RelinquishMagickMemory(kernelBufferPtr);
304 kernel=DestroyKernelInfo(kernel);
305 if (imageKernelBuffer == (cl_mem) NULL)
306 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
307 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
308 return(imageKernelBuffer);
311 static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
312 MagickCLDevice device,cl_command_queue queue,cl_mem imageBuffer,
313 cl_mem histogramBuffer,
Image *image,
const ChannelType channel,
338 histogramKernel = NULL;
340 outputReady = MagickFalse;
341 colorspace = image->colorspace;
342 method = image->intensity;
345 histogramKernel = AcquireOpenCLKernel(device,
"Histogram");
346 if (histogramKernel == NULL)
348 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
354 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
355 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(ChannelType),&channel);
356 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_uint),&colorspace);
357 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_uint),&method);
358 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_mem),(
void *)&histogramBuffer);
359 if (clStatus != CL_SUCCESS)
361 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
366 global_work_size[0] = image->columns;
367 global_work_size[1] = image->rows;
369 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
371 if (clStatus != CL_SUCCESS)
373 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
376 RecordProfileData(device,histogramKernel,event);
378 outputReady = MagickTrue;
382 if (histogramKernel!=NULL)
383 ReleaseOpenCLKernel(histogramKernel);
400 static Image *ComputeAddNoiseImage(
const Image *image,MagickCLEnv clEnv,
401 const NoiseType noise_type,
const double attenuate,
ExceptionInfo *exception)
423 numRandomNumberPerPixel,
452 filteredImageBuffer=NULL;
454 outputReady=MagickFalse;
456 device=RequestOpenCLDevice(clEnv);
457 queue=AcquireOpenCLCommandQueue(device);
458 if (queue == (cl_command_queue) NULL)
460 filteredImage=cloneImage(image,exception);
461 if (filteredImage == (
Image *) NULL)
463 if (filteredImage->number_channels != image->number_channels)
465 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
466 if (imageBuffer == (cl_mem) NULL)
468 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
469 if (filteredImageBuffer == (cl_mem) NULL)
474 numRandomNumberPerPixel=0;
485 case MultiplicativeGaussianNoise:
490 if (GetPixelRedTraits(image) != UndefinedPixelTrait)
491 numRandomNumberPerPixel+=(cl_uint) numRandPerChannel;
492 if (GetPixelGreenTraits(image) != UndefinedPixelTrait)
493 numRandomNumberPerPixel+=(cl_uint) numRandPerChannel;
494 if (GetPixelBlueTraits(image) != UndefinedPixelTrait)
495 numRandomNumberPerPixel+=(cl_uint) numRandPerChannel;
496 if (GetPixelAlphaTraits(image) != UndefinedPixelTrait)
497 numRandomNumberPerPixel+=(cl_uint) numRandPerChannel;
499 addNoiseKernel=AcquireOpenCLKernel(device,
"AddNoise");
500 if (addNoiseKernel == (cl_kernel) NULL)
502 (void)OpenCLThrowMagickException(device,exception,GetMagickModule(),
503 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
508 workItemCount=device->max_compute_units*2*256;
509 inputPixelCount=(cl_int) (image->columns*image->rows);
510 pixelsPerWorkitem=(inputPixelCount+workItemCount-1)/workItemCount;
511 pixelsPerWorkitem=((pixelsPerWorkitem+3)/4)*4;
513 gsize[0]=workItemCount;
515 randomInfo=AcquireRandomInfo();
516 s=GetRandomInfoSeed(randomInfo);
518 (void) GetPseudoRandomValue(randomInfo);
520 randomInfo=DestroyRandomInfo(randomInfo);
522 number_channels=(cl_uint) image->number_channels;
523 bufferLength=(cl_uint) (image->columns*image->rows*image->number_channels);
524 cl_attenuate=(cl_float) attenuate;
527 status =SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
528 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
529 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(ChannelType),(
void *)&image->channel_mask);
530 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_uint),(
void *)&bufferLength);
531 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_uint),(
void *)&pixelsPerWorkitem);
532 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(NoiseType),(
void *)&noise_type);
533 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_float),(
void *)&cl_attenuate);
534 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_uint),(
void *)&seed0);
535 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_uint),(
void *)&seed1);
536 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_uint),(
void *)&numRandomNumberPerPixel);
537 status|=SetOpenCLKernelArg(addNoiseKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
538 if (status != CL_SUCCESS)
540 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
541 ResourceLimitWarning,
"clSetKernelArg failed.",
".");
545 outputReady=EnqueueOpenCLKernel(queue,addNoiseKernel,1,(
const size_t *) NULL,gsize,
546 lsize,image,filteredImage,MagickFalse,exception);
550 if (imageBuffer != (cl_mem) NULL)
551 ReleaseOpenCLMemObject(imageBuffer);
552 if (filteredImageBuffer != (cl_mem) NULL)
553 ReleaseOpenCLMemObject(filteredImageBuffer);
554 if (addNoiseKernel != (cl_kernel) NULL)
555 ReleaseOpenCLKernel(addNoiseKernel);
556 if (queue != (cl_command_queue) NULL)
557 ReleaseOpenCLCommandQueue(device,queue);
558 if (device != (MagickCLDevice) NULL)
559 ReleaseOpenCLDevice(device);
560 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
561 filteredImage=DestroyImage(filteredImage);
563 return(filteredImage);
566 MagickPrivate
Image *AccelerateAddNoiseImage(
const Image *image,
567 const NoiseType noise_type,
const double attenuate,
ExceptionInfo *exception)
591 magick_unreferenced(image);
592 magick_unreferenced(noise_type);
593 magick_unreferenced(attenuate);
594 magick_unreferenced(exception);
595 return((
Image *)NULL);
610 static Image *ComputeBlurImage(
const Image* image,MagickCLEnv clEnv,
611 const double radius,
const double sigma,
ExceptionInfo *exception)
655 filteredImageBuffer=NULL;
656 tempImageBuffer=NULL;
657 imageKernelBuffer=NULL;
659 blurColumnKernel=NULL;
660 outputReady=MagickFalse;
662 assert(image != (
Image *) NULL);
663 assert(image->signature == MagickCoreSignature);
664 if (IsEventLogging() != MagickFalse)
665 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
667 device=RequestOpenCLDevice(clEnv);
668 queue=AcquireOpenCLCommandQueue(device);
669 filteredImage=cloneImage(image,exception);
670 if (filteredImage == (
Image *) NULL)
672 if (filteredImage->number_channels != image->number_channels)
674 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
675 if (imageBuffer == (cl_mem) NULL)
677 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
678 if (filteredImageBuffer == (cl_mem) NULL)
681 imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
683 if (imageKernelBuffer == (cl_mem) NULL)
686 length=image->columns*image->rows;
687 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
688 sizeof(cl_float4),(
void *) NULL);
689 if (tempImageBuffer == (cl_mem) NULL)
692 blurRowKernel=AcquireOpenCLKernel(device,
"BlurRow");
693 if (blurRowKernel == (cl_kernel) NULL)
695 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
696 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
700 number_channels=(cl_uint) image->number_channels;
701 imageColumns=(cl_uint) image->columns;
702 imageRows=(cl_uint) image->rows;
705 status =SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
706 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),&number_channels);
707 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(ChannelType),&image->channel_mask);
708 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
709 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
710 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
711 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
712 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_float4)*(chunkSize+kernelWidth),(
void *) NULL);
713 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
714 if (status != CL_SUCCESS)
716 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
717 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
721 gsize[0]=chunkSize*((image->columns+chunkSize-1)/chunkSize);
722 gsize[1]=image->rows;
726 outputReady=EnqueueOpenCLKernel(queue,blurRowKernel,2,(
size_t *) NULL,gsize,
727 lsize,image,filteredImage,MagickFalse,exception);
728 if (outputReady == MagickFalse)
731 blurColumnKernel=AcquireOpenCLKernel(device,
"BlurColumn");
732 if (blurColumnKernel == (cl_kernel) NULL)
734 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
735 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
740 status =SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
741 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),&number_channels);
742 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(ChannelType),&image->channel_mask);
743 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
744 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
745 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
746 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
747 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_float4)*(chunkSize+kernelWidth),(
void *) NULL);
748 status|=SetOpenCLKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
749 if (status != CL_SUCCESS)
751 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
752 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
756 gsize[0]=image->columns;
757 gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize);
761 outputReady=EnqueueOpenCLKernel(queue,blurColumnKernel,2,(
size_t *) NULL,gsize,
762 lsize,image,filteredImage,MagickFalse,exception);
766 if (imageBuffer != (cl_mem) NULL)
767 ReleaseOpenCLMemObject(imageBuffer);
768 if (filteredImageBuffer != (cl_mem) NULL)
769 ReleaseOpenCLMemObject(filteredImageBuffer);
770 if (tempImageBuffer != (cl_mem) NULL)
771 ReleaseOpenCLMemObject(tempImageBuffer);
772 if (imageKernelBuffer != (cl_mem) NULL)
773 ReleaseOpenCLMemObject(imageKernelBuffer);
774 if (blurRowKernel != (cl_kernel) NULL)
775 ReleaseOpenCLKernel(blurRowKernel);
776 if (blurColumnKernel != (cl_kernel) NULL)
777 ReleaseOpenCLKernel(blurColumnKernel);
778 if (queue != (cl_command_queue) NULL)
779 ReleaseOpenCLCommandQueue(device,queue);
780 if (device != (MagickCLDevice) NULL)
781 ReleaseOpenCLDevice(device);
782 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
783 filteredImage=DestroyImage(filteredImage);
785 return(filteredImage);
788 MagickPrivate
Image* AccelerateBlurImage(
const Image *image,
789 const double radius,
const double sigma,
ExceptionInfo *exception)
797 assert(image != NULL);
799 if (IsEventLogging() != MagickFalse)
800 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
802 if (checkAccelerateCondition(image) == MagickFalse)
803 return((
Image *) NULL);
805 clEnv=getOpenCLEnvironment(exception);
806 if (clEnv == (MagickCLEnv) NULL)
807 return((
Image *) NULL);
809 filteredImage=ComputeBlurImage(image,clEnv,radius,sigma,exception);
810 return(filteredImage);
825 static MagickBooleanType ComputeContrastImage(
Image *image,MagickCLEnv clEnv,
854 assert(image != (
Image *) NULL);
855 assert(image->signature == MagickCoreSignature);
856 if (IsEventLogging() != MagickFalse)
857 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
861 outputReady=MagickFalse;
863 device=RequestOpenCLDevice(clEnv);
864 queue=AcquireOpenCLCommandQueue(device);
865 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
866 if (imageBuffer == (cl_mem) NULL)
869 contrastKernel=AcquireOpenCLKernel(device,
"Contrast");
870 if (contrastKernel == (cl_kernel) NULL)
872 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
873 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
877 number_channels=(cl_uint) image->number_channels;
878 sign=sharpen != MagickFalse ? 1 : -1;
881 status =SetOpenCLKernelArg(contrastKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
882 status|=SetOpenCLKernelArg(contrastKernel,i++,
sizeof(cl_uint),&number_channels);
883 status|=SetOpenCLKernelArg(contrastKernel,i++,
sizeof(cl_int),&sign);
884 if (status != CL_SUCCESS)
886 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
887 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
891 gsize[0]=image->columns;
892 gsize[1]=image->rows;
894 outputReady=EnqueueOpenCLKernel(queue,contrastKernel,2,(
const size_t *) NULL,
895 gsize,(
const size_t *) NULL,image,(
Image *) NULL,MagickFalse,exception);
899 if (imageBuffer != (cl_mem) NULL)
900 ReleaseOpenCLMemObject(imageBuffer);
901 if (contrastKernel != (cl_kernel) NULL)
902 ReleaseOpenCLKernel(contrastKernel);
903 if (queue != (cl_command_queue) NULL)
904 ReleaseOpenCLCommandQueue(device,queue);
905 if (device != (MagickCLDevice) NULL)
906 ReleaseOpenCLDevice(device);
911 MagickPrivate MagickBooleanType AccelerateContrastImage(
Image *image,
920 assert(image != NULL);
922 if (IsEventLogging() != MagickFalse)
923 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
925 if (checkAccelerateCondition(image) == MagickFalse)
928 clEnv=getOpenCLEnvironment(exception);
929 if (clEnv == (MagickCLEnv) NULL)
932 status=ComputeContrastImage(image,clEnv,sharpen,exception);
948 static MagickBooleanType ComputeContrastStretchImage(
Image *image,
949 MagickCLEnv clEnv,
const double black_point,
const double white_point,
952 #define ContrastStretchImageTag "ContrastStretch/Image"
953 #define MaxRange(color) ((cl_float) ScaleQuantumToMap((Quantum) (color)))
1006 global_work_size[2];
1012 assert(image != (
Image *) NULL);
1013 assert(image->signature == MagickCoreSignature);
1014 if (IsEventLogging() != MagickFalse)
1015 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
1021 histogramBuffer = NULL;
1022 stretchMapBuffer = NULL;
1023 histogramKernel = NULL;
1024 stretchKernel = NULL;
1026 outputReady = MagickFalse;
1033 device = RequestOpenCLDevice(clEnv);
1034 queue = AcquireOpenCLCommandQueue(device);
1039 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*histogram));
1041 if (histogram == (cl_uint4 *) NULL)
1042 ThrowBinaryException(ResourceLimitError,
"MemoryAllocationFailed", image->filename);
1045 (void) memset(histogram,0,(MaxMap+1)*
sizeof(*histogram));
1061 image_view=AcquireAuthenticCacheView(image,exception);
1062 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
1064 if (inputPixels == (
void *) NULL)
1066 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
1072 if (ALIGNED(inputPixels,CLPixelPacket))
1074 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1078 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1081 length = image->columns * image->rows;
1082 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1083 if (clStatus != CL_SUCCESS)
1085 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1092 if (ALIGNED(histogram,cl_uint4))
1094 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
1095 hostPtr = histogram;
1099 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
1100 hostPtr = histogram;
1103 length = (MaxMap+1);
1104 histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(cl_uint4), hostPtr, &clStatus);
1105 if (clStatus != CL_SUCCESS)
1107 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1111 status = LaunchHistogramKernel(clEnv, device, queue, imageBuffer, histogramBuffer, image, image->channel_mask,exception);
1112 if (status == MagickFalse)
1116 if (ALIGNED(histogram,cl_uint4))
1118 length = (MaxMap+1);
1119 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
1123 length = (MaxMap+1);
1124 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length *
sizeof(cl_uint4), histogram, 0, NULL, NULL);
1126 if (clStatus != CL_SUCCESS)
1128 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
1133 if (ALIGNED(histogram,cl_uint4))
1135 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
1136 if (clStatus != CL_SUCCESS)
1138 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
".");
1144 #ifdef RECREATEBUFFER
1145 if (imageBuffer!=NULL)
1146 clEnv->library->clReleaseMemObject(imageBuffer);
1154 white.x=MaxRange(QuantumRange);
1155 if ((image->channel_mask & RedChannel) != 0)
1158 for (i=0; i <= (ssize_t) MaxMap; i++)
1160 intensity+=histogram[i].s[2];
1161 if (intensity > black_point)
1164 black.x=(cl_float) i;
1166 for (i=(ssize_t) MaxMap; i != 0; i--)
1168 intensity+=histogram[i].s[2];
1169 if (intensity > ((
double) image->columns*image->rows-white_point))
1172 white.x=(cl_float) i;
1175 white.y=MaxRange(QuantumRange);
1176 if ((image->channel_mask & GreenChannel) != 0)
1179 for (i=0; i <= (ssize_t) MaxMap; i++)
1181 intensity+=histogram[i].s[2];
1182 if (intensity > black_point)
1185 black.y=(cl_float) i;
1187 for (i=(ssize_t) MaxMap; i != 0; i--)
1189 intensity+=histogram[i].s[2];
1190 if (intensity > ((
double) image->columns*image->rows-white_point))
1193 white.y=(cl_float) i;
1196 white.z=MaxRange(QuantumRange);
1197 if ((image->channel_mask & BlueChannel) != 0)
1200 for (i=0; i <= (ssize_t) MaxMap; i++)
1202 intensity+=histogram[i].s[2];
1203 if (intensity > black_point)
1206 black.z=(cl_float) i;
1208 for (i=(ssize_t) MaxMap; i != 0; i--)
1210 intensity+=histogram[i].s[2];
1211 if (intensity > ((
double) image->columns*image->rows-white_point))
1214 white.z=(cl_float) i;
1217 white.w=MaxRange(QuantumRange);
1218 if ((image->channel_mask & AlphaChannel) != 0)
1221 for (i=0; i <= (ssize_t) MaxMap; i++)
1223 intensity+=histogram[i].s[2];
1224 if (intensity > black_point)
1227 black.w=(cl_float) i;
1229 for (i=(ssize_t) MaxMap; i != 0; i--)
1231 intensity+=histogram[i].s[2];
1232 if (intensity > ((
double) image->columns*image->rows-white_point))
1235 white.w=(cl_float) i;
1238 stretch_map=(
PixelPacket *) AcquireQuantumMemory(MaxMap+1UL,
1239 sizeof(*stretch_map));
1242 ThrowBinaryException(ResourceLimitError,
"MemoryAllocationFailed",
1248 (void) memset(stretch_map,0,(MaxMap+1)*
sizeof(*stretch_map));
1249 for (i=0; i <= (ssize_t) MaxMap; i++)
1251 if ((image->channel_mask & RedChannel) != 0)
1253 if (i < (ssize_t) black.x)
1254 stretch_map[i].red=(Quantum) 0;
1256 if (i > (ssize_t) white.x)
1257 stretch_map[i].red=QuantumRange;
1259 if (black.x != white.x)
1260 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap*
1261 (i-black.x)/(white.x-black.x)));
1263 if ((image->channel_mask & GreenChannel) != 0)
1265 if (i < (ssize_t) black.y)
1266 stretch_map[i].green=0;
1268 if (i > (ssize_t) white.y)
1269 stretch_map[i].green=QuantumRange;
1271 if (black.y != white.y)
1272 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap*
1273 (i-black.y)/(white.y-black.y)));
1275 if ((image->channel_mask & BlueChannel) != 0)
1277 if (i < (ssize_t) black.z)
1278 stretch_map[i].blue=0;
1280 if (i > (ssize_t) white.z)
1281 stretch_map[i].blue= QuantumRange;
1283 if (black.z != white.z)
1284 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap*
1285 (i-black.z)/(white.z-black.z)));
1287 if ((image->channel_mask & AlphaChannel) != 0)
1289 if (i < (ssize_t) black.w)
1290 stretch_map[i].alpha=0;
1292 if (i > (ssize_t) white.w)
1293 stretch_map[i].alpha=QuantumRange;
1295 if (black.w != white.w)
1296 stretch_map[i].alpha=ScaleMapToQuantum((MagickRealType) (MaxMap*
1297 (i-black.w)/(white.w-black.w)));
1304 if (((image->channel_mask & AlphaChannel) != 0) || (((image->channel_mask & IndexChannel) != 0) &&
1305 (image->colorspace == CMYKColorspace)))
1306 image->storage_class=DirectClass;
1307 if (image->storage_class == PseudoClass)
1312 for (i=0; i < (ssize_t) image->colors; i++)
1314 if ((image->channel_mask & RedChannel) != 0)
1316 if (black.x != white.x)
1317 image->colormap[i].red=stretch_map[
1318 ScaleQuantumToMap(image->colormap[i].red)].red;
1320 if ((image->channel_mask & GreenChannel) != 0)
1322 if (black.y != white.y)
1323 image->colormap[i].green=stretch_map[
1324 ScaleQuantumToMap(image->colormap[i].green)].green;
1326 if ((image->channel_mask & BlueChannel) != 0)
1328 if (black.z != white.z)
1329 image->colormap[i].blue=stretch_map[
1330 ScaleQuantumToMap(image->colormap[i].blue)].blue;
1332 if ((image->channel_mask & AlphaChannel) != 0)
1334 if (black.w != white.w)
1335 image->colormap[i].alpha=stretch_map[
1336 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
1351 #ifdef RECREATEBUFFER
1355 if (ALIGNED(inputPixels,CLPixelPacket))
1357 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
1361 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
1364 length = image->columns * image->rows;
1365 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1366 if (clStatus != CL_SUCCESS)
1368 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1376 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1377 hostPtr = stretch_map;
1381 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
1382 hostPtr = stretch_map;
1385 length = (MaxMap+1);
1386 stretchMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(
PixelPacket), hostPtr, &clStatus);
1387 if (clStatus != CL_SUCCESS)
1389 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1394 stretchKernel = AcquireOpenCLKernel(device,
"ContrastStretch");
1395 if (stretchKernel == NULL)
1397 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
1403 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1404 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(ChannelType),&image->channel_mask);
1405 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_mem),(
void *)&stretchMapBuffer);
1406 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_float4),&white);
1407 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_float4),&black);
1408 if (clStatus != CL_SUCCESS)
1410 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1415 global_work_size[0] = image->columns;
1416 global_work_size[1] = image->rows;
1418 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
1420 if (clStatus != CL_SUCCESS)
1422 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1425 RecordProfileData(device,stretchKernel,event);
1428 if (ALIGNED(inputPixels,CLPixelPacket))
1430 length = image->columns * image->rows;
1431 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1435 length = image->columns * image->rows;
1436 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
1438 if (clStatus != CL_SUCCESS)
1440 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
1444 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
1448 image_view=DestroyCacheView(image_view);
1450 if (imageBuffer!=NULL)
1451 clEnv->library->clReleaseMemObject(imageBuffer);
1453 if (stretchMapBuffer!=NULL)
1454 clEnv->library->clReleaseMemObject(stretchMapBuffer);
1455 if (stretch_map!=NULL)
1456 stretch_map=(
PixelPacket *) RelinquishMagickMemory(stretch_map);
1457 if (histogramBuffer!=NULL)
1458 clEnv->library->clReleaseMemObject(histogramBuffer);
1459 if (histogram!=NULL)
1460 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
1461 if (histogramKernel!=NULL)
1462 ReleaseOpenCLKernel(histogramKernel);
1463 if (stretchKernel!=NULL)
1464 ReleaseOpenCLKernel(stretchKernel);
1466 ReleaseOpenCLCommandQueue(device,queue);
1468 ReleaseOpenCLDevice(device);
1470 return(outputReady);
1473 MagickPrivate MagickBooleanType AccelerateContrastStretchImage(
1474 Image *image,
const double black_point,
const double white_point,
1483 assert(image != NULL);
1485 if (IsEventLogging() != MagickFalse)
1486 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
1488 if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
1489 (checkHistogramCondition(image,image->intensity) == MagickFalse))
1490 return(MagickFalse);
1492 clEnv=getOpenCLEnvironment(exception);
1493 if (clEnv == (MagickCLEnv) NULL)
1494 return(MagickFalse);
1496 status=ComputeContrastStretchImage(image,clEnv,black_point,white_point,
1513 static Image *ComputeConvolveImage(
const Image* image,MagickCLEnv clEnv,
1517 *filteredImage_view,
1534 filteredImageBuffer,
1559 global_work_size[3],
1561 localMemoryRequirement;
1580 filteredImageBuffer = NULL;
1581 convolutionKernel = NULL;
1585 filteredImage = NULL;
1586 filteredImage_view = NULL;
1587 outputReady = MagickFalse;
1589 device = RequestOpenCLDevice(clEnv);
1591 image_view=AcquireAuthenticCacheView(image,exception);
1592 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
1593 if (inputPixels == (
const void *) NULL)
1595 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
1604 if (ALIGNED(inputPixels,CLPixelPacket))
1606 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1610 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1613 length = image->columns * image->rows;
1614 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1615 if (clStatus != CL_SUCCESS)
1617 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1621 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
1622 assert(filteredImage != NULL);
1623 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1625 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"CloneImage failed.",
".");
1628 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
1629 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
1630 if (filteredPixels == (
void *) NULL)
1632 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
1636 if (ALIGNED(filteredPixels,CLPixelPacket))
1638 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
1639 hostPtr = filteredPixels;
1643 mem_flags = CL_MEM_WRITE_ONLY;
1647 length = image->columns * image->rows;
1648 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
1649 if (clStatus != CL_SUCCESS)
1651 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1655 kernelSize = (
unsigned int) (kernel->width * kernel->height);
1656 convolutionKernel = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize *
sizeof(
float), NULL, &clStatus);
1657 if (clStatus != CL_SUCCESS)
1659 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1663 queue = AcquireOpenCLCommandQueue(device);
1665 kernelBufferPtr = (
float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize *
sizeof(
float)
1666 , 0, NULL, NULL, &clStatus);
1667 if (clStatus != CL_SUCCESS)
1669 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
1672 for (i = 0; i < kernelSize; i++)
1674 kernelBufferPtr[i] = (float) kernel->values[i];
1676 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
1677 if (clStatus != CL_SUCCESS)
1679 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
".");
1685 localGroupSize[0] = 16;
1686 localGroupSize[1] = 16;
1687 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) *
sizeof(CLPixelPacket)
1688 + kernel->width*kernel->height*
sizeof(
float);
1690 if (localMemoryRequirement > device->local_memory_size)
1692 localGroupSize[0] = 8;
1693 localGroupSize[1] = 8;
1694 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) *
sizeof(CLPixelPacket)
1695 + kernel->width*kernel->height*
sizeof(
float);
1697 if (localMemoryRequirement <= device->local_memory_size)
1700 clkernel = AcquireOpenCLKernel(device,
"ConvolveOptimized");
1701 if (clkernel == NULL)
1703 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
1709 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1710 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
1711 imageWidth = (
unsigned int) image->columns;
1712 imageHeight = (
unsigned int) image->rows;
1713 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageWidth);
1714 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageHeight);
1715 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&convolutionKernel);
1716 filterWidth = (
unsigned int) kernel->width;
1717 filterHeight = (
unsigned int) kernel->height;
1718 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterWidth);
1719 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterHeight);
1720 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
1721 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&matte);
1722 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(ChannelType),(
void *)&image->channel_mask);
1723 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*
sizeof(CLPixelPacket),NULL);
1724 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*
sizeof(
float),NULL);
1725 if (clStatus != CL_SUCCESS)
1727 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1732 global_work_size[0] = ((image->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ;
1733 global_work_size[1] = ((image->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1];
1736 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event);
1737 if (clStatus != CL_SUCCESS)
1739 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1742 RecordProfileData(device,clkernel,event);
1747 clkernel = AcquireOpenCLKernel(device,
"Convolve");
1748 if (clkernel == NULL)
1750 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
1756 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1757 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
1758 imageWidth = (
unsigned int) image->columns;
1759 imageHeight = (
unsigned int) image->rows;
1760 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageWidth);
1761 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageHeight);
1762 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&convolutionKernel);
1763 filterWidth = (
unsigned int) kernel->width;
1764 filterHeight = (
unsigned int) kernel->height;
1765 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterWidth);
1766 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterHeight);
1767 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
1768 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&matte);
1769 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(ChannelType),(
void *)&image->channel_mask);
1770 if (clStatus != CL_SUCCESS)
1772 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
1776 localGroupSize[0] = 8;
1777 localGroupSize[1] = 8;
1778 global_work_size[0] = (image->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0];
1779 global_work_size[1] = (image->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1];
1780 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event);
1782 if (clStatus != CL_SUCCESS)
1784 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
1788 RecordProfileData(device,clkernel,event);
1790 if (ALIGNED(filteredPixels,CLPixelPacket))
1792 length = image->columns * image->rows;
1793 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
1797 length = image->columns * image->rows;
1798 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
1800 if (clStatus != CL_SUCCESS)
1802 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
1806 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
1810 image_view=DestroyCacheView(image_view);
1811 if (filteredImage_view != NULL)
1812 filteredImage_view=DestroyCacheView(filteredImage_view);
1813 if (imageBuffer != NULL)
1814 clEnv->library->clReleaseMemObject(imageBuffer);
1815 if (filteredImageBuffer != NULL)
1816 clEnv->library->clReleaseMemObject(filteredImageBuffer);
1817 if (convolutionKernel != NULL)
1818 clEnv->library->clReleaseMemObject(convolutionKernel);
1819 if (clkernel != NULL)
1820 ReleaseOpenCLKernel(clkernel);
1822 ReleaseOpenCLCommandQueue(device,queue);
1824 ReleaseOpenCLDevice(device);
1825 if (outputReady == MagickFalse)
1827 if (filteredImage != NULL)
1829 DestroyImage(filteredImage);
1830 filteredImage = NULL;
1834 return(filteredImage);
1837 MagickPrivate
Image *AccelerateConvolveImage(
const Image *image,
1855 magick_unreferenced(image);
1856 magick_unreferenced(kernel);
1857 magick_unreferenced(exception);
1858 return((
Image *)NULL);
1873 static Image *ComputeDespeckleImage(
const Image *image,MagickCLEnv clEnv,
1877 X[4] = {0, 1, 1,-1},
1878 Y[4] = {1, 0, 1, 1};
1881 *filteredImage_view,
1901 filteredImageBuffer,
1925 global_work_size[2];
1935 outputReady = MagickFalse;
1937 filteredImage = NULL;
1938 filteredImage_view = NULL;
1939 filteredPixels = NULL;
1941 filteredImageBuffer = NULL;
1945 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
1947 device = RequestOpenCLDevice(clEnv);
1948 queue = AcquireOpenCLCommandQueue(device);
1950 image_view=AcquireAuthenticCacheView(image,exception);
1951 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
1952 if (inputPixels == (
void *) NULL)
1954 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
1958 if (ALIGNED(inputPixels,CLPixelPacket))
1960 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
1964 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
1967 length = image->columns * image->rows;
1968 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
1969 if (clStatus != CL_SUCCESS)
1971 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1975 mem_flags = CL_MEM_READ_WRITE;
1976 length = image->columns * image->rows;
1977 for (k = 0; k < 2; k++)
1979 tempImageBuffer[k] = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), NULL, &clStatus);
1980 if (clStatus != CL_SUCCESS)
1982 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1987 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
1988 assert(filteredImage != NULL);
1989 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
1991 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"CloneImage failed.",
".");
1994 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
1995 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
1996 if (filteredPixels == (
void *) NULL)
1998 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
2002 if (ALIGNED(filteredPixels,CLPixelPacket))
2004 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
2005 hostPtr = filteredPixels;
2009 mem_flags = CL_MEM_WRITE_ONLY;
2013 length = image->columns * image->rows;
2014 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
2015 if (clStatus != CL_SUCCESS)
2017 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2021 hullPass1 = AcquireOpenCLKernel(device,
"HullPass1");
2022 hullPass2 = AcquireOpenCLKernel(device,
"HullPass2");
2024 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,
sizeof(cl_mem),(
void *)&imageBuffer);
2025 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,
sizeof(cl_mem),(
void *)(tempImageBuffer+1));
2026 imageWidth = (
unsigned int) image->columns;
2027 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,
sizeof(
unsigned int),(
void *)&imageWidth);
2028 imageHeight = (
unsigned int) image->rows;
2029 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,
sizeof(
unsigned int),(
void *)&imageHeight);
2030 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
2031 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,
sizeof(
int),(
void *)&matte);
2032 if (clStatus != CL_SUCCESS)
2034 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2038 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,
sizeof(cl_mem),(
void *)(tempImageBuffer+1));
2039 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,
sizeof(cl_mem),(
void *)tempImageBuffer);
2040 imageWidth = (
unsigned int) image->columns;
2041 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,
sizeof(
unsigned int),(
void *)&imageWidth);
2042 imageHeight = (
unsigned int) image->rows;
2043 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,
sizeof(
unsigned int),(
void *)&imageHeight);
2044 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
2045 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,
sizeof(
int),(
void *)&matte);
2046 if (clStatus != CL_SUCCESS)
2048 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2053 global_work_size[0] = image->columns;
2054 global_work_size[1] = image->rows;
2057 for (k = 0; k < 4; k++)
2066 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2067 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2068 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2069 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2070 if (clStatus != CL_SUCCESS)
2072 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2076 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2077 if (clStatus != CL_SUCCESS)
2079 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2082 RecordProfileData(device,hullPass1,event);
2085 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2086 if (clStatus != CL_SUCCESS)
2088 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2091 RecordProfileData(device,hullPass2,event);
2094 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,
sizeof(cl_mem),(
void *)(tempImageBuffer));
2095 offset.s[0] = -X[k];
2096 offset.s[1] = -Y[k];
2098 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2099 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2100 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2101 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2102 if (clStatus != CL_SUCCESS)
2104 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2108 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2109 if (clStatus != CL_SUCCESS)
2111 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2114 RecordProfileData(device,hullPass1,event);
2117 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2118 if (clStatus != CL_SUCCESS)
2120 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2123 RecordProfileData(device,hullPass2,event);
2125 offset.s[0] = -X[k];
2126 offset.s[1] = -Y[k];
2128 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2129 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2130 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2131 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2132 if (clStatus != CL_SUCCESS)
2134 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2138 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2139 if (clStatus != CL_SUCCESS)
2141 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2144 RecordProfileData(device,hullPass1,event);
2147 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2148 if (clStatus != CL_SUCCESS)
2150 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2153 RecordProfileData(device,hullPass2,event);
2158 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2159 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2160 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2161 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2164 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
2166 if (clStatus != CL_SUCCESS)
2168 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2172 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2173 if (clStatus != CL_SUCCESS)
2175 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2178 RecordProfileData(device,hullPass1,event);
2181 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2182 if (clStatus != CL_SUCCESS)
2184 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2187 RecordProfileData(device,hullPass2,event);
2190 if (ALIGNED(filteredPixels,CLPixelPacket))
2192 length = image->columns * image->rows;
2193 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2197 length = image->columns * image->rows;
2198 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
2200 if (clStatus != CL_SUCCESS)
2202 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
2206 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
2210 image_view=DestroyCacheView(image_view);
2211 if (filteredImage_view != NULL)
2212 filteredImage_view=DestroyCacheView(filteredImage_view);
2215 ReleaseOpenCLCommandQueue(device,queue);
2217 ReleaseOpenCLDevice(device);
2218 if (imageBuffer!=NULL)
2219 clEnv->library->clReleaseMemObject(imageBuffer);
2220 for (k = 0; k < 2; k++)
2222 if (tempImageBuffer[k]!=NULL)
2223 clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
2225 if (filteredImageBuffer!=NULL)
2226 clEnv->library->clReleaseMemObject(filteredImageBuffer);
2227 if (hullPass1!=NULL)
2228 ReleaseOpenCLKernel(hullPass1);
2229 if (hullPass2!=NULL)
2230 ReleaseOpenCLKernel(hullPass2);
2231 if (outputReady == MagickFalse && filteredImage != NULL)
2232 filteredImage=DestroyImage(filteredImage);
2234 return(filteredImage);
2237 MagickPrivate
Image *AccelerateDespeckleImage(
const Image* image,
2246 assert(image != NULL);
2249 if (checkAccelerateConditionRGBA(image) == MagickFalse)
2250 return((
Image *) NULL);
2252 clEnv=getOpenCLEnvironment(exception);
2253 if (clEnv == (MagickCLEnv) NULL)
2254 return((
Image *) NULL);
2256 filteredImage=ComputeDespeckleImage(image,clEnv,exception);
2257 return(filteredImage);
2272 static MagickBooleanType ComputeEqualizeImage(
Image *image,MagickCLEnv clEnv,
2275 #define EqualizeImageTag "Equalize/Image"
2327 global_work_size[2];
2333 assert(image != (
Image *) NULL);
2334 assert(image->signature == MagickCoreSignature);
2335 if (IsEventLogging() != MagickFalse)
2336 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2343 histogramBuffer = NULL;
2344 equalizeMapBuffer = NULL;
2345 histogramKernel = NULL;
2346 equalizeKernel = NULL;
2348 outputReady = MagickFalse;
2353 device = RequestOpenCLDevice(clEnv);
2354 queue = AcquireOpenCLCommandQueue(device);
2359 histogram=(cl_uint4 *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*histogram));
2360 if (histogram == (cl_uint4 *) NULL)
2361 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
2364 (void) memset(histogram,0,(MaxMap+1)*
sizeof(*histogram));
2369 image_view=AcquireAuthenticCacheView(image,exception);
2370 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
2372 if (inputPixels == (
void *) NULL)
2374 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
2380 if (ALIGNED(inputPixels,CLPixelPacket))
2382 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2386 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
2389 length = image->columns * image->rows;
2390 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
2391 if (clStatus != CL_SUCCESS)
2393 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2400 if (ALIGNED(histogram,cl_uint4))
2402 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
2403 hostPtr = histogram;
2407 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
2408 hostPtr = histogram;
2411 length = (MaxMap+1);
2412 histogramBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(cl_uint4), hostPtr, &clStatus);
2413 if (clStatus != CL_SUCCESS)
2415 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2419 status = LaunchHistogramKernel(clEnv, device, queue, imageBuffer, histogramBuffer, image, image->channel_mask, exception);
2420 if (status == MagickFalse)
2424 if (ALIGNED(histogram,cl_uint4))
2426 length = (MaxMap+1);
2427 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(cl_uint4), 0, NULL, NULL, &clStatus);
2431 length = (MaxMap+1);
2432 clStatus = clEnv->library->clEnqueueReadBuffer(queue, histogramBuffer, CL_TRUE, 0, length *
sizeof(cl_uint4), histogram, 0, NULL, NULL);
2434 if (clStatus != CL_SUCCESS)
2436 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
2441 if (ALIGNED(histogram,cl_uint4))
2443 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
2444 if (clStatus != CL_SUCCESS)
2446 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
".");
2452 #ifdef RECREATEBUFFER
2453 if (imageBuffer!=NULL)
2454 clEnv->library->clReleaseMemObject(imageBuffer);
2458 equalize_map=(
PixelPacket *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*equalize_map));
2460 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
2462 map=(cl_float4 *) AcquireQuantumMemory(MaxMap+1UL,
sizeof(*map));
2463 if (map == (cl_float4 *) NULL)
2464 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
2469 (void) memset(&intensity,0,
sizeof(intensity));
2470 for (i=0; i <= (ssize_t) MaxMap; i++)
2472 if ((image->channel_mask & SyncChannels) != 0)
2474 intensity.x+=histogram[i].s[2];
2478 if ((image->channel_mask & RedChannel) != 0)
2479 intensity.x+=histogram[i].s[2];
2480 if ((image->channel_mask & GreenChannel) != 0)
2481 intensity.y+=histogram[i].s[1];
2482 if ((image->channel_mask & BlueChannel) != 0)
2483 intensity.z+=histogram[i].s[0];
2484 if ((image->channel_mask & AlphaChannel) != 0)
2485 intensity.w+=histogram[i].s[3];
2489 white=map[(int) MaxMap];
2490 (void) memset(equalize_map,0,(MaxMap+1)*
sizeof(*equalize_map));
2491 for (i=0; i <= (ssize_t) MaxMap; i++)
2493 if ((image->channel_mask & SyncChannels) != 0)
2495 if (white.x != black.x)
2496 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2497 (map[i].x-black.x))/(white.x-black.x)));
2500 if (((image->channel_mask & RedChannel) != 0) && (white.x != black.x))
2501 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2502 (map[i].x-black.x))/(white.x-black.x)));
2503 if (((image->channel_mask & GreenChannel) != 0) && (white.y != black.y))
2504 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2505 (map[i].y-black.y))/(white.y-black.y)));
2506 if (((image->channel_mask & BlueChannel) != 0) && (white.z != black.z))
2507 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2508 (map[i].z-black.z))/(white.z-black.z)));
2509 if (((image->channel_mask & AlphaChannel) != 0) && (white.w != black.w))
2510 equalize_map[i].alpha=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2511 (map[i].w-black.w))/(white.w-black.w)));
2514 if (image->storage_class == PseudoClass)
2519 for (i=0; i < (ssize_t) image->colors; i++)
2521 if ((image->channel_mask & SyncChannels) != 0)
2523 if (white.x != black.x)
2525 image->colormap[i].red=equalize_map[
2526 ScaleQuantumToMap(image->colormap[i].red)].red;
2527 image->colormap[i].green=equalize_map[
2528 ScaleQuantumToMap(image->colormap[i].green)].red;
2529 image->colormap[i].blue=equalize_map[
2530 ScaleQuantumToMap(image->colormap[i].blue)].red;
2531 image->colormap[i].alpha=equalize_map[
2532 ScaleQuantumToMap(image->colormap[i].alpha)].red;
2536 if (((image->channel_mask & RedChannel) != 0) && (white.x != black.x))
2537 image->colormap[i].red=equalize_map[
2538 ScaleQuantumToMap(image->colormap[i].red)].red;
2539 if (((image->channel_mask & GreenChannel) != 0) && (white.y != black.y))
2540 image->colormap[i].green=equalize_map[
2541 ScaleQuantumToMap(image->colormap[i].green)].green;
2542 if (((image->channel_mask & BlueChannel) != 0) && (white.z != black.z))
2543 image->colormap[i].blue=equalize_map[
2544 ScaleQuantumToMap(image->colormap[i].blue)].blue;
2545 if (((image->channel_mask & AlphaChannel) != 0) && (white.w != black.w))
2546 image->colormap[i].alpha=equalize_map[
2547 ScaleQuantumToMap(image->colormap[i].alpha)].alpha;
2560 #ifdef RECREATEBUFFER
2564 if (ALIGNED(inputPixels,CLPixelPacket))
2566 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
2570 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
2573 length = image->columns * image->rows;
2574 imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
2575 if (clStatus != CL_SUCCESS)
2577 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2585 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
2586 hostPtr = equalize_map;
2590 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
2591 hostPtr = equalize_map;
2594 length = (MaxMap+1);
2595 equalizeMapBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(
PixelPacket), hostPtr, &clStatus);
2596 if (clStatus != CL_SUCCESS)
2598 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2603 equalizeKernel = AcquireOpenCLKernel(device,
"Equalize");
2604 if (equalizeKernel == NULL)
2606 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2612 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2613 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(ChannelType),&image->channel_mask);
2614 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_mem),(
void *)&equalizeMapBuffer);
2615 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_float4),&white);
2616 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_float4),&black);
2617 if (clStatus != CL_SUCCESS)
2619 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
2624 global_work_size[0] = image->columns;
2625 global_work_size[1] = image->rows;
2627 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2629 if (clStatus != CL_SUCCESS)
2631 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
2634 RecordProfileData(device,equalizeKernel,event);
2637 if (ALIGNED(inputPixels,CLPixelPacket))
2639 length = image->columns * image->rows;
2640 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
2644 length = image->columns * image->rows;
2645 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
2647 if (clStatus != CL_SUCCESS)
2649 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
2653 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
2657 image_view=DestroyCacheView(image_view);
2659 if (imageBuffer!=NULL)
2660 clEnv->library->clReleaseMemObject(imageBuffer);
2662 map=(cl_float4 *) RelinquishMagickMemory(map);
2663 if (equalizeMapBuffer!=NULL)
2664 clEnv->library->clReleaseMemObject(equalizeMapBuffer);
2665 if (equalize_map!=NULL)
2666 equalize_map=(
PixelPacket *) RelinquishMagickMemory(equalize_map);
2667 if (histogramBuffer!=NULL)
2668 clEnv->library->clReleaseMemObject(histogramBuffer);
2669 if (histogram!=NULL)
2670 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
2671 if (histogramKernel!=NULL)
2672 ReleaseOpenCLKernel(histogramKernel);
2673 if (equalizeKernel!=NULL)
2674 ReleaseOpenCLKernel(equalizeKernel);
2676 ReleaseOpenCLCommandQueue(device, queue);
2678 ReleaseOpenCLDevice(device);
2680 return(outputReady);
2683 MagickPrivate MagickBooleanType AccelerateEqualizeImage(
Image *image,
2692 assert(image != NULL);
2694 if (IsEventLogging() != MagickFalse)
2695 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2697 if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
2698 (checkHistogramCondition(image,image->intensity) == MagickFalse))
2699 return(MagickFalse);
2701 clEnv=getOpenCLEnvironment(exception);
2702 if (clEnv == (MagickCLEnv) NULL)
2703 return(MagickFalse);
2705 status=ComputeEqualizeImage(image,clEnv,exception);
2721 static MagickBooleanType ComputeFunctionImage(
Image *image,MagickCLEnv clEnv,
2722 const MagickFunction
function,
const size_t number_parameters,
2743 *parametersBufferPtr;
2755 assert(image != (
Image *) NULL);
2756 assert(image->signature == MagickCoreSignature);
2757 if (IsEventLogging() != MagickFalse)
2758 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2760 outputReady=MagickFalse;
2762 functionKernel=NULL;
2763 parametersBuffer=NULL;
2765 device=RequestOpenCLDevice(clEnv);
2766 queue=AcquireOpenCLCommandQueue(device);
2767 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
2768 if (imageBuffer == (cl_mem) NULL)
2771 parametersBufferPtr=(
float *) AcquireQuantumMemory(number_parameters,
2773 if (parametersBufferPtr == (
float *) NULL)
2775 for (i=0; i<number_parameters; i++)
2776 parametersBufferPtr[i]=(
float) parameters[i];
2777 parametersBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
2778 CL_MEM_COPY_HOST_PTR,number_parameters*
sizeof(*parametersBufferPtr),
2779 parametersBufferPtr);
2780 parametersBufferPtr=RelinquishMagickMemory(parametersBufferPtr);
2781 if (parametersBuffer == (cl_mem) NULL)
2783 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2784 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
2788 functionKernel=AcquireOpenCLKernel(device,
"ComputeFunction");
2789 if (functionKernel == (cl_kernel) NULL)
2791 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2792 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2796 number_channels=(cl_uint) image->number_channels;
2797 number_params=(cl_uint) number_parameters;
2800 status =SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2801 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
2802 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(ChannelType),(
void *)&image->channel_mask);
2803 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(MagickFunction),(
void *)&
function);
2804 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_uint),(
void *)&number_params);
2805 status|=SetOpenCLKernelArg(functionKernel,i++,
sizeof(cl_mem),(
void *)¶metersBuffer);
2806 if (status != CL_SUCCESS)
2808 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2809 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
2813 gsize[0]=image->columns;
2814 gsize[1]=image->rows;
2815 outputReady=EnqueueOpenCLKernel(queue,functionKernel,2,(
const size_t *) NULL,
2816 gsize,(
const size_t *) NULL,image,(
const Image *) NULL,MagickFalse,
2821 if (imageBuffer != (cl_mem) NULL)
2822 ReleaseOpenCLMemObject(imageBuffer);
2823 if (parametersBuffer != (cl_mem) NULL)
2824 ReleaseOpenCLMemObject(parametersBuffer);
2825 if (functionKernel != (cl_kernel) NULL)
2826 ReleaseOpenCLKernel(functionKernel);
2827 if (queue != (cl_command_queue) NULL)
2828 ReleaseOpenCLCommandQueue(device,queue);
2829 if (device != (MagickCLDevice) NULL)
2830 ReleaseOpenCLDevice(device);
2831 return(outputReady);
2834 MagickPrivate MagickBooleanType AccelerateFunctionImage(
Image *image,
2835 const MagickFunction
function,
const size_t number_parameters,
2844 assert(image != NULL);
2846 if (IsEventLogging() != MagickFalse)
2847 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2849 if (checkAccelerateCondition(image) == MagickFalse)
2850 return(MagickFalse);
2852 clEnv=getOpenCLEnvironment(exception);
2853 if (clEnv == (MagickCLEnv) NULL)
2854 return(MagickFalse);
2856 status=ComputeFunctionImage(image,clEnv,
function,number_parameters,
2857 parameters,exception);
2873 static MagickBooleanType ComputeGrayscaleImage(
Image *image,MagickCLEnv clEnv,
2903 assert(image != (
Image *) NULL);
2904 assert(image->signature == MagickCoreSignature);
2905 if (IsEventLogging() != MagickFalse)
2906 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2908 outputReady=MagickFalse;
2910 grayscaleKernel=NULL;
2912 device=RequestOpenCLDevice(clEnv);
2913 queue=AcquireOpenCLCommandQueue(device);
2914 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
2915 if (imageBuffer == (cl_mem) NULL)
2918 grayscaleKernel=AcquireOpenCLKernel(device,
"Grayscale");
2919 if (grayscaleKernel == (cl_kernel) NULL)
2921 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2922 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
2926 number_channels=(cl_uint) image->number_channels;
2927 intensityMethod=(cl_uint) method;
2928 colorspace=(cl_uint) image->colorspace;
2931 status =SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2932 status|=SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_uint),&number_channels);
2933 status|=SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_uint),&colorspace);
2934 status|=SetOpenCLKernelArg(grayscaleKernel,i++,
sizeof(cl_uint),&intensityMethod);
2935 if (status != CL_SUCCESS)
2937 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
2938 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
2942 gsize[0]=image->columns;
2943 gsize[1]=image->rows;
2944 outputReady=EnqueueOpenCLKernel(queue,grayscaleKernel,2,
2945 (
const size_t *) NULL,gsize,(
const size_t *) NULL,image,(
Image *) NULL,
2946 MagickFalse,exception);
2950 if (imageBuffer != (cl_mem) NULL)
2951 ReleaseOpenCLMemObject(imageBuffer);
2952 if (grayscaleKernel != (cl_kernel) NULL)
2953 ReleaseOpenCLKernel(grayscaleKernel);
2954 if (queue != (cl_command_queue) NULL)
2955 ReleaseOpenCLCommandQueue(device,queue);
2956 if (device != (MagickCLDevice) NULL)
2957 ReleaseOpenCLDevice(device);
2959 return(outputReady);
2962 MagickPrivate MagickBooleanType AccelerateGrayscaleImage(
Image* image,
2971 assert(image != NULL);
2973 if (IsEventLogging() != MagickFalse)
2974 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2976 if ((checkAccelerateCondition(image) == MagickFalse) ||
2977 (checkPixelIntensity(image,method) == MagickFalse))
2978 return(MagickFalse);
2980 if (image->number_channels < 3)
2981 return(MagickFalse);
2983 if ((GetPixelRedTraits(image) == UndefinedPixelTrait) ||
2984 (GetPixelGreenTraits(image) == UndefinedPixelTrait) ||
2985 (GetPixelBlueTraits(image) == UndefinedPixelTrait))
2986 return(MagickFalse);
2988 clEnv=getOpenCLEnvironment(exception);
2989 if (clEnv == (MagickCLEnv) NULL)
2990 return(MagickFalse);
2992 status=ComputeGrayscaleImage(image,clEnv,method,exception);
3008 static Image *ComputeLocalContrastImage(
const Image *image,MagickCLEnv clEnv,
3009 const double radius,
const double strength,
ExceptionInfo *exception)
3012 *filteredImage_view,
3030 filteredImageBuffer,
3063 filteredImage = NULL;
3064 filteredImage_view = NULL;
3066 filteredImageBuffer = NULL;
3067 tempImageBuffer = NULL;
3068 imageKernelBuffer = NULL;
3069 blurRowKernel = NULL;
3070 blurColumnKernel = NULL;
3072 outputReady = MagickFalse;
3074 device = RequestOpenCLDevice(clEnv);
3075 queue = AcquireOpenCLCommandQueue(device);
3079 image_view=AcquireAuthenticCacheView(image,exception);
3080 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
3081 if (inputPixels == (
const void *) NULL)
3083 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
3090 if (ALIGNED(inputPixels,CLPixelPacket))
3092 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
3096 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
3099 length = image->columns * image->rows;
3100 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
3101 if (clStatus != CL_SUCCESS)
3103 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3110 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3111 assert(filteredImage != NULL);
3112 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
3114 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"CloneImage failed.",
".");
3117 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
3118 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
3119 if (filteredPixels == (
void *) NULL)
3121 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
3125 if (ALIGNED(filteredPixels,CLPixelPacket))
3127 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
3128 hostPtr = filteredPixels;
3132 mem_flags = CL_MEM_WRITE_ONLY;
3137 length = image->columns * image->rows;
3138 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
3139 if (clStatus != CL_SUCCESS)
3141 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3149 length = image->columns * image->rows;
3150 tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length *
sizeof(
float), NULL, &clStatus);
3151 if (clStatus != CL_SUCCESS)
3153 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3160 blurRowKernel = AcquireOpenCLKernel(device,
"LocalContrastBlurRow");
3161 if (blurRowKernel == NULL)
3163 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
3167 blurColumnKernel = AcquireOpenCLKernel(device,
"LocalContrastBlurApplyColumn");
3168 if (blurColumnKernel == NULL)
3170 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
3176 imageColumns = (
unsigned int) image->columns;
3177 imageRows = (
unsigned int) image->rows;
3178 iRadius = (cl_int) (image->rows > image->columns ? image->rows : image->columns) * 0.002f * fabs(radius);
3180 passes = (((1.0f * imageRows) * imageColumns * iRadius) + 3999999999) / 4000000000.0f;
3181 passes = (passes < 1) ? 1: passes;
3185 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
3186 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
3187 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
3188 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_int),(
void *)&iRadius);
3189 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
3190 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
3192 if (clStatus != CL_SUCCESS)
3194 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
3202 for (x = 0; x < passes; ++x) {
3208 gsize[1] = (image->rows + passes - 1) / passes;
3212 goffset[1] = x * gsize[1];
3214 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, 0, NULL, &event);
3215 if (clStatus != CL_SUCCESS)
3217 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
3220 clEnv->library->clFlush(queue);
3221 RecordProfileData(device,blurRowKernel,event);
3226 cl_float FStrength = strength;
3228 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
3229 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
3230 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
3231 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&iRadius);
3232 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_float),(
void *)&FStrength);
3233 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
3234 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
3236 if (clStatus != CL_SUCCESS)
3238 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
3246 for (x = 0; x < passes; ++x) {
3251 gsize[0] = ((image->columns + 3) / 4) * 4;
3252 gsize[1] = ((((image->rows + 63) / 64) + (passes + 1)) / passes) * 64;
3256 goffset[1] = x * gsize[1];
3258 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, 0, NULL, &event);
3259 if (clStatus != CL_SUCCESS)
3261 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
3264 clEnv->library->clFlush(queue);
3265 RecordProfileData(device,blurColumnKernel,event);
3271 if (ALIGNED(filteredPixels,CLPixelPacket))
3273 length = image->columns * image->rows;
3274 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3278 length = image->columns * image->rows;
3279 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
3281 if (clStatus != CL_SUCCESS)
3283 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
3287 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
3291 image_view=DestroyCacheView(image_view);
3292 if (filteredImage_view != NULL)
3293 filteredImage_view=DestroyCacheView(filteredImage_view);
3295 if (imageBuffer!=NULL)
3296 clEnv->library->clReleaseMemObject(imageBuffer);
3297 if (filteredImageBuffer!=NULL)
3298 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3299 if (tempImageBuffer!=NULL)
3300 clEnv->library->clReleaseMemObject(tempImageBuffer);
3301 if (imageKernelBuffer!=NULL)
3302 clEnv->library->clReleaseMemObject(imageKernelBuffer);
3303 if (blurRowKernel!=NULL)
3304 ReleaseOpenCLKernel(blurRowKernel);
3305 if (blurColumnKernel!=NULL)
3306 ReleaseOpenCLKernel(blurColumnKernel);
3308 ReleaseOpenCLCommandQueue(device, queue);
3310 ReleaseOpenCLDevice(device);
3311 if (outputReady == MagickFalse)
3313 if (filteredImage != NULL)
3315 DestroyImage(filteredImage);
3316 filteredImage = NULL;
3320 return(filteredImage);
3323 MagickPrivate
Image *AccelerateLocalContrastImage(
const Image *image,
3324 const double radius,
const double strength,
ExceptionInfo *exception)
3332 assert(image != NULL);
3335 if (checkAccelerateConditionRGBA(image) == MagickFalse)
3336 return((
Image *) NULL);
3338 clEnv=getOpenCLEnvironment(exception);
3339 if (clEnv == (MagickCLEnv) NULL)
3340 return((
Image *) NULL);
3342 filteredImage=ComputeLocalContrastImage(image,clEnv,radius,strength,
3344 return(filteredImage);
3359 static MagickBooleanType ComputeModulateImage(
Image *image,MagickCLEnv clEnv,
3360 const double percent_brightness,
const double percent_hue,
3361 const double percent_saturation,
const ColorspaceType colorspace,
3406 assert(image != (
Image *) NULL);
3407 assert(image->signature == MagickCoreSignature);
3408 if (IsEventLogging() != MagickFalse)
3409 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
3413 modulateKernel = NULL;
3418 device = RequestOpenCLDevice(clEnv);
3419 queue = AcquireOpenCLCommandQueue(device);
3421 outputReady = MagickFalse;
3427 image_view=AcquireAuthenticCacheView(image,exception);
3428 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception);
3429 if (inputPixels == (
void *) NULL)
3431 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,
"UnableToReadPixelCache.",
"`%s'",image->filename);
3439 if (ALIGNED(inputPixels,CLPixelPacket))
3441 mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
3445 mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
3448 length = image->columns * image->rows;
3449 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags, length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
3450 if (clStatus != CL_SUCCESS)
3452 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3456 modulateKernel = AcquireOpenCLKernel(device,
"Modulate");
3457 if (modulateKernel == NULL)
3459 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
3463 bright=percent_brightness;
3465 saturation=percent_saturation;
3469 clStatus=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
3470 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&bright);
3471 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&hue);
3472 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&saturation);
3473 clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,
sizeof(cl_float),&color);
3474 if (clStatus != CL_SUCCESS)
3476 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
3481 size_t global_work_size[2];
3482 global_work_size[0] = image->columns;
3483 global_work_size[1] = image->rows;
3485 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
3486 if (clStatus != CL_SUCCESS)
3488 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
".");
3491 RecordProfileData(device,modulateKernel,event);
3494 if (ALIGNED(inputPixels,CLPixelPacket))
3496 length = image->columns * image->rows;
3497 clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
3501 length = image->columns * image->rows;
3502 clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length *
sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
3504 if (clStatus != CL_SUCCESS)
3506 (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
".");
3510 outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
3514 image_view=DestroyCacheView(image_view);
3516 if (imageBuffer!=NULL)
3517 clEnv->library->clReleaseMemObject(imageBuffer);
3518 if (modulateKernel!=NULL)
3519 ReleaseOpenCLKernel(modulateKernel);
3521 ReleaseOpenCLCommandQueue(device,queue);
3523 ReleaseOpenCLDevice(device);
3529 MagickPrivate MagickBooleanType AccelerateModulateImage(
Image *image,
3530 const double percent_brightness,
const double percent_hue,
3531 const double percent_saturation,
const ColorspaceType colorspace,
3540 assert(image != NULL);
3542 if (IsEventLogging() != MagickFalse)
3543 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
3545 if (checkAccelerateConditionRGBA(image) == MagickFalse)
3546 return(MagickFalse);
3548 if ((colorspace != HSLColorspace) && (colorspace != UndefinedColorspace))
3549 return(MagickFalse);
3551 clEnv=getOpenCLEnvironment(exception);
3552 if (clEnv == (MagickCLEnv) NULL)
3553 return(MagickFalse);
3555 status=ComputeModulateImage(image,clEnv,percent_brightness,percent_hue,
3556 percent_saturation,colorspace,exception);
3572 static Image* ComputeMotionBlurImage(
const Image *image,MagickCLEnv clEnv,
3573 const double *kernel,
const size_t width,
const OffsetInfo *offset,
3577 *filteredImage_view,
3596 filteredImageBuffer,
3629 global_work_size[2],
3642 assert(image != (
Image *) NULL);
3643 assert(image->signature == MagickCoreSignature);
3644 if (IsEventLogging() != MagickFalse)
3645 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
3647 outputReady = MagickFalse;
3648 filteredImage = NULL;
3649 filteredImage_view = NULL;
3651 filteredImageBuffer = NULL;
3652 imageKernelBuffer = NULL;
3653 motionBlurKernel = NULL;
3656 device = RequestOpenCLDevice(clEnv);
3660 image_view=AcquireAuthenticCacheView(image,exception);
3661 inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,
3662 image->rows,exception);
3663 if (inputPixels == (
const void *) NULL)
3665 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
3666 "UnableToReadPixelCache.",
"`%s'",image->filename);
3675 if (ALIGNED(inputPixels,CLPixelPacket))
3677 mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
3681 mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR;
3686 length = image->columns * image->rows;
3687 imageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
3688 length *
sizeof(CLPixelPacket), (
void*)inputPixels, &clStatus);
3689 if (clStatus != CL_SUCCESS)
3691 (void) ThrowMagickException(exception, GetMagickModule(),
3692 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3697 filteredImage = CloneImage(image,image->columns,image->rows,
3698 MagickTrue,exception);
3699 assert(filteredImage != NULL);
3700 if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
3702 (void) ThrowMagickException(exception, GetMagickModule(),
3703 ResourceLimitError,
"CloneImage failed.",
".");
3706 filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
3707 filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
3708 if (filteredPixels == (
void *) NULL)
3710 (void) ThrowMagickException(exception,GetMagickModule(),CacheError,
3711 "UnableToReadPixelCache.",
"`%s'",filteredImage->filename);
3715 if (ALIGNED(filteredPixels,CLPixelPacket))
3717 mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
3718 hostPtr = filteredPixels;
3722 mem_flags = CL_MEM_WRITE_ONLY;
3728 length = image->columns * image->rows;
3729 filteredImageBuffer = clEnv->library->clCreateBuffer(device->context, mem_flags,
3730 length *
sizeof(CLPixelPacket), hostPtr, &clStatus);
3731 if (clStatus != CL_SUCCESS)
3733 (void) ThrowMagickException(exception, GetMagickModule(),
3734 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3739 imageKernelBuffer = clEnv->library->clCreateBuffer(device->context,
3740 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width *
sizeof(
float), NULL,
3742 if (clStatus != CL_SUCCESS)
3744 (void) ThrowMagickException(exception, GetMagickModule(),
3745 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3749 queue = AcquireOpenCLCommandQueue(device);
3750 kernelBufferPtr = (
float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
3751 CL_TRUE, CL_MAP_WRITE, 0, width *
sizeof(
float), 0, NULL, NULL, &clStatus);
3752 if (clStatus != CL_SUCCESS)
3754 (void) ThrowMagickException(exception, GetMagickModule(),
3755 ResourceLimitError,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
3758 for (i = 0; i < width; i++)
3760 kernelBufferPtr[i] = (float) kernel[i];
3762 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr,
3764 if (clStatus != CL_SUCCESS)
3766 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3767 "clEnv->library->clEnqueueUnmapMemObject failed.",
".");
3771 offsetBuffer = clEnv->library->clCreateBuffer(device->context,
3772 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width *
sizeof(cl_int2), NULL,
3774 if (clStatus != CL_SUCCESS)
3776 (void) ThrowMagickException(exception, GetMagickModule(),
3777 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3781 offsetBufferPtr = (
int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
3782 CL_MAP_WRITE, 0, width *
sizeof(cl_int2), 0, NULL, NULL, &clStatus);
3783 if (clStatus != CL_SUCCESS)
3785 (void) ThrowMagickException(exception, GetMagickModule(),
3786 ResourceLimitError,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
3789 for (i = 0; i < width; i++)
3791 offsetBufferPtr[2*i] = (int)offset[i].x;
3792 offsetBufferPtr[2*i+1] = (int)offset[i].y;
3794 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
3796 if (clStatus != CL_SUCCESS)
3798 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3799 "clEnv->library->clEnqueueUnmapMemObject failed.",
".");
3807 motionBlurKernel = AcquireOpenCLKernel(device,
"MotionBlur");
3808 if (motionBlurKernel == NULL)
3810 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3811 "AcquireOpenCLKernel failed.",
".");
3819 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3820 (
void *)&imageBuffer);
3821 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3822 (
void *)&filteredImageBuffer);
3823 imageWidth = (
unsigned int) image->columns;
3824 imageHeight = (
unsigned int) image->rows;
3825 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3827 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3829 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3830 (
void *)&imageKernelBuffer);
3831 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3833 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3834 (
void *)&offsetBuffer);
3836 GetPixelInfo(image,&bias);
3837 biasPixel.s[0] = bias.red;
3838 biasPixel.s[1] = bias.green;
3839 biasPixel.s[2] = bias.blue;
3840 biasPixel.s[3] = bias.alpha;
3841 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_float4), &biasPixel);
3843 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(ChannelType), &image->channel_mask);
3844 matte = (image->alpha_trait > CopyPixelTrait)?1:0;
3845 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int), &matte);
3846 if (clStatus != CL_SUCCESS)
3848 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3849 "clEnv->library->clSetKernelArg failed.",
".");
3856 local_work_size[0] = 16;
3857 local_work_size[1] = 16;
3858 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3859 (
unsigned int) image->columns,(
unsigned int) local_work_size[0]);
3860 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3861 (
unsigned int) image->rows,(
unsigned int) local_work_size[1]);
3862 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
3863 global_work_size, local_work_size, 0, NULL, &event);
3865 if (clStatus != CL_SUCCESS)
3867 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3868 "clEnv->library->clEnqueueNDRangeKernel failed.",
".");
3871 RecordProfileData(device,motionBlurKernel,event);
3873 if (ALIGNED(filteredPixels,CLPixelPacket))
3875 length = image->columns * image->rows;
3876 clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE,
3877 CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(CLPixelPacket), 0, NULL,
3882 length = image->columns * image->rows;
3883 clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0,
3884 length *
sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
3886 if (clStatus != CL_SUCCESS)
3888 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3889 "Reading output image from CL buffer failed.",
".");
3892 outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
3896 image_view=DestroyCacheView(image_view);
3897 if (filteredImage_view != NULL)
3898 filteredImage_view=DestroyCacheView(filteredImage_view);
3900 if (filteredImageBuffer!=NULL)
3901 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3902 if (imageBuffer!=NULL)
3903 clEnv->library->clReleaseMemObject(imageBuffer);
3904 if (imageKernelBuffer!=NULL)
3905 clEnv->library->clReleaseMemObject(imageKernelBuffer);
3906 if (motionBlurKernel!=NULL)
3907 ReleaseOpenCLKernel(motionBlurKernel);
3909 ReleaseOpenCLCommandQueue(device,queue);
3911 ReleaseOpenCLDevice(device);
3912 if (outputReady == MagickFalse && filteredImage != NULL)
3913 filteredImage=DestroyImage(filteredImage);
3915 return(filteredImage);
3918 MagickPrivate
Image *AccelerateMotionBlurImage(
const Image *image,
3919 const double* kernel,
const size_t width,
const OffsetInfo *offset,
3928 assert(image != NULL);
3929 assert(kernel != (
double *) NULL);
3933 if (checkAccelerateConditionRGBA(image) == MagickFalse)
3934 return((
Image *) NULL);
3936 clEnv=getOpenCLEnvironment(exception);
3937 if (clEnv == (MagickCLEnv) NULL)
3938 return((
Image *) NULL);
3940 filteredImage=ComputeMotionBlurImage(image,clEnv,kernel,width,offset,
3942 return(filteredImage);
3957 static MagickBooleanType resizeHorizontalFilter(MagickCLDevice device,
3958 cl_command_queue queue,
const Image *image,
Image *filteredImage,
3959 cl_mem imageBuffer,cl_uint number_channels,cl_uint columns,cl_uint rows,
3960 cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows,
3961 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3971 workgroupSize = 256;
3975 resizeFilterSupport,
3976 resizeFilterWindowSupport,
3990 gammaAccumulatorLocalMemorySize,
3993 imageCacheLocalMemorySize,
3994 pixelAccumulatorLocalMemorySize,
3996 totalLocalMemorySize,
3997 weightAccumulatorLocalMemorySize;
4003 horizontalKernel=NULL;
4004 outputReady=MagickFalse;
4009 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
4010 support=scale*GetResizeFilterSupport(resizeFilter);
4017 support=(float) 0.5;
4020 scale=PerceptibleReciprocal(scale);
4022 if (resizedColumns < workgroupSize)
4025 pixelPerWorkgroup=32;
4029 chunkSize=workgroupSize;
4030 pixelPerWorkgroup=workgroupSize;
4033 DisableMSCWarning(4127)
4038 numCachedPixels=(int) ceil((pixelPerWorkgroup-1)/xFactor+2*support);
4039 imageCacheLocalMemorySize=numCachedPixels*
sizeof(CLQuantum)*
4041 totalLocalMemorySize=imageCacheLocalMemorySize;
4044 pixelAccumulatorLocalMemorySize=chunkSize*
sizeof(cl_float4);
4045 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
4048 weightAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
4049 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
4052 if ((number_channels == 4) || (number_channels == 2))
4053 gammaAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
4055 gammaAccumulatorLocalMemorySize=
sizeof(float);
4056 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
4058 if (totalLocalMemorySize <= device->local_memory_size)
4062 pixelPerWorkgroup=pixelPerWorkgroup/2;
4063 chunkSize=chunkSize/2;
4064 if ((pixelPerWorkgroup == 0) || (chunkSize == 0))
4072 resizeFilterType=(int)GetResizeFilterWeightingType(resizeFilter);
4073 resizeWindowType=(int)GetResizeFilterWindowWeightingType(resizeFilter);
4075 horizontalKernel=AcquireOpenCLKernel(device,
"ResizeHorizontalFilter");
4076 if (horizontalKernel == (cl_kernel) NULL)
4078 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4079 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4083 resizeFilterScale=(float) GetResizeFilterScale(resizeFilter);
4084 resizeFilterSupport=(float) GetResizeFilterSupport(resizeFilter);
4085 resizeFilterWindowSupport=(float) GetResizeFilterWindowSupport(resizeFilter);
4086 resizeFilterBlur=(float) GetResizeFilterBlur(resizeFilter);
4089 status =SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_mem),(
void*)&imageBuffer);
4090 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&number_channels);
4091 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&columns);
4092 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&rows);
4093 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_mem),(
void*)&resizedImageBuffer);
4094 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&resizedColumns);
4095 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_uint),(
void*)&resizedRows);
4096 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&xFactor);
4097 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
int),(
void*)&resizeFilterType);
4098 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
int),(
void*)&resizeWindowType);
4099 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(cl_mem),(
void*)&resizeFilterCubicCoefficients);
4100 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterScale);
4101 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterSupport);
4102 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterWindowSupport);
4103 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
float),(
void*)&resizeFilterBlur);
4104 status|=SetOpenCLKernelArg(horizontalKernel,i++,imageCacheLocalMemorySize,NULL);
4105 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
int),&numCachedPixels);
4106 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
unsigned int),&pixelPerWorkgroup);
4107 status|=SetOpenCLKernelArg(horizontalKernel,i++,
sizeof(
unsigned int),&chunkSize);
4108 status|=SetOpenCLKernelArg(horizontalKernel,i++,pixelAccumulatorLocalMemorySize,NULL);
4109 status|=SetOpenCLKernelArg(horizontalKernel,i++,weightAccumulatorLocalMemorySize,NULL);
4110 status|=SetOpenCLKernelArg(horizontalKernel,i++,gammaAccumulatorLocalMemorySize,NULL);
4112 if (status != CL_SUCCESS)
4114 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4115 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
4119 gsize[0]=(resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*
4121 gsize[1]=resizedRows;
4122 lsize[0]=workgroupSize;
4124 outputReady=EnqueueOpenCLKernel(queue,horizontalKernel,2,
4125 (
const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
4130 if (horizontalKernel != (cl_kernel) NULL)
4131 ReleaseOpenCLKernel(horizontalKernel);
4133 return(outputReady);
4136 static MagickBooleanType resizeVerticalFilter(MagickCLDevice device,
4137 cl_command_queue queue,
const Image *image,
Image * filteredImage,
4138 cl_mem imageBuffer,cl_uint number_channels,cl_uint columns,cl_uint rows,
4139 cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows,
4140 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
4150 workgroupSize = 256;
4154 resizeFilterSupport,
4155 resizeFilterWindowSupport,
4169 gammaAccumulatorLocalMemorySize,
4172 imageCacheLocalMemorySize,
4173 pixelAccumulatorLocalMemorySize,
4175 totalLocalMemorySize,
4176 weightAccumulatorLocalMemorySize;
4182 verticalKernel=NULL;
4183 outputReady=MagickFalse;
4188 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
4189 support=scale*GetResizeFilterSupport(resizeFilter);
4196 support=(float) 0.5;
4199 scale=PerceptibleReciprocal(scale);
4201 if (resizedRows < workgroupSize)
4204 pixelPerWorkgroup=32;
4208 chunkSize=workgroupSize;
4209 pixelPerWorkgroup=workgroupSize;
4212 DisableMSCWarning(4127)
4217 numCachedPixels=(int)ceil((pixelPerWorkgroup-1)/yFactor+2*support);
4218 imageCacheLocalMemorySize=numCachedPixels*
sizeof(CLQuantum)*
4220 totalLocalMemorySize=imageCacheLocalMemorySize;
4223 pixelAccumulatorLocalMemorySize=chunkSize*
sizeof(cl_float4);
4224 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
4227 weightAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
4228 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
4231 if ((number_channels == 4) || (number_channels == 2))
4232 gammaAccumulatorLocalMemorySize=chunkSize*
sizeof(float);
4234 gammaAccumulatorLocalMemorySize=
sizeof(float);
4235 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
4237 if (totalLocalMemorySize <= device->local_memory_size)
4241 pixelPerWorkgroup=pixelPerWorkgroup/2;
4242 chunkSize=chunkSize/2;
4243 if ((pixelPerWorkgroup == 0) || (chunkSize == 0))
4251 resizeFilterType=(int)GetResizeFilterWeightingType(resizeFilter);
4252 resizeWindowType=(int)GetResizeFilterWindowWeightingType(resizeFilter);
4254 verticalKernel=AcquireOpenCLKernel(device,
"ResizeVerticalFilter");
4255 if (verticalKernel == (cl_kernel) NULL)
4257 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4258 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4262 resizeFilterScale=(float) GetResizeFilterScale(resizeFilter);
4263 resizeFilterSupport=(float) GetResizeFilterSupport(resizeFilter);
4264 resizeFilterBlur=(float) GetResizeFilterBlur(resizeFilter);
4265 resizeFilterWindowSupport=(float) GetResizeFilterWindowSupport(resizeFilter);
4268 status =SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_mem),(
void*)&imageBuffer);
4269 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&number_channels);
4270 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&columns);
4271 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&rows);
4272 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_mem),(
void*)&resizedImageBuffer);
4273 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&resizedColumns);
4274 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_uint),(
void*)&resizedRows);
4275 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&yFactor);
4276 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
int),(
void*)&resizeFilterType);
4277 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
int),(
void*)&resizeWindowType);
4278 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(cl_mem),(
void*)&resizeFilterCubicCoefficients);
4279 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterScale);
4280 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterSupport);
4281 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterWindowSupport);
4282 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
float),(
void*)&resizeFilterBlur);
4283 status|=SetOpenCLKernelArg(verticalKernel,i++,imageCacheLocalMemorySize, NULL);
4284 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
int), &numCachedPixels);
4285 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
unsigned int), &pixelPerWorkgroup);
4286 status|=SetOpenCLKernelArg(verticalKernel,i++,
sizeof(
unsigned int), &chunkSize);
4287 status|=SetOpenCLKernelArg(verticalKernel,i++,pixelAccumulatorLocalMemorySize, NULL);
4288 status|=SetOpenCLKernelArg(verticalKernel,i++,weightAccumulatorLocalMemorySize, NULL);
4289 status|=SetOpenCLKernelArg(verticalKernel,i++,gammaAccumulatorLocalMemorySize, NULL);
4291 if (status != CL_SUCCESS)
4293 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4294 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
4298 gsize[0]=resizedColumns;
4299 gsize[1]=(resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*
4302 lsize[1]=workgroupSize;
4303 outputReady=EnqueueOpenCLKernel(queue,verticalKernel,2,(
const size_t *) NULL,
4304 gsize,lsize,image,filteredImage,MagickFalse,exception);
4308 if (verticalKernel != (cl_kernel) NULL)
4309 ReleaseOpenCLKernel(verticalKernel);
4311 return(outputReady);
4314 static Image *ComputeResizeImage(
const Image* image,MagickCLEnv clEnv,
4315 const size_t resizedColumns,
const size_t resizedRows,
4322 cubicCoefficientsBuffer,
4323 filteredImageBuffer,
4331 *resizeFilterCoefficient;
4334 coefficientBuffer[7],
4355 filteredImageBuffer=NULL;
4356 tempImageBuffer=NULL;
4357 cubicCoefficientsBuffer=NULL;
4358 outputReady=MagickFalse;
4360 device=RequestOpenCLDevice(clEnv);
4361 queue=AcquireOpenCLCommandQueue(device);
4362 filteredImage=CloneImage(image,resizedColumns,resizedRows,MagickTrue,
4364 if (filteredImage == (
Image *) NULL)
4366 if (filteredImage->number_channels != image->number_channels)
4368 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4369 if (imageBuffer == (cl_mem) NULL)
4371 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4372 if (filteredImageBuffer == (cl_mem) NULL)
4375 resizeFilterCoefficient=GetResizeFilterCoefficient(resizeFilter);
4376 for (i = 0; i < 7; i++)
4377 coefficientBuffer[i]=(
float) resizeFilterCoefficient[i];
4378 cubicCoefficientsBuffer=CreateOpenCLBuffer(device,CL_MEM_COPY_HOST_PTR |
4379 CL_MEM_READ_ONLY,
sizeof(coefficientBuffer),&coefficientBuffer);
4380 if (cubicCoefficientsBuffer == (cl_mem) NULL)
4382 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4383 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4387 number_channels=(cl_uint) image->number_channels;
4388 xFactor=(
float) resizedColumns/(float) image->columns;
4389 yFactor=(
float) resizedRows/(float) image->rows;
4390 if (xFactor > yFactor)
4392 length=resizedColumns*image->rows*number_channels;
4393 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
4394 sizeof(CLQuantum),(
void *) NULL);
4395 if (tempImageBuffer == (cl_mem) NULL)
4397 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4398 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4402 outputReady=resizeHorizontalFilter(device,queue,image,filteredImage,
4403 imageBuffer,number_channels,(cl_uint) image->columns,
4404 (cl_uint) image->rows,tempImageBuffer,(cl_uint) resizedColumns,
4405 (cl_uint) image->rows,resizeFilter,cubicCoefficientsBuffer,xFactor,
4407 if (outputReady == MagickFalse)
4410 outputReady=resizeVerticalFilter(device,queue,image,filteredImage,
4411 tempImageBuffer,number_channels,(cl_uint) resizedColumns,
4412 (cl_uint) image->rows,filteredImageBuffer,(cl_uint) resizedColumns,
4413 (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,yFactor,
4415 if (outputReady == MagickFalse)
4420 length=image->columns*resizedRows*number_channels;
4421 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
4422 sizeof(CLQuantum),(
void *) NULL);
4423 if (tempImageBuffer == (cl_mem) NULL)
4425 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4426 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4430 outputReady=resizeVerticalFilter(device,queue,image,filteredImage,
4431 imageBuffer,number_channels,(cl_uint) image->columns,
4432 (cl_int) image->rows,tempImageBuffer,(cl_uint) image->columns,
4433 (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,yFactor,
4435 if (outputReady == MagickFalse)
4438 outputReady=resizeHorizontalFilter(device,queue,image,filteredImage,
4439 tempImageBuffer,number_channels,(cl_uint) image->columns,
4440 (cl_uint) resizedRows,filteredImageBuffer,(cl_uint) resizedColumns,
4441 (cl_uint) resizedRows,resizeFilter,cubicCoefficientsBuffer,xFactor,
4443 if (outputReady == MagickFalse)
4449 if (imageBuffer != (cl_mem) NULL)
4450 ReleaseOpenCLMemObject(imageBuffer);
4451 if (filteredImageBuffer != (cl_mem) NULL)
4452 ReleaseOpenCLMemObject(filteredImageBuffer);
4453 if (tempImageBuffer != (cl_mem) NULL)
4454 ReleaseOpenCLMemObject(tempImageBuffer);
4455 if (cubicCoefficientsBuffer != (cl_mem) NULL)
4456 ReleaseOpenCLMemObject(cubicCoefficientsBuffer);
4457 if (queue != (cl_command_queue) NULL)
4458 ReleaseOpenCLCommandQueue(device,queue);
4459 if (device != (MagickCLDevice) NULL)
4460 ReleaseOpenCLDevice(device);
4461 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4462 filteredImage=DestroyImage(filteredImage);
4464 return(filteredImage);
4467 static MagickBooleanType gpuSupportedResizeWeighting(
4468 ResizeWeightingFunctionType f)
4475 if (supportedResizeWeighting[i] == LastWeightingFunction)
4477 if (supportedResizeWeighting[i] == f)
4480 return(MagickFalse);
4483 MagickPrivate
Image *AccelerateResizeImage(
const Image *image,
4484 const size_t resizedColumns,
const size_t resizedRows,
4493 assert(image != NULL);
4496 if (checkAccelerateCondition(image) == MagickFalse)
4497 return((
Image *) NULL);
4499 if ((gpuSupportedResizeWeighting(GetResizeFilterWeightingType(
4500 resizeFilter)) == MagickFalse) ||
4501 (gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(
4502 resizeFilter)) == MagickFalse))
4503 return((
Image *) NULL);
4505 clEnv=getOpenCLEnvironment(exception);
4506 if (clEnv == (MagickCLEnv) NULL)
4507 return((
Image *) NULL);
4509 filteredImage=ComputeResizeImage(image,clEnv,resizedColumns,resizedRows,
4510 resizeFilter,exception);
4511 return(filteredImage);
4526 static Image* ComputeRotationalBlurImage(
const Image *image,MagickCLEnv clEnv,
4540 filteredImageBuffer,
4545 rotationalBlurKernel;
4571 assert(image != (
Image *) NULL);
4572 assert(image->signature == MagickCoreSignature);
4573 if (IsEventLogging() != MagickFalse)
4574 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
4578 filteredImageBuffer=NULL;
4579 sinThetaBuffer=NULL;
4580 cosThetaBuffer=NULL;
4581 rotationalBlurKernel=NULL;
4582 outputReady=MagickFalse;
4584 device=RequestOpenCLDevice(clEnv);
4585 queue=AcquireOpenCLCommandQueue(device);
4586 filteredImage=cloneImage(image,exception);
4587 if (filteredImage == (
Image *) NULL)
4589 if (filteredImage->number_channels != image->number_channels)
4591 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4592 if (imageBuffer == (cl_mem) NULL)
4594 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4595 if (filteredImageBuffer == (cl_mem) NULL)
4598 blurCenter.x=(float) (image->columns-1)/2.0;
4599 blurCenter.y=(float) (image->rows-1)/2.0;
4600 blurRadius=hypot(blurCenter.x,blurCenter.y);
4601 cossin_theta_size=(
unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt(
4602 (
double) blurRadius)+2UL);
4604 cosThetaPtr=AcquireQuantumMemory(cossin_theta_size,
sizeof(
float));
4605 if (cosThetaPtr == (
float *) NULL)
4607 sinThetaPtr=AcquireQuantumMemory(cossin_theta_size,
sizeof(
float));
4608 if (sinThetaPtr == (
float *) NULL)
4610 cosThetaPtr=RelinquishMagickMemory(cosThetaPtr);
4614 theta=DegreesToRadians(angle)/(double) (cossin_theta_size-1);
4615 offset=theta*(float) (cossin_theta_size-1)/2.0;
4616 for (i=0; i < (ssize_t) cossin_theta_size; i++)
4618 cosThetaPtr[i]=(float)cos((
double) (theta*i-offset));
4619 sinThetaPtr[i]=(float)sin((
double) (theta*i-offset));
4622 sinThetaBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
4623 CL_MEM_COPY_HOST_PTR,cossin_theta_size*
sizeof(
float),sinThetaPtr);
4624 sinThetaPtr=RelinquishMagickMemory(sinThetaPtr);
4625 cosThetaBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
4626 CL_MEM_COPY_HOST_PTR,cossin_theta_size*
sizeof(
float),cosThetaPtr);
4627 cosThetaPtr=RelinquishMagickMemory(cosThetaPtr);
4628 if ((sinThetaBuffer == (cl_mem) NULL) || (cosThetaBuffer == (cl_mem) NULL))
4630 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4631 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4635 rotationalBlurKernel=AcquireOpenCLKernel(device,
"RotationalBlur");
4636 if (rotationalBlurKernel == (cl_kernel) NULL)
4638 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4639 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4643 number_channels=(cl_uint) image->number_channels;
4646 status =SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4647 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_uint),&number_channels);
4648 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(ChannelType), &image->channel_mask);
4649 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_float2), &blurCenter);
4650 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&cosThetaBuffer);
4651 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&sinThetaBuffer);
4652 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_uint), &cossin_theta_size);
4653 status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4654 if (status != CL_SUCCESS)
4656 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4657 ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
4661 gsize[0]=image->columns;
4662 gsize[1]=image->rows;
4663 outputReady=EnqueueOpenCLKernel(queue,rotationalBlurKernel,2,
4664 (
const size_t *) NULL,gsize,(
const size_t *) NULL,image,filteredImage,
4665 MagickFalse,exception);
4669 if (imageBuffer != (cl_mem) NULL)
4670 ReleaseOpenCLMemObject(imageBuffer);
4671 if (filteredImageBuffer != (cl_mem) NULL)
4672 ReleaseOpenCLMemObject(filteredImageBuffer);
4673 if (sinThetaBuffer != (cl_mem) NULL)
4674 ReleaseOpenCLMemObject(sinThetaBuffer);
4675 if (cosThetaBuffer != (cl_mem) NULL)
4676 ReleaseOpenCLMemObject(cosThetaBuffer);
4677 if (rotationalBlurKernel != (cl_kernel) NULL)
4678 ReleaseOpenCLKernel(rotationalBlurKernel);
4679 if (queue != (cl_command_queue) NULL)
4680 ReleaseOpenCLCommandQueue(device,queue);
4681 if (device != (MagickCLDevice) NULL)
4682 ReleaseOpenCLDevice(device);
4683 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4684 filteredImage=DestroyImage(filteredImage);
4686 return(filteredImage);
4689 MagickPrivate
Image* AccelerateRotationalBlurImage(
const Image *image,
4698 assert(image != NULL);
4700 if (IsEventLogging() != MagickFalse)
4701 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
4703 if (checkAccelerateCondition(image) == MagickFalse)
4704 return((
Image *) NULL);
4706 clEnv=getOpenCLEnvironment(exception);
4707 if (clEnv == (MagickCLEnv) NULL)
4708 return((
Image *) NULL);
4710 filteredImage=ComputeRotationalBlurImage(image,clEnv,angle,exception);
4711 return filteredImage;
4726 static Image *ComputeUnsharpMaskImage(
const Image *image,MagickCLEnv clEnv,
4727 const double radius,
const double sigma,
const double gain,
4738 unsharpMaskBlurColumnKernel;
4741 filteredImageBuffer,
4778 filteredImageBuffer=NULL;
4779 tempImageBuffer=NULL;
4780 imageKernelBuffer=NULL;
4782 unsharpMaskBlurColumnKernel=NULL;
4783 outputReady=MagickFalse;
4785 device=RequestOpenCLDevice(clEnv);
4786 queue=AcquireOpenCLCommandQueue(device);
4787 filteredImage=cloneImage(image,exception);
4788 if (filteredImage == (
Image *) NULL)
4790 if (filteredImage->number_channels != image->number_channels)
4792 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4793 if (imageBuffer == (cl_mem) NULL)
4795 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4796 if (filteredImageBuffer == (cl_mem) NULL)
4799 imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
4802 length=image->columns*image->rows;
4803 tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
4804 sizeof(cl_float4),NULL);
4805 if (tempImageBuffer == (cl_mem) NULL)
4807 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4808 ResourceLimitWarning,
"CreateOpenCLBuffer failed.",
".");
4812 blurRowKernel=AcquireOpenCLKernel(device,
"BlurRow");
4813 if (blurRowKernel == (cl_kernel) NULL)
4815 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4816 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4820 unsharpMaskBlurColumnKernel=AcquireOpenCLKernel(device,
4821 "UnsharpMaskBlurColumn");
4822 if (unsharpMaskBlurColumnKernel == (cl_kernel) NULL)
4824 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4825 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4829 number_channels=(cl_uint) image->number_channels;
4830 imageColumns=(cl_uint) image->columns;
4831 imageRows=(cl_uint) image->rows;
4836 status =SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4837 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),&number_channels);
4838 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(ChannelType),&image->channel_mask);
4839 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4840 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
4841 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
4842 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
4843 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_float4)*(chunkSize+kernelWidth),(
void *) NULL);
4844 status|=SetOpenCLKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
4845 if (status != CL_SUCCESS)
4847 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4848 ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
4852 gsize[0]=chunkSize*((image->columns+chunkSize-1)/chunkSize);
4853 gsize[1]=image->rows;
4856 outputReady=EnqueueOpenCLKernel(queue,blurRowKernel,2,
4857 (
const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
4862 fThreshold=(float) threshold;
4865 status =SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4866 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
4867 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),&number_channels);
4868 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(ChannelType),&image->channel_mask);
4869 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
4870 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
4871 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,(chunkSize+kernelWidth-1)*
sizeof(cl_float4),NULL);
4872 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,kernelWidth*
sizeof(
float),NULL);
4873 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4874 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
4875 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
float),(
void *)&fGain);
4876 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
float),(
void *)&fThreshold);
4877 status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4878 if (status != CL_SUCCESS)
4880 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4881 ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
".");
4885 gsize[0]=image->columns;
4886 gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize);
4889 outputReady=EnqueueOpenCLKernel(queue,unsharpMaskBlurColumnKernel,2,
4890 (
const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
4895 if (imageBuffer != (cl_mem) NULL)
4896 ReleaseOpenCLMemObject(imageBuffer);
4897 if (filteredImageBuffer != (cl_mem) NULL)
4898 ReleaseOpenCLMemObject(filteredImageBuffer);
4899 if (tempImageBuffer != (cl_mem) NULL)
4900 ReleaseOpenCLMemObject(tempImageBuffer);
4901 if (imageKernelBuffer != (cl_mem) NULL)
4902 ReleaseOpenCLMemObject(imageKernelBuffer);
4903 if (blurRowKernel != (cl_kernel) NULL)
4904 ReleaseOpenCLKernel(blurRowKernel);
4905 if (unsharpMaskBlurColumnKernel != (cl_kernel) NULL)
4906 ReleaseOpenCLKernel(unsharpMaskBlurColumnKernel);
4907 if (queue != (cl_command_queue) NULL)
4908 ReleaseOpenCLCommandQueue(device,queue);
4909 if (device != (MagickCLDevice) NULL)
4910 ReleaseOpenCLDevice(device);
4911 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
4912 filteredImage=DestroyImage(filteredImage);
4914 return(filteredImage);
4917 static Image *ComputeUnsharpMaskImageSingle(
const Image *image,
4918 MagickCLEnv clEnv,
const double radius,
const double sigma,
const double gain,
4931 filteredImageBuffer,
4961 filteredImageBuffer=NULL;
4962 imageKernelBuffer=NULL;
4963 unsharpMaskKernel=NULL;
4964 outputReady=MagickFalse;
4966 device=RequestOpenCLDevice(clEnv);
4967 queue=AcquireOpenCLCommandQueue(device);
4968 filteredImage=cloneImage(image,exception);
4969 if (filteredImage == (
Image *) NULL)
4971 if (filteredImage->number_channels != image->number_channels)
4973 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
4974 if (imageBuffer == (cl_mem) NULL)
4976 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
4977 if (filteredImageBuffer == (cl_mem) NULL)
4980 imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
4983 unsharpMaskKernel=AcquireOpenCLKernel(device,
"UnsharpMask");
4984 if (unsharpMaskKernel == NULL)
4986 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
4987 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
4991 imageColumns=(cl_uint) image->columns;
4992 imageRows=(cl_uint) image->rows;
4993 number_channels=(cl_uint) image->number_channels;
4995 fThreshold=(float) threshold;
4998 status =SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4999 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
5000 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(ChannelType),(
void *)&image->channel_mask);
5001 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
5002 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&kernelWidth);
5003 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&imageColumns);
5004 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&imageRows);
5005 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_float4)*(8 * (32 + kernelWidth)),(
void *) NULL);
5006 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(
float),(
void *)&fGain);
5007 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(
float),(
void *)&fThreshold);
5008 status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
5009 if (status != CL_SUCCESS)
5011 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
5012 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
5016 gsize[0]=((image->columns + 7) / 8)*8;
5017 gsize[1]=((image->rows + 31) / 32)*32;
5020 outputReady=EnqueueOpenCLKernel(queue,unsharpMaskKernel,2,(
const size_t *) NULL,
5021 gsize,lsize,image,filteredImage,MagickFalse,exception);
5025 if (imageBuffer != (cl_mem) NULL)
5026 ReleaseOpenCLMemObject(imageBuffer);
5027 if (filteredImageBuffer != (cl_mem) NULL)
5028 ReleaseOpenCLMemObject(filteredImageBuffer);
5029 if (imageKernelBuffer != (cl_mem) NULL)
5030 ReleaseOpenCLMemObject(imageKernelBuffer);
5031 if (unsharpMaskKernel != (cl_kernel) NULL)
5032 ReleaseOpenCLKernel(unsharpMaskKernel);
5033 if (queue != (cl_command_queue) NULL)
5034 ReleaseOpenCLCommandQueue(device,queue);
5035 if (device != (MagickCLDevice) NULL)
5036 ReleaseOpenCLDevice(device);
5037 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
5038 filteredImage=DestroyImage(filteredImage);
5040 return(filteredImage);
5043 MagickPrivate
Image *AccelerateUnsharpMaskImage(
const Image *image,
5044 const double radius,
const double sigma,
const double gain,
5053 assert(image != NULL);
5056 if (checkAccelerateCondition(image) == MagickFalse)
5057 return((
Image *) NULL);
5059 clEnv=getOpenCLEnvironment(exception);
5060 if (clEnv == (MagickCLEnv) NULL)
5061 return((
Image *) NULL);
5064 filteredImage=ComputeUnsharpMaskImageSingle(image,clEnv,radius,sigma,gain,
5065 threshold,exception);
5067 filteredImage=ComputeUnsharpMaskImage(image,clEnv,radius,sigma,gain,
5068 threshold,exception);
5069 return(filteredImage);
5072 static Image *ComputeWaveletDenoiseImage(
const Image *image,MagickCLEnv clEnv,
5084 SIZE=TILESIZE-2*PAD;
5096 filteredImageBuffer,
5124 filteredImageBuffer=NULL;
5127 outputReady=MagickFalse;
5129 device=RequestOpenCLDevice(clEnv);
5131 if (strcmp(
"Intel(R) HD Graphics",device->name) == 0)
5133 queue=AcquireOpenCLCommandQueue(device);
5134 filteredImage=CloneImage(image,0,0,MagickTrue,
5136 if (filteredImage == (
Image *) NULL)
5138 if (filteredImage->number_channels != image->number_channels)
5140 imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
5141 if (imageBuffer == (cl_mem) NULL)
5143 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
5144 if (filteredImageBuffer == (cl_mem) NULL)
5147 denoiseKernel=AcquireOpenCLKernel(device,
"WaveletDenoise");
5148 if (denoiseKernel == (cl_kernel) NULL)
5150 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
5151 ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
".");
5155 number_channels=(cl_uint)image->number_channels;
5156 width=(cl_uint)image->columns;
5157 height=(cl_uint)image->rows;
5158 max_channels=number_channels;
5159 if ((max_channels == 4) || (max_channels == 2))
5160 max_channels=max_channels-1;
5162 passes=(((1.0f*image->columns)*image->rows)+1999999.0f)/2000000.0f;
5163 passes=(passes < 1) ? 1 : passes;
5166 status =SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
5167 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
5168 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&number_channels);
5169 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&max_channels);
5170 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_float),(
void *)&thresh);
5171 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_int),(
void *)&PASSES);
5172 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&width);
5173 status|=SetOpenCLKernelArg(denoiseKernel,i++,
sizeof(cl_uint),(
void *)&height);
5174 if (status != CL_SUCCESS)
5176 (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
5177 ResourceLimitWarning,
"SetOpenCLKernelArg failed.",
".");
5181 for (x = 0; x < passes; ++x)
5183 gsize[0]=((width+(SIZE-1))/SIZE)*TILESIZE;
5184 gsize[1]=((((height+(SIZE-1))/SIZE)+passes-1)/passes)*4;
5188 goffset[1]=x*gsize[1];
5190 outputReady=EnqueueOpenCLKernel(queue,denoiseKernel,2,goffset,gsize,lsize,
5191 image,filteredImage,MagickTrue,exception);
5192 if (outputReady == MagickFalse)
5198 if (imageBuffer != (cl_mem) NULL)
5199 ReleaseOpenCLMemObject(imageBuffer);
5200 if (filteredImageBuffer != (cl_mem) NULL)
5201 ReleaseOpenCLMemObject(filteredImageBuffer);
5202 if (denoiseKernel != (cl_kernel) NULL)
5203 ReleaseOpenCLKernel(denoiseKernel);
5204 if (queue != (cl_command_queue) NULL)
5205 ReleaseOpenCLCommandQueue(device,queue);
5206 if (device != (MagickCLDevice) NULL)
5207 ReleaseOpenCLDevice(device);
5208 if ((outputReady == MagickFalse) && (filteredImage != (
Image *) NULL))
5209 filteredImage=DestroyImage(filteredImage);
5211 return(filteredImage);
5214 MagickPrivate
Image *AccelerateWaveletDenoiseImage(
const Image *image,
5223 assert(image != NULL);
5226 if (checkAccelerateCondition(image) == MagickFalse)
5227 return((
Image *) NULL);
5229 clEnv=getOpenCLEnvironment(exception);
5230 if (clEnv == (MagickCLEnv) NULL)
5231 return((
Image *) NULL);
5233 filteredImage=ComputeWaveletDenoiseImage(image,clEnv,threshold,exception);
5235 return(filteredImage);