44 #include "magick/studio.h"
45 #include "magick/accelerate-private.h"
46 #include "magick/accelerate-kernels-private.h"
47 #include "magick/artifact.h"
48 #include "magick/cache.h"
49 #include "magick/cache-private.h"
50 #include "magick/cache-view.h"
51 #include "magick/color-private.h"
52 #include "magick/delegate-private.h"
53 #include "magick/enhance.h"
54 #include "magick/exception.h"
55 #include "magick/exception-private.h"
56 #include "magick/gem.h"
57 #include "magick/hashmap.h"
58 #include "magick/image.h"
59 #include "magick/image-private.h"
60 #include "magick/list.h"
61 #include "magick/memory_.h"
62 #include "magick/monitor-private.h"
63 #include "magick/opencl.h"
64 #include "magick/opencl-private.h"
65 #include "magick/option.h"
66 #include "magick/pixel-private.h"
67 #include "magick/prepress.h"
68 #include "magick/quantize.h"
69 #include "magick/random_.h"
70 #include "magick/random-private.h"
71 #include "magick/registry.h"
72 #include "magick/resize.h"
73 #include "magick/resize-private.h"
74 #include "magick/semaphore.h"
75 #include "magick/splay-tree.h"
76 #include "magick/statistic.h"
77 #include "magick/string_.h"
78 #include "magick/string-private.h"
79 #include "magick/token.h"
81 #ifdef MAGICKCORE_CLPERFMARKER
82 #include "CLPerfMarker.h"
85 #define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
86 #define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
88 #if defined(MAGICKCORE_OPENCL_SUPPORT)
93 #define ALIGNED(pointer,type) ((((size_t)(pointer)) & (sizeof(type)-1)) == 0)
98 static const ResizeWeightingFunctionType supportedResizeWeighting[] =
100 BoxWeightingFunction,
101 TriangleWeightingFunction,
102 HanningWeightingFunction,
103 HammingWeightingFunction,
104 BlackmanWeightingFunction,
105 CubicBCWeightingFunction,
106 SincWeightingFunction,
107 SincFastWeightingFunction,
108 LastWeightingFunction
114 static Image *ComputeUnsharpMaskImageSingle(
const Image *image,
115 const double radius,
const double sigma,
const double gain,
116 const double threshold,
int blurOnly,
ExceptionInfo *exception);
122 static MagickBooleanType checkAccelerateCondition(
const Image* image,
123 const ChannelType channel)
126 if (image->storage_class != DirectClass)
130 if (image->colorspace != RGBColorspace &&
131 image->colorspace != sRGBColorspace &&
132 image->colorspace != LinearGRAYColorspace &&
133 image->colorspace != GRAYColorspace)
137 if (((channel & RedChannel) == 0) ||
138 ((channel & GreenChannel) == 0) ||
139 ((channel & BlueChannel) == 0))
143 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
144 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
148 if ((image->clip_mask != (
Image *) NULL) || (image->mask != (
Image *) NULL))
154 static MagickBooleanType checkHistogramCondition(
Image *image,
155 const ChannelType channel)
158 if ((channel & SyncChannels) == 0)
161 if (image->intensity == Rec601LuminancePixelIntensityMethod ||
162 image->intensity == Rec709LuminancePixelIntensityMethod)
165 if (image->colorspace != sRGBColorspace)
171 static MagickBooleanType checkOpenCLEnvironment(
ExceptionInfo* exception)
179 clEnv=GetDefaultOpenCLEnv();
181 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,
182 sizeof(MagickBooleanType),&flag,exception);
183 if (flag != MagickFalse)
186 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED,
187 sizeof(MagickBooleanType),&flag,exception);
188 if (flag == MagickFalse)
190 if (InitOpenCLEnv(clEnv,exception) == MagickFalse)
193 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,
194 sizeof(MagickBooleanType),&flag,exception);
195 if (flag != MagickFalse)
204 inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
205 const unsigned int orgGlobalSize,
const unsigned int localGroupSize)
207 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
210 static MagickBooleanType paramMatchesValue(
MagickCLEnv clEnv,
211 MagickOpenCLEnvParam param,
const char *value,
ExceptionInfo *exception)
219 status=GetMagickOpenCLEnvParam(clEnv,param,
sizeof(val),&val,exception);
220 if (status != MagickFalse)
222 status=strcmp(value,val) == 0 ? MagickTrue : MagickFalse;
223 RelinquishMagickMemory(val);
240 static Image *ComputeAddNoiseImage(
const Image *image,
241 const ChannelType channel,
const NoiseType noise_type,
289 **magick_restrict random_info;
297 numRandomNumberPerPixel;
299 #if defined(MAGICKCORE_OPENMP_SUPPORT)
304 outputReady = MagickFalse;
306 filteredImage = NULL;
309 filteredImageBuffer = NULL;
311 addNoiseKernel = NULL;
313 clEnv = GetDefaultOpenCLEnv();
314 context = GetOpenCLContext(clEnv);
315 queue = AcquireOpenCLCommandQueue(clEnv);
317 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
318 if (filteredImage == (
Image *) NULL)
321 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
322 if (imageBuffer == (cl_mem) NULL)
324 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
325 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
328 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
329 if (filteredImageBuffer == (cl_mem) NULL)
331 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
332 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
337 numRandomNumberPerPixel = 0;
339 unsigned int numRandPerChannel = 0;
347 numRandPerChannel = 1;
350 case MultiplicativeGaussianNoise:
352 numRandPerChannel = 2;
356 if ((channel & RedChannel) != 0)
357 numRandomNumberPerPixel+=numRandPerChannel;
358 if ((channel & GreenChannel) != 0)
359 numRandomNumberPerPixel+=numRandPerChannel;
360 if ((channel & BlueChannel) != 0)
361 numRandomNumberPerPixel+=numRandPerChannel;
362 if ((channel & OpacityChannel) != 0)
363 numRandomNumberPerPixel+=numRandPerChannel;
368 option=GetImageArtifact(image,
"attenuate");
369 if (option != (
char *) NULL)
370 attenuate=StringToDouble(option,(
char **) NULL);
371 random_info=AcquireRandomInfoTLS();
372 #if defined(MAGICKCORE_OPENMP_SUPPORT)
373 key=GetRandomSecretKey(random_info[0]);
377 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,
"AddNoise");
380 cl_uint computeUnitCount;
381 cl_uint workItemCount;
382 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(cl_uint), &computeUnitCount, NULL);
383 workItemCount = computeUnitCount * 2 * 256;
384 inputPixelCount = (cl_int) (image->columns * image->rows);
385 pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount;
386 pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4;
388 local_work_size[0] = 256;
389 global_work_size[0] = workItemCount;
393 const unsigned long* s = GetRandomInfoSeed(randomInfo);
395 GetPseudoRandomValue(randomInfo);
397 randomInfo = DestroyRandomInfo(randomInfo);
401 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(cl_mem),(
void *)&imageBuffer);
402 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
403 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(cl_uint),(
void *)&inputPixelCount);
404 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(cl_uint),(
void *)&pixelsPerWorkitem);
405 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(ChannelType),(
void *)&channel);
406 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(NoiseType),(
void *)&noise_type);
408 option=GetImageArtifact(image,
"attenuate");
409 if (option != (
char *) NULL)
410 attenuate=(float)StringToDouble(option,(
char **) NULL);
411 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(
float),(
void *)&attenuate);
412 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(cl_uint),(
void *)&seed0);
413 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(cl_uint),(
void *)&seed1);
414 clEnv->library->clSetKernelArg(addNoiseKernel,k++,
sizeof(
unsigned int),(
void *)&numRandomNumberPerPixel);
416 events=GetOpenCLEvents(image,&event_count);
417 clStatus=clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,NULL,event_count,events,&event);
418 events=(cl_event *) RelinquishMagickMemory(events);
419 if (clStatus != CL_SUCCESS)
421 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
424 if (RecordProfileData(clEnv,AddNoiseKernel,event) == MagickFalse)
426 AddOpenCLEvent(image,event);
427 AddOpenCLEvent(filteredImage,event);
429 clEnv->library->clReleaseEvent(event);
430 outputReady=MagickTrue;
433 OpenCLLogException(__FUNCTION__,__LINE__,exception);
435 if (imageBuffer != (cl_mem) NULL)
436 clEnv->library->clReleaseMemObject(imageBuffer);
437 if (filteredImageBuffer != (cl_mem) NULL)
438 clEnv->library->clReleaseMemObject(filteredImageBuffer);
439 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
440 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
441 if ((outputReady == MagickFalse) && (filteredImage != NULL))
442 filteredImage=(
Image *) DestroyImage(filteredImage);
444 return(filteredImage);
447 MagickPrivate
Image *AccelerateAddNoiseImage(
const Image *image,
448 const ChannelType channel,
const NoiseType noise_type,
467 magick_unreferenced(image);
468 magick_unreferenced(channel);
469 magick_unreferenced(noise_type);
470 magick_unreferenced(exception);
471 return((
Image *)NULL);
486 static Image *ComputeBlurImage(
const Image* image,
const ChannelType channel,
487 const double radius,
const double sigma,
ExceptionInfo *exception)
490 geometry[MaxTextExtent];
545 filteredImage = NULL;
547 tempImageBuffer = NULL;
548 filteredImageBuffer = NULL;
549 imageKernelBuffer = NULL;
550 blurRowKernel = NULL;
551 blurColumnKernel = NULL;
555 outputReady = MagickFalse;
557 clEnv = GetDefaultOpenCLEnv();
558 context = GetOpenCLContext(clEnv);
559 queue = AcquireOpenCLCommandQueue(clEnv);
561 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
562 if (filteredImage == (
Image *) NULL)
565 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
566 if (imageBuffer == (cl_mem) NULL)
568 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
569 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
572 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
573 if (filteredImageBuffer == (cl_mem) NULL)
575 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
576 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
582 (void) FormatLocaleString(geometry,MaxTextExtent,
"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
583 kernel=AcquireKernelInfo(geometry);
586 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"MemoryAllocationFailed.",
".");
591 kernelBufferPtr = (
float *)AcquireMagickMemory(kernel->width *
sizeof(
float));
592 if (kernelBufferPtr == (
float *) NULL)
594 (void)OpenCLThrowMagickException(exception,GetMagickModule(),
595 ResourceLimitWarning,
"AcquireMagickMemory failed.",
"'%s'",
".");
598 for (i = 0; i < kernel->width; i++)
599 kernelBufferPtr[i] = (
float)kernel->values[i];
601 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, kernel->width *
sizeof(
float), kernelBufferPtr, &clStatus);
602 RelinquishMagickMemory(kernelBufferPtr);
603 if (clStatus != CL_SUCCESS)
605 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
615 length = image->columns * image->rows;
616 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 *
sizeof(
float), NULL, &clStatus);
617 if (clStatus != CL_SUCCESS)
619 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
626 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"BlurRow");
627 if (blurRowKernel == NULL)
629 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
633 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"BlurColumn");
634 if (blurColumnKernel == NULL)
636 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
646 imageColumns = (
unsigned int) image->columns;
647 imageRows = (
unsigned int) image->rows;
651 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
652 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
653 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(ChannelType),&channel);
654 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
655 kernelWidth = (
unsigned int) kernel->width;
656 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&kernelWidth);
657 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
658 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
659 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(CLPixelPacket)*(chunkSize+kernel->width),(
void *) NULL);
660 if (clStatus != CL_SUCCESS)
662 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
672 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
673 gsize[1] = image->rows;
674 wsize[0] = chunkSize;
677 events=GetOpenCLEvents(image,&event_count);
678 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, event_count, events, &event);
679 events=(cl_event *) RelinquishMagickMemory(events);
680 if (clStatus != CL_SUCCESS)
682 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
685 if (RecordProfileData(clEnv,BlurRowKernel,event) == MagickFalse)
687 AddOpenCLEvent(image,event);
688 AddOpenCLEvent(filteredImage,event);
690 clEnv->library->clReleaseEvent(event);
699 imageColumns = (
unsigned int) image->columns;
700 imageRows = (
unsigned int) image->rows;
704 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
705 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
706 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(ChannelType),&channel);
707 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
708 kernelWidth = (
unsigned int) kernel->width;
709 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&kernelWidth);
710 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
711 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
712 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_float4)*(chunkSize+kernel->width),(
void *) NULL);
713 if (clStatus != CL_SUCCESS)
715 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
725 gsize[0] = image->columns;
726 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
728 wsize[1] = chunkSize;
730 events=GetOpenCLEvents(image,&event_count);
731 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, event_count, events, &event);
732 events=(cl_event *) RelinquishMagickMemory(events);
733 if (clStatus != CL_SUCCESS)
735 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
738 if (RecordProfileData(clEnv,BlurColumnKernel,event) == MagickFalse)
740 AddOpenCLEvent(image,event);
741 AddOpenCLEvent(filteredImage,event);
743 clEnv->library->clReleaseEvent(event);
749 outputReady=MagickTrue;
752 OpenCLLogException(__FUNCTION__,__LINE__,exception);
754 if (imageBuffer != (cl_mem) NULL)
755 clEnv->library->clReleaseMemObject(imageBuffer);
756 if (filteredImageBuffer != (cl_mem) NULL)
757 clEnv->library->clReleaseMemObject(filteredImageBuffer);
758 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
759 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
760 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
761 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
762 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
763 if (kernel!=NULL) DestroyKernelInfo(kernel);
764 if ((outputReady == MagickFalse) && (filteredImage != NULL))
765 filteredImage=(
Image *) DestroyImage(filteredImage);
766 return(filteredImage);
769 MagickPrivate
Image* AccelerateBlurImage(
const Image *image,
770 const ChannelType channel,
const double radius,
const double sigma,
776 assert(image != NULL);
779 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
780 (checkAccelerateCondition(image, channel) == MagickFalse))
783 filteredImage=ComputeBlurImage(image, channel, radius, sigma, exception);
784 return(filteredImage);
799 static MagickBooleanType LaunchCompositeKernel(
const Image *image,
800 MagickCLEnv clEnv,cl_command_queue queue,cl_mem imageBuffer,
801 const unsigned int inputWidth,
const unsigned int inputHeight,
802 const unsigned int inputMatte,
const ChannelType channel,
803 const CompositeOperator compose,
const cl_mem compositeImageBuffer,
804 const unsigned int compositeWidth,
const unsigned int compositeHeight,
805 const unsigned int compositeMatte,
const float destination_dissolve,
806 const float source_dissolve)
833 compositeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
837 clStatus = clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(cl_mem), (
void*)&imageBuffer);
838 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&inputWidth);
839 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&inputHeight);
840 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&inputMatte);
841 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(cl_mem), (
void*)&compositeImageBuffer);
842 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&compositeWidth);
843 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&compositeHeight);
844 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&compositeMatte);
845 composeOp = (
unsigned int)compose;
846 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
unsigned int), (
void*)&composeOp);
847 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(ChannelType), (
void*)&channel);
848 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
float), (
void*)&destination_dissolve);
849 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++,
sizeof(
float), (
void*)&source_dissolve);
851 if (clStatus != CL_SUCCESS)
854 local_work_size[0] = 64;
855 local_work_size[1] = 1;
857 global_work_size[0] = padGlobalWorkgroupSizeToLocalWorkgroupSize(inputWidth,
858 (
unsigned int)local_work_size[0]);
859 global_work_size[1] = inputHeight;
860 events=GetOpenCLEvents(image,&event_count);
861 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL,
862 global_work_size, local_work_size, event_count, events, &event);
863 events=(cl_event *) RelinquishMagickMemory(events);
864 if (clStatus == CL_SUCCESS)
865 AddOpenCLEvent(image,event);
866 clEnv->library->clReleaseEvent(event);
868 RelinquishOpenCLKernel(clEnv, compositeKernel);
870 return((clStatus == CL_SUCCESS) ? MagickTrue : MagickFalse);
873 static MagickBooleanType ComputeCompositeImage(
Image *image,
874 const ChannelType channel,
const CompositeOperator compose,
875 const Image *compositeImage,
const ssize_t magick_unused(x_offset),
876 const ssize_t magick_unused(y_offset),
const float destination_dissolve,
886 compositeImageBuffer,
896 magick_unreferenced(x_offset);
897 magick_unreferenced(y_offset);
899 status = MagickFalse;
900 outputReady = MagickFalse;
902 compositeImageBuffer = NULL;
904 clEnv = GetDefaultOpenCLEnv();
905 context = GetOpenCLContext(clEnv);
906 queue = AcquireOpenCLCommandQueue(clEnv);
908 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
909 if (imageBuffer == (cl_mem) NULL)
911 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
912 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
916 compositeImageBuffer = GetAuthenticOpenCLBuffer(compositeImage,exception);
917 if (compositeImageBuffer == (cl_mem) NULL)
919 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
920 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
924 status = LaunchCompositeKernel(image,clEnv, queue, imageBuffer,
925 (
unsigned int)image->columns,
926 (
unsigned int)image->rows,
927 (
unsigned int)image->matte,
928 channel, compose, compositeImageBuffer,
929 (
unsigned int)compositeImage->columns,
930 (
unsigned int)compositeImage->rows,
931 (
unsigned int)compositeImage->matte,
932 destination_dissolve, source_dissolve);
934 if (status == MagickFalse)
937 outputReady = MagickTrue;
941 if (imageBuffer != (cl_mem) NULL)
942 clEnv->library->clReleaseMemObject(imageBuffer);
943 if (compositeImageBuffer != (cl_mem) NULL)
944 clEnv->library->clReleaseMemObject(compositeImageBuffer);
946 RelinquishOpenCLCommandQueue(clEnv, queue);
951 MagickPrivate MagickBooleanType AccelerateCompositeImage(
Image *image,
952 const ChannelType channel,
const CompositeOperator compose,
953 const Image *composite,
const ssize_t x_offset,
const ssize_t y_offset,
954 const float destination_dissolve,
const float source_dissolve,
960 assert(image != NULL);
963 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
964 (checkAccelerateCondition(image, channel) == MagickFalse))
971 || image->columns != composite->columns
972 || image->rows != composite->rows)
976 case ColorDodgeCompositeOp:
977 case BlendCompositeOp:
984 status = ComputeCompositeImage(image, channel, compose, composite,
985 x_offset, y_offset, destination_dissolve, source_dissolve, exception);
1002 static MagickBooleanType ComputeContrastImage(
Image *image,
1036 global_work_size[2];
1042 outputReady = MagickFalse;
1046 filterKernel = NULL;
1049 clEnv = GetDefaultOpenCLEnv();
1050 context = GetOpenCLContext(clEnv);
1052 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
1053 if (imageBuffer == (cl_mem) NULL)
1055 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1056 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
1060 filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"Contrast");
1061 if (filterKernel == NULL)
1063 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
1068 clStatus=clEnv->library->clSetKernelArg(filterKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1070 uSharpen = (sharpen == MagickFalse)?0:1;
1071 clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,
sizeof(cl_uint),&uSharpen);
1072 if (clStatus != CL_SUCCESS)
1074 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
1078 global_work_size[0] = image->columns;
1079 global_work_size[1] = image->rows;
1081 queue = AcquireOpenCLCommandQueue(clEnv);
1082 events=GetOpenCLEvents(image,&event_count);
1083 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
1084 events=(cl_event *) RelinquishMagickMemory(events);
1085 if (clStatus != CL_SUCCESS)
1087 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
1090 if (RecordProfileData(clEnv,ContrastKernel,event) == MagickFalse)
1091 AddOpenCLEvent(image,event);
1092 clEnv->library->clReleaseEvent(event);
1093 outputReady=MagickTrue;
1096 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1099 if (imageBuffer != (cl_mem) NULL)
1100 clEnv->library->clReleaseMemObject(imageBuffer);
1101 if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel);
1102 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1103 return(outputReady);
1106 MagickPrivate MagickBooleanType AccelerateContrastImage(
Image *image,
1112 assert(image != NULL);
1115 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1116 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
1117 return(MagickFalse);
1119 status = ComputeContrastImage(image,sharpen,exception);
1135 static MagickBooleanType LaunchHistogramKernel(
MagickCLEnv clEnv,
1136 cl_command_queue queue,cl_mem imageBuffer,cl_mem histogramBuffer,
1163 global_work_size[2];
1165 histogramKernel = NULL;
1167 outputReady = MagickFalse;
1168 method = image->intensity;
1169 colorspace = image->colorspace;
1172 histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"Histogram");
1173 if (histogramKernel == NULL)
1175 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
1181 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1182 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(ChannelType),&channel);
1183 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_int),&method);
1184 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_int),&colorspace);
1185 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,
sizeof(cl_mem),(
void *)&histogramBuffer);
1186 if (clStatus != CL_SUCCESS)
1188 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
1193 global_work_size[0] = image->columns;
1194 global_work_size[1] = image->rows;
1196 events=GetOpenCLEvents(image,&event_count);
1197 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
1198 events=(cl_event *) RelinquishMagickMemory(events);
1200 if (clStatus != CL_SUCCESS)
1202 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
1205 if (RecordProfileData(clEnv,HistogramKernel,event) == MagickFalse)
1206 AddOpenCLEvent(image,event);
1207 clEnv->library->clReleaseEvent(event);
1209 outputReady = MagickTrue;
1212 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1214 if (histogramKernel!=NULL)
1215 RelinquishOpenCLKernel(clEnv, histogramKernel);
1217 return(outputReady);
1220 MagickPrivate MagickBooleanType ComputeContrastStretchImageChannel(
Image *image,
1221 const ChannelType channel,
const double black_point,
const double white_point,
1224 #define ContrastStretchImageTag "ContrastStretch/Image"
1225 #define MaxRange(color) ((MagickRealType) ScaleQuantumToMap((Quantum) (color)))
1280 global_work_size[2];
1285 histogramBuffer = NULL;
1286 stretchMapBuffer = NULL;
1287 histogramKernel = NULL;
1288 stretchKernel = NULL;
1291 outputReady = MagickFalse;
1294 assert(image != (
Image *) NULL);
1295 assert(image->signature == MagickCoreSignature);
1296 if (IsEventLogging() != MagickFalse)
1297 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
1304 clEnv = GetDefaultOpenCLEnv();
1305 context = GetOpenCLContext(clEnv);
1306 queue = AcquireOpenCLCommandQueue(clEnv);
1311 length = (MaxMap+1);
1312 histogram=(cl_uint4 *) AcquireQuantumMemory(length,
sizeof(*histogram));
1314 if (histogram == (cl_uint4 *) NULL)
1315 ThrowBinaryException(ResourceLimitError,
"MemoryAllocationFailed", image->filename);
1318 (void) memset(histogram,0,length*
sizeof(*histogram));
1327 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
1328 if (imageBuffer == (cl_mem) NULL)
1330 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1331 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
1336 histogramBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, length *
sizeof(cl_uint4), histogram, &clStatus);
1337 if (clStatus != CL_SUCCESS)
1339 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1343 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, channel, exception);
1344 if (status == MagickFalse)
1348 events=GetOpenCLEvents(image,&event_count);
1349 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(cl_uint4), event_count, events, NULL, &clStatus);
1350 events=(cl_event *) RelinquishMagickMemory(events);
1351 if (clStatus != CL_SUCCESS)
1353 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
"'%s'",
".");
1358 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
1359 if (clStatus != CL_SUCCESS)
1361 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
"'%s'",
".");
1370 white.z=MaxRange(QuantumRange);
1371 if ((channel & RedChannel) != 0)
1374 for (i=0; i <= (ssize_t) MaxMap; i++)
1376 intensity+=histogram[i].s[2];
1377 if (intensity > black_point)
1380 black.z=(MagickRealType) i;
1382 for (i=(ssize_t) MaxMap; i != 0; i--)
1384 intensity+=histogram[i].s[2];
1385 if (intensity > ((
double) image->columns*image->rows-white_point))
1388 white.z=(MagickRealType) i;
1391 white.y=MaxRange(QuantumRange);
1392 if ((channel & GreenChannel) != 0)
1395 for (i=0; i <= (ssize_t) MaxMap; i++)
1397 intensity+=histogram[i].s[2];
1398 if (intensity > black_point)
1401 black.y=(MagickRealType) i;
1403 for (i=(ssize_t) MaxMap; i != 0; i--)
1405 intensity+=histogram[i].s[2];
1406 if (intensity > ((
double) image->columns*image->rows-white_point))
1409 white.y=(MagickRealType) i;
1412 white.x=MaxRange(QuantumRange);
1413 if ((channel & BlueChannel) != 0)
1416 for (i=0; i <= (ssize_t) MaxMap; i++)
1418 intensity+=histogram[i].s[2];
1419 if (intensity > black_point)
1422 black.x=(MagickRealType) i;
1424 for (i=(ssize_t) MaxMap; i != 0; i--)
1426 intensity+=histogram[i].s[2];
1427 if (intensity > ((
double) image->columns*image->rows-white_point))
1430 white.x=(MagickRealType) i;
1433 white.w=MaxRange(QuantumRange);
1434 if ((channel & OpacityChannel) != 0)
1437 for (i=0; i <= (ssize_t) MaxMap; i++)
1439 intensity+=histogram[i].s[2];
1440 if (intensity > black_point)
1443 black.w=(MagickRealType) i;
1445 for (i=(ssize_t) MaxMap; i != 0; i--)
1447 intensity+=histogram[i].s[2];
1448 if (intensity > ((
double) image->columns*image->rows-white_point))
1451 white.w=(MagickRealType) i;
1478 stretch_map=(
PixelPacket *) AcquireQuantumMemory(length,
1479 sizeof(*stretch_map));
1482 ThrowBinaryException(ResourceLimitError,
"MemoryAllocationFailed",
1488 (void) memset(stretch_map,0,length*
sizeof(*stretch_map));
1489 for (i=0; i <= (ssize_t) MaxMap; i++)
1491 if ((channel & RedChannel) != 0)
1493 if (i < (ssize_t) black.z)
1494 stretch_map[i].red=(Quantum) 0;
1496 if (i > (ssize_t) white.z)
1497 stretch_map[i].red=QuantumRange;
1499 if (black.z != white.z)
1500 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap*
1501 (i-black.z)/(white.z-black.z)));
1503 if ((channel & GreenChannel) != 0)
1505 if (i < (ssize_t) black.y)
1506 stretch_map[i].green=0;
1508 if (i > (ssize_t) white.y)
1509 stretch_map[i].green=QuantumRange;
1511 if (black.y != white.y)
1512 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap*
1513 (i-black.y)/(white.y-black.y)));
1515 if ((channel & BlueChannel) != 0)
1517 if (i < (ssize_t) black.x)
1518 stretch_map[i].blue=0;
1520 if (i > (ssize_t) white.x)
1521 stretch_map[i].blue= QuantumRange;
1523 if (black.x != white.x)
1524 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap*
1525 (i-black.x)/(white.x-black.x)));
1527 if ((channel & OpacityChannel) != 0)
1529 if (i < (ssize_t) black.w)
1530 stretch_map[i].opacity=0;
1532 if (i > (ssize_t) white.w)
1533 stretch_map[i].opacity=QuantumRange;
1535 if (black.w != white.w)
1536 stretch_map[i].opacity=ScaleMapToQuantum((MagickRealType) (MaxMap*
1537 (i-black.w)/(white.w-black.w)));
1559 if (((channel & OpacityChannel) != 0) || (((channel & IndexChannel) != 0) &&
1560 (image->colorspace == CMYKColorspace)))
1561 image->storage_class=DirectClass;
1562 if (image->storage_class == PseudoClass)
1567 for (i=0; i < (ssize_t) image->colors; i++)
1569 if ((channel & RedChannel) != 0)
1571 if (black.z != white.z)
1572 image->colormap[i].red=stretch_map[
1573 ScaleQuantumToMap(image->colormap[i].red)].red;
1575 if ((channel & GreenChannel) != 0)
1577 if (black.y != white.y)
1578 image->colormap[i].green=stretch_map[
1579 ScaleQuantumToMap(image->colormap[i].green)].green;
1581 if ((channel & BlueChannel) != 0)
1583 if (black.x != white.x)
1584 image->colormap[i].blue=stretch_map[
1585 ScaleQuantumToMap(image->colormap[i].blue)].blue;
1587 if ((channel & OpacityChannel) != 0)
1589 if (black.w != white.w)
1590 image->colormap[i].opacity=stretch_map[
1591 ScaleQuantumToMap(image->colormap[i].opacity)].opacity;
1598 stretchMapBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, length, stretch_map, &clStatus);
1599 if (clStatus != CL_SUCCESS)
1601 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1606 stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"ContrastStretch");
1607 if (stretchKernel == NULL)
1609 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
1615 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1616 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(ChannelType),&channel);
1617 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_mem),(
void *)&stretchMapBuffer);
1618 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_float4),&white);
1619 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,
sizeof(cl_float4),&black);
1620 if (clStatus != CL_SUCCESS)
1622 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
1627 global_work_size[0] = image->columns;
1628 global_work_size[1] = image->rows;
1630 events=GetOpenCLEvents(image,&event_count);
1631 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
1632 events=(cl_event *) RelinquishMagickMemory(events);
1634 if (clStatus != CL_SUCCESS)
1636 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
1640 if (RecordProfileData(clEnv,ContrastStretchKernel,event) == MagickFalse)
1641 AddOpenCLEvent(image, event);
1642 clEnv->library->clReleaseEvent(event);
1644 outputReady=MagickTrue;
1647 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1649 if (imageBuffer != (cl_mem) NULL)
1650 clEnv->library->clReleaseMemObject(imageBuffer);
1652 if (stretchMapBuffer!=NULL)
1653 clEnv->library->clReleaseMemObject(stretchMapBuffer);
1654 if (stretch_map!=NULL)
1655 stretch_map=(
PixelPacket *) RelinquishMagickMemory(stretch_map);
1658 if (histogramBuffer!=NULL)
1659 clEnv->library->clReleaseMemObject(histogramBuffer);
1660 if (histogram!=NULL)
1661 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
1664 if (histogramKernel!=NULL)
1665 RelinquishOpenCLKernel(clEnv, histogramKernel);
1666 if (stretchKernel!=NULL)
1667 RelinquishOpenCLKernel(clEnv, stretchKernel);
1670 RelinquishOpenCLCommandQueue(clEnv, queue);
1672 return(outputReady);
1675 MagickPrivate MagickBooleanType AccelerateContrastStretchImageChannel(
1676 Image *image,
const ChannelType channel,
const double black_point,
1682 assert(image != NULL);
1685 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1686 (checkAccelerateCondition(image, channel) == MagickFalse) ||
1687 (checkHistogramCondition(image, channel) == MagickFalse))
1688 return(MagickFalse);
1690 status=ComputeContrastStretchImageChannel(image,channel, black_point, white_point, exception);
1706 static Image *ComputeConvolveImage(
const Image* image,
1726 filteredImageBuffer,
1733 deviceLocalMemorySize;
1751 global_work_size[3],
1753 localMemoryRequirement;
1769 filteredImageBuffer = NULL;
1770 convolutionKernel = NULL;
1774 filteredImage = NULL;
1775 outputReady = MagickFalse;
1777 clEnv = GetDefaultOpenCLEnv();
1779 context = GetOpenCLContext(clEnv);
1781 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
1782 if (filteredImage == (
Image *) NULL)
1785 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
1786 if (imageBuffer == (cl_mem) NULL)
1788 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1789 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
1792 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,exception);
1793 if (filteredImageBuffer == (cl_mem) NULL)
1795 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1796 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
1800 kernelSize = (
unsigned int) (kernel->width * kernel->height);
1801 convolutionKernel = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize *
sizeof(
float), NULL, &clStatus);
1802 if (clStatus != CL_SUCCESS)
1804 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
1808 queue = AcquireOpenCLCommandQueue(clEnv);
1810 kernelBufferPtr = (
float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize *
sizeof(
float)
1811 , 0, NULL, NULL, &clStatus);
1812 if (clStatus != CL_SUCCESS)
1814 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
1817 for (i = 0; i < kernelSize; i++)
1819 kernelBufferPtr[i] = (float) kernel->values[i];
1821 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
1822 if (clStatus != CL_SUCCESS)
1824 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
"'%s'",
".");
1828 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
1832 localGroupSize[0] = 16;
1833 localGroupSize[1] = 16;
1834 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) *
sizeof(CLPixelPacket)
1835 + kernel->width*kernel->height*
sizeof(
float);
1837 if (localMemoryRequirement > deviceLocalMemorySize)
1839 localGroupSize[0] = 8;
1840 localGroupSize[1] = 8;
1841 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) *
sizeof(CLPixelPacket)
1842 + kernel->width*kernel->height*
sizeof(
float);
1844 if (localMemoryRequirement <= deviceLocalMemorySize)
1847 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"ConvolveOptimized");
1848 if (clkernel == NULL)
1850 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
1856 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1857 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
1858 imageWidth = (
unsigned int) image->columns;
1859 imageHeight = (
unsigned int) image->rows;
1860 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageWidth);
1861 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageHeight);
1862 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&convolutionKernel);
1863 filterWidth = (
unsigned int) kernel->width;
1864 filterHeight = (
unsigned int) kernel->height;
1865 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterWidth);
1866 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterHeight);
1867 matte = (image->matte==MagickTrue)?1:0;
1868 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&matte);
1869 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(ChannelType),(
void *)&channel);
1870 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*
sizeof(CLPixelPacket),NULL);
1871 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*
sizeof(
float),NULL);
1872 if (clStatus != CL_SUCCESS)
1874 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
1879 global_work_size[0] = ((image->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ;
1880 global_work_size[1] = ((image->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1];
1883 events = GetOpenCLEvents(image, &event_count);
1884 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, event_count, events, &event);
1885 events=(cl_event *) RelinquishMagickMemory(events);
1886 if (clStatus != CL_SUCCESS)
1888 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
1891 if (RecordProfileData(clEnv,ConvolveKernel,event) == MagickFalse)
1893 AddOpenCLEvent(image, event);
1894 AddOpenCLEvent(filteredImage, event);
1896 clEnv->library->clReleaseEvent(event);
1901 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"Convolve");
1902 if (clkernel == NULL)
1904 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
1910 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
1911 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
1912 imageWidth = (
unsigned int) image->columns;
1913 imageHeight = (
unsigned int) image->rows;
1914 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageWidth);
1915 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&imageHeight);
1916 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&convolutionKernel);
1917 filterWidth = (
unsigned int) kernel->width;
1918 filterHeight = (
unsigned int) kernel->height;
1919 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterWidth);
1920 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&filterHeight);
1921 matte = (image->matte==MagickTrue)?1:0;
1922 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&matte);
1923 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(ChannelType),(
void *)&channel);
1924 if (clStatus != CL_SUCCESS)
1926 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
1930 localGroupSize[0] = 8;
1931 localGroupSize[1] = 8;
1932 global_work_size[0] = (image->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0];
1933 global_work_size[1] = (image->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1];
1934 events=GetOpenCLEvents(image,&event_count);
1935 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, event_count, events, &event);
1936 events=(cl_event *) RelinquishMagickMemory(events);
1938 if (clStatus != CL_SUCCESS)
1940 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
1943 if (RecordProfileData(clEnv,ConvolveKernel,event) == MagickFalse)
1945 AddOpenCLEvent(image,event);
1946 AddOpenCLEvent(filteredImage,event);
1948 clEnv->library->clReleaseEvent(event);
1951 outputReady = MagickTrue;
1954 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1956 if (imageBuffer != (cl_mem) NULL)
1957 clEnv->library->clReleaseMemObject(imageBuffer);
1959 if (filteredImageBuffer != (cl_mem) NULL)
1960 clEnv->library->clReleaseMemObject(filteredImageBuffer);
1962 if (convolutionKernel != NULL)
1963 clEnv->library->clReleaseMemObject(convolutionKernel);
1965 if (clkernel != NULL)
1966 RelinquishOpenCLKernel(clEnv, clkernel);
1969 RelinquishOpenCLCommandQueue(clEnv, queue);
1971 if ((outputReady == MagickFalse) && (filteredImage != NULL))
1972 filteredImage=(
Image *) DestroyImage(filteredImage);
1974 return(filteredImage);
1977 MagickPrivate
Image *AccelerateConvolveImageChannel(
const Image *image,
1983 assert(image != NULL);
1987 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1988 (checkAccelerateCondition(image, channel) == MagickFalse))
1991 filteredImage=ComputeConvolveImage(image, channel, kernel, exception);
1992 return(filteredImage);
2007 static Image *ComputeDespeckleImage(
const Image *image,
2011 X[4] = {0, 1, 1,-1},
2012 Y[4] = {1, 0, 1, 1};
2031 filteredImageBuffer,
2055 global_work_size[2];
2061 outputReady = MagickFalse;
2063 filteredImage = NULL;
2066 filteredImageBuffer = NULL;
2070 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
2071 clEnv = GetDefaultOpenCLEnv();
2072 context = GetOpenCLContext(clEnv);
2073 queue = AcquireOpenCLCommandQueue(clEnv);
2076 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
2077 if (filteredImage == (
Image *) NULL)
2080 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
2081 if (imageBuffer == (cl_mem) NULL)
2083 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2084 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
2087 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage, exception);
2088 if (filteredImageBuffer == (cl_mem) NULL)
2090 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2091 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
2095 hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"HullPass1");
2096 hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"HullPass2");
2098 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,
sizeof(cl_mem),(
void *)&imageBuffer);
2099 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,
sizeof(cl_mem),(
void *)(tempImageBuffer+1));
2100 imageWidth = (
unsigned int) image->columns;
2101 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,
sizeof(
unsigned int),(
void *)&imageWidth);
2102 imageHeight = (
unsigned int) image->rows;
2103 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,
sizeof(
unsigned int),(
void *)&imageHeight);
2104 matte = (image->matte==MagickFalse)?0:1;
2105 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,
sizeof(
int),(
void *)&matte);
2106 if (clStatus != CL_SUCCESS)
2108 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2112 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,
sizeof(cl_mem),(
void *)(tempImageBuffer+1));
2113 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,
sizeof(cl_mem),(
void *)tempImageBuffer);
2114 imageWidth = (
unsigned int) image->columns;
2115 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,
sizeof(
unsigned int),(
void *)&imageWidth);
2116 imageHeight = (
unsigned int) image->rows;
2117 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,
sizeof(
unsigned int),(
void *)&imageHeight);
2118 matte = (image->matte==MagickFalse)?0:1;
2119 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,
sizeof(
int),(
void *)&matte);
2120 if (clStatus != CL_SUCCESS)
2122 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2127 global_work_size[0] = image->columns;
2128 global_work_size[1] = image->rows;
2130 events=GetOpenCLEvents(image,&event_count);
2131 for (k = 0; k < 4; k++)
2140 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2141 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2142 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2143 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2144 if (clStatus != CL_SUCCESS)
2146 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2150 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2151 if (clStatus != CL_SUCCESS)
2153 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2156 RecordProfileData(clEnv,HullPass1Kernel,event);
2157 clEnv->library->clReleaseEvent(event);
2159 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2160 if (clStatus != CL_SUCCESS)
2162 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2165 RecordProfileData(clEnv,HullPass2Kernel,event);
2166 clEnv->library->clReleaseEvent(event);
2170 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,
sizeof(cl_mem),(
void *)(tempImageBuffer));
2171 offset.s[0] = -X[k];
2172 offset.s[1] = -Y[k];
2174 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2175 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2176 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2177 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2178 if (clStatus != CL_SUCCESS)
2180 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2184 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2185 if (clStatus != CL_SUCCESS)
2187 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2190 RecordProfileData(clEnv,HullPass1Kernel,event);
2191 clEnv->library->clReleaseEvent(event);
2193 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2194 if (clStatus != CL_SUCCESS)
2196 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2199 RecordProfileData(clEnv,HullPass2Kernel,event);
2200 clEnv->library->clReleaseEvent(event);
2202 offset.s[0] = -X[k];
2203 offset.s[1] = -Y[k];
2205 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2206 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2207 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2208 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2209 if (clStatus != CL_SUCCESS)
2211 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2215 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2216 if (clStatus != CL_SUCCESS)
2218 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2221 RecordProfileData(clEnv,HullPass1Kernel,event);
2222 clEnv->library->clReleaseEvent(event);
2224 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2225 if (clStatus != CL_SUCCESS)
2227 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2230 RecordProfileData(clEnv,HullPass2Kernel,event);
2231 clEnv->library->clReleaseEvent(event);
2236 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,
sizeof(cl_int2),(
void *)&offset);
2237 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,
sizeof(
int),(
void *)&polarity);
2238 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,
sizeof(cl_int2),(
void *)&offset);
2239 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,
sizeof(
int),(
void *)&polarity);
2242 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
2244 if (clStatus != CL_SUCCESS)
2246 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2250 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2251 if (clStatus != CL_SUCCESS)
2253 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2256 RecordProfileData(clEnv,HullPass1Kernel,event);
2257 clEnv->library->clReleaseEvent(event);
2259 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2260 if (clStatus != CL_SUCCESS)
2262 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2265 if ((k == 3) && (RecordProfileData(clEnv,HullPass2Kernel,event) == MagickFalse))
2267 AddOpenCLEvent(image,event);
2268 AddOpenCLEvent(filteredImage,event);
2270 clEnv->library->clReleaseEvent(event);
2273 outputReady=MagickTrue;
2276 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2278 if (imageBuffer != (cl_mem) NULL)
2279 clEnv->library->clReleaseMemObject(imageBuffer);
2280 if (filteredImageBuffer != (cl_mem) NULL)
2281 clEnv->library->clReleaseMemObject(filteredImageBuffer);
2282 events=(cl_event *) RelinquishMagickMemory(events);
2283 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2284 for (k = 0; k < 2; k++)
2286 if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
2288 if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
2289 if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
2290 if ((outputReady == MagickFalse) && (filteredImage != NULL))
2291 filteredImage=(
Image *) DestroyImage(filteredImage);
2292 return(filteredImage);
2295 MagickPrivate
Image *AccelerateDespeckleImage(
const Image* image,
2301 assert(image != NULL);
2304 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2305 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
2308 filteredImage=ComputeDespeckleImage(image,exception);
2309 return(filteredImage);
2324 MagickPrivate MagickBooleanType ComputeEqualizeImage(
Image *image,
2327 #define EqualizeImageTag "Equalize/Image"
2382 global_work_size[2];
2388 histogramBuffer = NULL;
2389 equalizeMapBuffer = NULL;
2390 histogramKernel = NULL;
2391 equalizeKernel = NULL;
2394 outputReady = MagickFalse;
2396 assert(image != (
Image *) NULL);
2397 assert(image->signature == MagickCoreSignature);
2398 if (IsEventLogging() != MagickFalse)
2399 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2404 clEnv = GetDefaultOpenCLEnv();
2405 context = GetOpenCLContext(clEnv);
2406 queue = AcquireOpenCLCommandQueue(clEnv);
2412 histogram=(cl_uint4 *) AcquireQuantumMemory(length,
sizeof(*histogram));
2413 if (histogram == (cl_uint4 *) NULL)
2414 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
2417 (void) memset(histogram,0,length*
sizeof(*histogram));
2419 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
2420 if (imageBuffer == (cl_mem) NULL)
2422 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2423 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
2428 histogramBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, length *
sizeof(cl_uint4), histogram, &clStatus);
2429 if (clStatus != CL_SUCCESS)
2431 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2435 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, channel, exception);
2436 if (status == MagickFalse)
2440 events=GetOpenCLEvents(image,&event_count);
2441 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length *
sizeof(cl_uint4), event_count, events, NULL, &clStatus);
2442 events=(cl_event *) RelinquishMagickMemory(events);
2443 if (clStatus != CL_SUCCESS)
2445 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"Reading output image from CL buffer failed.",
"'%s'",
".");
2450 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
2451 if (clStatus != CL_SUCCESS)
2453 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
"'%s'",
".");
2458 equalize_map=(
PixelPacket *) AcquireQuantumMemory(length,
sizeof(*equalize_map));
2460 ThrowBinaryException(ResourceLimitWarning,
"MemoryAllocationFailed", image->filename);
2462 map=(cl_float4 *) AcquireQuantumMemory(length,
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 ((channel & SyncChannels) != 0)
2474 intensity.z+=histogram[i].s[2];
2478 if ((channel & RedChannel) != 0)
2479 intensity.z+=histogram[i].s[2];
2480 if ((channel & GreenChannel) != 0)
2481 intensity.y+=histogram[i].s[1];
2482 if ((channel & BlueChannel) != 0)
2483 intensity.x+=histogram[i].s[0];
2484 if ((channel & OpacityChannel) != 0)
2485 intensity.w+=histogram[i].s[3];
2496 white=map[(int) MaxMap];
2497 (void) memset(equalize_map,0,length*
sizeof(*equalize_map));
2498 for (i=0; i <= (ssize_t) MaxMap; i++)
2500 if ((channel & SyncChannels) != 0)
2502 if (white.z != black.z)
2503 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2504 (map[i].z-black.z))/(white.z-black.z)));
2507 if (((channel & RedChannel) != 0) && (white.z != black.z))
2508 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2509 (map[i].z-black.z))/(white.z-black.z)));
2510 if (((channel & GreenChannel) != 0) && (white.y != black.y))
2511 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2512 (map[i].y-black.y))/(white.y-black.y)));
2513 if (((channel & BlueChannel) != 0) && (white.x != black.x))
2514 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2515 (map[i].x-black.x))/(white.x-black.x)));
2516 if (((channel & OpacityChannel) != 0) && (white.w != black.w))
2517 equalize_map[i].opacity=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2518 (map[i].w-black.w))/(white.w-black.w)));
2528 if (image->storage_class == PseudoClass)
2533 for (i=0; i < (ssize_t) image->colors; i++)
2535 if ((channel & SyncChannels) != 0)
2537 if (white.z != black.z)
2539 image->colormap[i].red=equalize_map[
2540 ScaleQuantumToMap(image->colormap[i].red)].red;
2541 image->colormap[i].green=equalize_map[
2542 ScaleQuantumToMap(image->colormap[i].green)].red;
2543 image->colormap[i].blue=equalize_map[
2544 ScaleQuantumToMap(image->colormap[i].blue)].red;
2545 image->colormap[i].opacity=equalize_map[
2546 ScaleQuantumToMap(image->colormap[i].opacity)].red;
2550 if (((channel & RedChannel) != 0) && (white.z != black.z))
2551 image->colormap[i].red=equalize_map[
2552 ScaleQuantumToMap(image->colormap[i].red)].red;
2553 if (((channel & GreenChannel) != 0) && (white.y != black.y))
2554 image->colormap[i].green=equalize_map[
2555 ScaleQuantumToMap(image->colormap[i].green)].green;
2556 if (((channel & BlueChannel) != 0) && (white.x != black.x))
2557 image->colormap[i].blue=equalize_map[
2558 ScaleQuantumToMap(image->colormap[i].blue)].blue;
2559 if (((channel & OpacityChannel) != 0) &&
2560 (white.w != black.w))
2561 image->colormap[i].opacity=equalize_map[
2562 ScaleQuantumToMap(image->colormap[i].opacity)].opacity;
2567 equalizeMapBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, length *
sizeof(
PixelPacket), equalize_map, &clStatus);
2568 if (clStatus != CL_SUCCESS)
2570 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
2575 equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"Equalize");
2576 if (equalizeKernel == NULL)
2578 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
2584 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2585 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(ChannelType),&channel);
2586 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_mem),(
void *)&equalizeMapBuffer);
2587 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_float4),&white);
2588 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,
sizeof(cl_float4),&black);
2589 if (clStatus != CL_SUCCESS)
2591 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2596 global_work_size[0] = image->columns;
2597 global_work_size[1] = image->rows;
2599 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2601 if (clStatus != CL_SUCCESS)
2603 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2606 if (RecordProfileData(clEnv,EqualizeKernel,event) == MagickFalse)
2607 AddOpenCLEvent(image,event);
2608 clEnv->library->clReleaseEvent(event);
2611 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2613 if (imageBuffer != (cl_mem) NULL)
2614 clEnv->library->clReleaseMemObject(imageBuffer);
2617 map=(cl_float4 *) RelinquishMagickMemory(map);
2619 if (equalizeMapBuffer!=NULL)
2620 clEnv->library->clReleaseMemObject(equalizeMapBuffer);
2621 if (equalize_map!=NULL)
2622 equalize_map=(
PixelPacket *) RelinquishMagickMemory(equalize_map);
2624 if (histogramBuffer!=NULL)
2625 clEnv->library->clReleaseMemObject(histogramBuffer);
2626 if (histogram!=NULL)
2627 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
2629 if (histogramKernel!=NULL)
2630 RelinquishOpenCLKernel(clEnv, histogramKernel);
2631 if (equalizeKernel!=NULL)
2632 RelinquishOpenCLKernel(clEnv, equalizeKernel);
2635 RelinquishOpenCLCommandQueue(clEnv, queue);
2637 return(outputReady);
2640 MagickPrivate MagickBooleanType AccelerateEqualizeImage(
Image *image,
2646 assert(image != NULL);
2649 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2650 (checkAccelerateCondition(image, channel) == MagickFalse) ||
2651 (checkHistogramCondition(image, channel) == MagickFalse))
2652 return(MagickFalse);
2654 status=ComputeEqualizeImage(image,channel,exception);
2670 static MagickBooleanType ComputeFunctionImage(
Image *image,
2671 const ChannelType channel,
const MagickFunction
function,
2672 const size_t number_parameters,
const double *parameters,
2698 *parametersBufferPtr;
2713 status = MagickFalse;
2719 parametersBuffer = NULL;
2721 clEnv = GetDefaultOpenCLEnv();
2722 context = GetOpenCLContext(clEnv);
2724 queue = AcquireOpenCLCommandQueue(clEnv);
2726 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
2727 if (imageBuffer == (cl_mem) NULL)
2729 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2730 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
2736 parametersBufferPtr = (
float*)AcquireMagickMemory(number_parameters *
sizeof(
float));
2738 for (i = 0; i < number_parameters; i++)
2739 parametersBufferPtr[i] = (
float)parameters[i];
2741 parametersBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, number_parameters *
sizeof(
float), parametersBufferPtr, &clStatus);
2742 parametersBufferPtr=RelinquishMagickMemory(parametersBufferPtr);
2745 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"ComputeFunction");
2746 if (clkernel == NULL)
2748 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
2754 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2755 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(ChannelType),(
void *)&channel);
2756 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(MagickFunction),(
void *)&
function);
2757 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(
unsigned int),(
void *)&number_parameters);
2758 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,
sizeof(cl_mem),(
void *)¶metersBuffer);
2759 if (clStatus != CL_SUCCESS)
2761 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2765 globalWorkSize[0] = image->columns;
2766 globalWorkSize[1] = image->rows;
2768 events=GetOpenCLEvents(image,&event_count);
2769 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, event_count, events, &event);
2770 events=(cl_event *) RelinquishMagickMemory(events);
2771 if (clStatus != CL_SUCCESS)
2773 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2776 if (RecordProfileData(clEnv,ComputeFunctionKernel,event) == MagickFalse)
2777 AddOpenCLEvent(image,event);
2778 clEnv->library->clReleaseEvent(event);
2779 status = MagickTrue;
2782 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2784 if (imageBuffer != (cl_mem) NULL)
2785 clEnv->library->clReleaseMemObject(imageBuffer);
2786 if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel);
2787 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2788 if (parametersBuffer != NULL) clEnv->library->clReleaseMemObject(parametersBuffer);
2793 MagickPrivate MagickBooleanType AccelerateFunctionImage(
Image *image,
2794 const ChannelType channel,
const MagickFunction
function,
2795 const size_t number_parameters,
const double *parameters,
2801 assert(image != NULL);
2804 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2805 (checkAccelerateCondition(image, channel) == MagickFalse))
2806 return(MagickFalse);
2808 status=ComputeFunctionImage(image, channel,
function, number_parameters, parameters, exception);
2824 MagickBooleanType ComputeGrayscaleImage(
Image *image,
2865 grayscaleKernel = NULL;
2867 assert(image != (
Image *) NULL);
2868 assert(image->signature == MagickCoreSignature);
2869 if (IsEventLogging() != MagickFalse)
2870 (void) LogMagickEvent(TraceEvent,GetMagickModule(),
"%s",image->filename);
2875 clEnv = GetDefaultOpenCLEnv();
2876 context = GetOpenCLContext(clEnv);
2877 queue = AcquireOpenCLCommandQueue(clEnv);
2879 outputReady = MagickFalse;
2881 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
2882 if (imageBuffer == (cl_mem) NULL)
2884 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2885 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
2889 intensityMethod = method;
2890 colorspace = image->colorspace;
2892 grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"Grayscale");
2893 if (grayscaleKernel == NULL)
2895 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
2900 clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
2901 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,
sizeof(cl_int),&intensityMethod);
2902 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,
sizeof(cl_int),&colorspace);
2903 if (clStatus != CL_SUCCESS)
2905 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
2906 printf(
"no kernel\n");
2911 size_t global_work_size[2];
2912 global_work_size[0] = image->columns;
2913 global_work_size[1] = image->rows;
2915 events=GetOpenCLEvents(image,&event_count);
2916 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, grayscaleKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
2917 events=(cl_event *) RelinquishMagickMemory(events);
2918 if (clStatus != CL_SUCCESS)
2920 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
2923 if (RecordProfileData(clEnv,GrayScaleKernel,event) == MagickFalse)
2924 AddOpenCLEvent(image,event);
2925 clEnv->library->clReleaseEvent(event);
2928 outputReady=MagickTrue;
2931 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2933 if (imageBuffer != (cl_mem) NULL)
2934 clEnv->library->clReleaseMemObject(imageBuffer);
2935 if (grayscaleKernel!=NULL)
2936 RelinquishOpenCLKernel(clEnv, grayscaleKernel);
2938 RelinquishOpenCLCommandQueue(clEnv, queue);
2940 return(outputReady);
2943 MagickPrivate MagickBooleanType AccelerateGrayscaleImage(
Image* image,
2949 assert(image != NULL);
2952 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2953 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
2954 return(MagickFalse);
2956 if (method == Rec601LuminancePixelIntensityMethod || method == Rec709LuminancePixelIntensityMethod)
2957 return(MagickFalse);
2959 if (image->colorspace != sRGBColorspace)
2960 return(MagickFalse);
2962 status=ComputeGrayscaleImage(image,method,exception);
2978 static Image *ComputeLocalContrastImage(
const Image *image,
2979 const double radius,
const double strength,
ExceptionInfo *exception)
2999 filteredImageBuffer,
3026 filteredImage = NULL;
3029 filteredImageBuffer = NULL;
3030 tempImageBuffer = NULL;
3031 blurRowKernel = NULL;
3032 blurColumnKernel = NULL;
3034 outputReady = MagickFalse;
3036 clEnv = GetDefaultOpenCLEnv();
3037 context = GetOpenCLContext(clEnv);
3038 queue = AcquireOpenCLCommandQueue(clEnv);
3040 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3041 if (filteredImage == (
Image *) NULL)
3044 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
3045 if (imageBuffer == (cl_mem) NULL)
3047 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3048 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3051 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
3052 if (filteredImageBuffer == (cl_mem) NULL)
3054 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3055 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3062 length = image->columns * image->rows;
3063 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length *
sizeof(
float), NULL, &clStatus);
3064 if (clStatus != CL_SUCCESS)
3066 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3073 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"LocalContrastBlurRow");
3074 if (blurRowKernel == NULL)
3076 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
3080 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"LocalContrastBlurApplyColumn");
3081 if (blurColumnKernel == NULL)
3083 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
3089 imageColumns = (
unsigned int) image->columns;
3090 imageRows = (
unsigned int) image->rows;
3091 iRadius = (cl_int) (image->rows > image->columns ? image->rows : image->columns) * 0.002f * fabs(radius);
3093 passes = (((1.0f * imageRows) * imageColumns * iRadius) + 3999999999) / 4000000000.0f;
3094 passes = (passes < 1) ? 1: passes;
3098 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
3099 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
3100 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
3101 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_int),(
void *)&iRadius);
3102 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
3103 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
3105 if (clStatus != CL_SUCCESS)
3107 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
3115 for (x = 0; x < passes; ++x) {
3121 gsize[1] = (image->rows + passes - 1) / passes;
3125 goffset[1] = x * gsize[1];
3127 events=GetOpenCLEvents(image,&event_count);
3128 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, event_count, events, &event);
3129 events=(cl_event *) RelinquishMagickMemory(events);
3130 if (clStatus != CL_SUCCESS)
3132 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
3135 clEnv->library->clFlush(queue);
3136 if (RecordProfileData(clEnv,LocalContrastBlurRowKernel,event) == MagickFalse)
3138 AddOpenCLEvent(image,event);
3139 AddOpenCLEvent(filteredImage, event);
3141 clEnv->library->clReleaseEvent(event);
3146 cl_float FStrength = strength;
3148 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
3149 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
3150 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
3151 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&iRadius);
3152 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(cl_float),(
void *)&FStrength);
3153 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
3154 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
3156 if (clStatus != CL_SUCCESS)
3158 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
3166 for (x = 0; x < passes; ++x) {
3171 gsize[0] = ((image->columns + 3) / 4) * 4;
3172 gsize[1] = ((((image->rows + 63) / 64) + (passes + 1)) / passes) * 64;
3176 goffset[1] = x * gsize[1];
3178 events=GetOpenCLEvents(image,&event_count);
3179 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, event_count, events, &event);
3180 events=(cl_event *) RelinquishMagickMemory(events);
3181 if (clStatus != CL_SUCCESS)
3183 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
3186 clEnv->library->clFlush(queue);
3187 if (RecordProfileData(clEnv, LocalContrastBlurApplyColumnKernel, event) == MagickFalse)
3189 AddOpenCLEvent(image,event);
3190 AddOpenCLEvent(filteredImage,event);
3192 clEnv->library->clReleaseEvent(event);
3197 outputReady = MagickTrue;
3201 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3203 if (imageBuffer != (cl_mem) NULL)
3204 clEnv->library->clReleaseMemObject(imageBuffer);
3205 if (filteredImageBuffer != (cl_mem) NULL)
3206 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3207 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
3208 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
3209 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
3210 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3211 if ((outputReady == MagickFalse) && (filteredImage != NULL))
3212 filteredImage=(
Image *) DestroyImage(filteredImage);
3213 return(filteredImage);
3216 MagickPrivate
Image *AccelerateLocalContrastImage(
const Image *image,
3217 const double radius,
const double strength,
ExceptionInfo *exception)
3222 assert(image != NULL);
3225 if ((checkOpenCLEnvironment(exception) == MagickFalse))
3228 filteredImage=ComputeLocalContrastImage(image,radius,strength,exception);
3230 return(filteredImage);
3245 MagickBooleanType ComputeModulateImage(
Image *image,
3246 double percent_brightness,
double percent_hue,
double percent_saturation,
3289 modulateKernel = NULL;
3292 assert(image != (
Image *)NULL);
3293 assert(image->signature == MagickCoreSignature);
3294 if (IsEventLogging() != MagickFalse)
3295 (void) LogMagickEvent(TraceEvent, GetMagickModule(),
"%s", image->filename);
3300 clEnv = GetDefaultOpenCLEnv();
3301 context = GetOpenCLContext(clEnv);
3302 queue = AcquireOpenCLCommandQueue(clEnv);
3304 outputReady = MagickFalse;
3306 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
3307 if (imageBuffer == (cl_mem) NULL)
3309 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3310 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3314 modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"Modulate");
3315 if (modulateKernel == NULL)
3317 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
3321 bright = percent_brightness;
3323 saturation = percent_saturation;
3327 clStatus = clEnv->library->clSetKernelArg(modulateKernel, i++,
sizeof(cl_mem), (
void *)&imageBuffer);
3328 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++,
sizeof(cl_float), &bright);
3329 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++,
sizeof(cl_float), &hue);
3330 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++,
sizeof(cl_float), &saturation);
3331 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++,
sizeof(cl_float), &color);
3332 if (clStatus != CL_SUCCESS)
3334 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
3335 printf(
"no kernel\n");
3340 size_t global_work_size[2];
3341 global_work_size[0] = image->columns;
3342 global_work_size[1] = image->rows;
3344 events=GetOpenCLEvents(image,&event_count);
3345 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
3346 events=(cl_event *) RelinquishMagickMemory(events);
3347 if (clStatus != CL_SUCCESS)
3349 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
3352 if (RecordProfileData(clEnv, ModulateKernel, event) == MagickFalse)
3353 AddOpenCLEvent(image,event);
3354 clEnv->library->clReleaseEvent(event);
3357 outputReady=MagickTrue;
3360 OpenCLLogException(__FUNCTION__, __LINE__, exception);
3362 if (imageBuffer != (cl_mem) NULL)
3363 clEnv->library->clReleaseMemObject(imageBuffer);
3364 if (modulateKernel != NULL)
3365 RelinquishOpenCLKernel(clEnv, modulateKernel);
3367 RelinquishOpenCLCommandQueue(clEnv, queue);
3369 return(outputReady);
3372 MagickPrivate MagickBooleanType AccelerateModulateImage(
Image *image,
3373 double percent_brightness,
double percent_hue,
double percent_saturation,
3379 assert(image != NULL);
3382 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3383 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
3384 return(MagickFalse);
3386 if ((colorspace != HSLColorspace && colorspace != UndefinedColorspace))
3387 return(MagickFalse);
3389 status = ComputeModulateImage(image, percent_brightness, percent_hue, percent_saturation, colorspace, exception);
3405 static Image* ComputeMotionBlurImage(
const Image *image,
3406 const ChannelType channel,
const double *kernel,
const size_t width,
3428 filteredImageBuffer,
3458 global_work_size[2],
3467 outputReady = MagickFalse;
3469 filteredImage = NULL;
3471 filteredImageBuffer = NULL;
3472 imageKernelBuffer = NULL;
3473 motionBlurKernel = NULL;
3476 clEnv = GetDefaultOpenCLEnv();
3477 context = GetOpenCLContext(clEnv);
3479 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3480 if (filteredImage == (
Image *) NULL)
3483 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
3484 if (imageBuffer == (cl_mem) NULL)
3486 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3487 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3490 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage, exception);
3491 if (filteredImageBuffer == (cl_mem) NULL)
3493 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3494 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3498 imageKernelBuffer = clEnv->library->clCreateBuffer(context,
3499 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width *
sizeof(
float), NULL,
3501 if (clStatus != CL_SUCCESS)
3503 (void) ThrowMagickException(exception, GetMagickModule(),
3504 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3508 queue = AcquireOpenCLCommandQueue(clEnv);
3509 events=GetOpenCLEvents(image,&event_count);
3511 kernelBufferPtr = (
float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
3512 CL_TRUE, CL_MAP_WRITE, 0, width *
sizeof(
float), event_count, events, NULL, &clStatus);
3513 events=(cl_event *) RelinquishMagickMemory(events);
3514 if (clStatus != CL_SUCCESS)
3516 (void) ThrowMagickException(exception, GetMagickModule(),
3517 ResourceLimitError,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
3520 for (i = 0; i < width; i++)
3522 kernelBufferPtr[i] = (float) kernel[i];
3524 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr,
3526 if (clStatus != CL_SUCCESS)
3528 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3529 "clEnv->library->clEnqueueUnmapMemObject failed.",
"'%s'",
".");
3533 offsetBuffer = clEnv->library->clCreateBuffer(context,
3534 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width *
sizeof(cl_int2), NULL,
3536 if (clStatus != CL_SUCCESS)
3538 (void) ThrowMagickException(exception, GetMagickModule(),
3539 ResourceLimitError,
"clEnv->library->clCreateBuffer failed.",
".");
3543 offsetBufferPtr = (
int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
3544 CL_MAP_WRITE, 0, width *
sizeof(cl_int2), 0, NULL, NULL, &clStatus);
3545 if (clStatus != CL_SUCCESS)
3547 (void) ThrowMagickException(exception, GetMagickModule(),
3548 ResourceLimitError,
"clEnv->library->clEnqueueMapBuffer failed.",
".");
3551 for (i = 0; i < width; i++)
3553 offsetBufferPtr[2*i] = (int)offset[i].x;
3554 offsetBufferPtr[2*i+1] = (int)offset[i].y;
3556 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
3558 if (clStatus != CL_SUCCESS)
3560 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3561 "clEnv->library->clEnqueueUnmapMemObject failed.",
"'%s'",
".");
3569 motionBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
3571 if (motionBlurKernel == NULL)
3573 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3574 "AcquireOpenCLKernel failed.",
"'%s'",
".");
3582 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3583 (
void *)&imageBuffer);
3584 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3585 (
void *)&filteredImageBuffer);
3586 imageWidth = (
unsigned int) image->columns;
3587 imageHeight = (
unsigned int) image->rows;
3588 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3590 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3592 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3593 (
void *)&imageKernelBuffer);
3594 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int),
3596 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_mem),
3597 (
void *)&offsetBuffer);
3599 GetMagickPixelPacket(image,&bias);
3600 biasPixel.s[0] = bias.red;
3601 biasPixel.s[1] = bias.green;
3602 biasPixel.s[2] = bias.blue;
3603 biasPixel.s[3] = bias.opacity;
3604 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(cl_float4), &biasPixel);
3606 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(ChannelType), &channel);
3607 matte = (image->matte != MagickFalse)?1:0;
3608 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,
sizeof(
unsigned int), &matte);
3609 if (clStatus != CL_SUCCESS)
3611 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3612 "clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
3619 local_work_size[0] = 16;
3620 local_work_size[1] = 16;
3621 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3622 (
unsigned int) image->columns,(
unsigned int) local_work_size[0]);
3623 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3624 (
unsigned int) image->rows,(
unsigned int) local_work_size[1]);
3625 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
3626 global_work_size, local_work_size, 0, NULL, &event);
3628 if (clStatus != CL_SUCCESS)
3630 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3631 "clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
3634 if (RecordProfileData(clEnv,MotionBlurKernel,event) == MagickFalse)
3636 AddOpenCLEvent(image, event);
3637 AddOpenCLEvent(filteredImage, event);
3639 clEnv->library->clReleaseEvent(event);
3641 outputReady = MagickTrue;
3645 if (imageBuffer != (cl_mem) NULL)
3646 clEnv->library->clReleaseMemObject(imageBuffer);
3647 if (filteredImageBuffer != (cl_mem) NULL)
3648 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3649 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
3650 if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, motionBlurKernel);
3651 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3652 if ((outputReady == MagickFalse) && (filteredImage != NULL))
3653 filteredImage=(
Image *) DestroyImage(filteredImage);
3655 return(filteredImage);
3658 MagickPrivate
Image *AccelerateMotionBlurImage(
const Image *image,
3659 const ChannelType channel,
const double* kernel,
const size_t width,
3665 assert(image != NULL);
3666 assert(kernel != (
double *) NULL);
3670 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3671 (checkAccelerateCondition(image, channel) == MagickFalse))
3674 filteredImage=ComputeMotionBlurImage(image, channel, kernel, width,
3676 return(filteredImage);
3691 static Image *ComputeRadialBlurImage(
const Image *image,
3692 const ChannelType channel,
const double angle,
ExceptionInfo *exception)
3711 filteredImageBuffer,
3747 global_work_size[2];
3754 outputReady = MagickFalse;
3756 filteredImage = NULL;
3758 filteredImageBuffer = NULL;
3759 sinThetaBuffer = NULL;
3760 cosThetaBuffer = NULL;
3762 radialBlurKernel = NULL;
3765 clEnv = GetDefaultOpenCLEnv();
3766 context = GetOpenCLContext(clEnv);
3768 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3769 if (filteredImage == (
Image *) NULL)
3772 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
3773 if (imageBuffer == (cl_mem) NULL)
3775 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3776 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3779 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage, exception);
3780 if (filteredImageBuffer == (cl_mem) NULL)
3782 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3783 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
3787 blurCenter.s[0] = (float) (image->columns-1)/2.0;
3788 blurCenter.s[1] = (float) (image->rows-1)/2.0;
3789 blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]);
3790 cossin_theta_size=(
unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((
double)blurRadius)+2UL);
3793 sinThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size *
sizeof(
float), NULL, &clStatus);
3794 if (clStatus != CL_SUCCESS)
3796 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3799 cosThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size *
sizeof(
float), NULL, &clStatus);
3800 if (clStatus != CL_SUCCESS)
3802 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
3806 queue = AcquireOpenCLCommandQueue(clEnv);
3807 events=GetOpenCLEvents(image,&event_count);
3809 sinThetaPtr = (
float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*
sizeof(
float), event_count, events, NULL, &clStatus);
3810 events=(cl_event *) RelinquishMagickMemory(events);
3811 if (clStatus != CL_SUCCESS)
3813 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnqueuemapBuffer failed.",
".");
3817 cosThetaPtr = (
float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*
sizeof(
float), 0, NULL, NULL, &clStatus);
3818 if (clStatus != CL_SUCCESS)
3820 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnqueuemapBuffer failed.",
".");
3824 theta=DegreesToRadians(angle)/(MagickRealType) (cossin_theta_size-1);
3825 offset=theta*(MagickRealType) (cossin_theta_size-1)/2.0;
3826 for (i=0; i < (ssize_t) cossin_theta_size; i++)
3828 cosThetaPtr[i]=(float)cos((
double) (theta*i-offset));
3829 sinThetaPtr[i]=(float)sin((
double) (theta*i-offset));
3832 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL);
3833 clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL);
3834 if (clStatus != CL_SUCCESS)
3836 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueUnmapMemObject failed.",
"'%s'",
".");
3841 radialBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"RadialBlur");
3842 if (radialBlurKernel == NULL)
3844 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
3851 clStatus=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
3852 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
3854 GetMagickPixelPacket(image,&bias);
3855 biasPixel.s[0] = bias.red;
3856 biasPixel.s[1] = bias.green;
3857 biasPixel.s[2] = bias.blue;
3858 biasPixel.s[3] = bias.opacity;
3859 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(cl_float4), &biasPixel);
3860 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(ChannelType), &channel);
3862 matte = (image->matte != MagickFalse)?1:0;
3863 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(
unsigned int), &matte);
3865 clStatus=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(cl_float2), &blurCenter);
3867 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(cl_mem),(
void *)&cosThetaBuffer);
3868 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(cl_mem),(
void *)&sinThetaBuffer);
3869 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,
sizeof(
unsigned int), &cossin_theta_size);
3870 if (clStatus != CL_SUCCESS)
3872 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
3877 global_work_size[0] = image->columns;
3878 global_work_size[1] = image->rows;
3880 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, radialBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
3881 if (clStatus != CL_SUCCESS)
3883 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
3886 if (RecordProfileData(clEnv,RadialBlurKernel,event) == MagickFalse)
3888 AddOpenCLEvent(image,event);
3889 AddOpenCLEvent(filteredImage,event);
3891 clEnv->library->clReleaseEvent(event);
3893 outputReady = MagickTrue;
3896 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3898 if (imageBuffer != (cl_mem) NULL)
3899 clEnv->library->clReleaseMemObject(imageBuffer);
3900 if (filteredImageBuffer != (cl_mem) NULL)
3901 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3902 if (sinThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(sinThetaBuffer);
3903 if (cosThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(cosThetaBuffer);
3904 if (radialBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, radialBlurKernel);
3905 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3906 if ((outputReady == MagickFalse) && (filteredImage != NULL))
3907 filteredImage=(
Image *) DestroyImage(filteredImage);
3908 return filteredImage;
3911 MagickPrivate
Image *AccelerateRadialBlurImage(
const Image *image,
3912 const ChannelType channel,
const double angle,
ExceptionInfo *exception)
3917 assert(image != NULL);
3920 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3921 (checkAccelerateCondition(image, channel) == MagickFalse))
3924 filteredImage=ComputeRadialBlurImage(image, channel, angle, exception);
3925 return filteredImage;
3940 static MagickBooleanType resizeHorizontalFilter(
const Image *image,
3941 const Image *filteredImage,cl_mem imageBuffer,
const unsigned int imageColumns,
3942 const unsigned int imageRows,
const unsigned int matte,cl_mem resizedImage,
3943 const unsigned int resizedColumns,
const unsigned int resizedRows,
3944 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3945 const float xFactor,
MagickCLEnv clEnv,cl_command_queue queue,
3964 workgroupSize = 256;
3968 resizeFilterSupport,
3969 resizeFilterWindowSupport,
3982 status = MagickFalse;
3985 deviceLocalMemorySize,
3986 gammaAccumulatorLocalMemorySize,
3987 global_work_size[2],
3988 imageCacheLocalMemorySize,
3989 pixelAccumulatorLocalMemorySize,
3991 totalLocalMemorySize,
3992 weightAccumulatorLocalMemorySize;
3999 horizontalKernel = NULL;
4000 status = MagickFalse;
4005 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
4006 support=scale*GetResizeFilterSupport(resizeFilter);
4013 support=(MagickRealType) 0.5;
4016 scale=PerceptibleReciprocal(scale);
4018 if (resizedColumns < workgroupSize)
4021 pixelPerWorkgroup = 32;
4025 chunkSize = workgroupSize;
4026 pixelPerWorkgroup = workgroupSize;
4030 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
4032 DisableMSCWarning(4127)
4037 cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
4038 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
4039 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
4040 imageCacheLocalMemorySize = numCachedPixels *
sizeof(CLPixelPacket);
4041 totalLocalMemorySize = imageCacheLocalMemorySize;
4044 pixelAccumulatorLocalMemorySize = chunkSize *
sizeof(cl_float4);
4045 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
4048 weightAccumulatorLocalMemorySize = chunkSize *
sizeof(float);
4049 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
4053 gammaAccumulatorLocalMemorySize =
sizeof(float);
4055 gammaAccumulatorLocalMemorySize = chunkSize *
sizeof(float);
4056 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
4058 if (totalLocalMemorySize <= deviceLocalMemorySize)
4062 pixelPerWorkgroup = pixelPerWorkgroup/2;
4063 chunkSize = chunkSize/2;
4064 if (pixelPerWorkgroup == 0
4073 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
4074 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
4076 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"ResizeHorizontalFilter");
4077 if (horizontalKernel == NULL)
4079 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
4084 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(cl_mem), (
void*)&imageBuffer);
4085 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&imageColumns);
4086 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&imageRows);
4087 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&matte);
4088 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&xFactor);
4089 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(cl_mem), (
void*)&resizedImage);
4091 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&resizedColumns);
4092 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&resizedRows);
4094 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
int), (
void*)&resizeFilterType);
4095 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
int), (
void*)&resizeWindowType);
4096 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(cl_mem), (
void*)&resizeFilterCubicCoefficients);
4098 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
4099 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterScale);
4101 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
4102 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterSupport);
4104 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
4105 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterWindowSupport);
4107 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
4108 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterBlur);
4111 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
4112 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
int), &numCachedPixels);
4113 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), &pixelPerWorkgroup);
4114 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), &chunkSize);
4117 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
4118 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
4119 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
4121 if (clStatus != CL_SUCCESS)
4123 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
4127 global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
4128 global_work_size[1] = resizedRows;
4130 local_work_size[0] = workgroupSize;
4131 local_work_size[1] = 1;
4132 events=GetOpenCLEvents(image,&event_count);
4133 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, event_count, events, &event);
4134 events=(cl_event *) RelinquishMagickMemory(events);
4135 if (clStatus != CL_SUCCESS)
4137 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
4140 if (RecordProfileData(clEnv,ResizeHorizontalKernel,event) == MagickFalse)
4142 AddOpenCLEvent(image,event);
4143 AddOpenCLEvent(filteredImage,event);
4145 clEnv->library->clReleaseEvent(event);
4146 status = MagickTrue;
4150 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4152 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
4157 static MagickBooleanType resizeVerticalFilter(
const Image *image,
4158 const Image *filteredImage,cl_mem imageBuffer,
const unsigned int imageColumns,
4159 const unsigned int imageRows,
const unsigned int matte,cl_mem resizedImage,
4160 const unsigned int resizedColumns,
const unsigned int resizedRows,
4161 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
4162 const float yFactor,
MagickCLEnv clEnv,cl_command_queue queue,
4181 workgroupSize = 256;
4185 resizeFilterSupport,
4186 resizeFilterWindowSupport,
4199 status = MagickFalse;
4202 deviceLocalMemorySize,
4203 gammaAccumulatorLocalMemorySize,
4204 global_work_size[2],
4205 imageCacheLocalMemorySize,
4206 pixelAccumulatorLocalMemorySize,
4208 totalLocalMemorySize,
4209 weightAccumulatorLocalMemorySize;
4216 horizontalKernel = NULL;
4217 status = MagickFalse;
4222 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
4223 support=scale*GetResizeFilterSupport(resizeFilter);
4230 support=(MagickRealType) 0.5;
4233 scale=PerceptibleReciprocal(scale);
4235 if (resizedRows < workgroupSize)
4238 pixelPerWorkgroup = 32;
4242 chunkSize = workgroupSize;
4243 pixelPerWorkgroup = workgroupSize;
4247 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
4249 DisableMSCWarning(4127)
4254 cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
4255 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
4256 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
4257 imageCacheLocalMemorySize = numCachedPixels *
sizeof(CLPixelPacket);
4258 totalLocalMemorySize = imageCacheLocalMemorySize;
4261 pixelAccumulatorLocalMemorySize = chunkSize *
sizeof(cl_float4);
4262 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
4265 weightAccumulatorLocalMemorySize = chunkSize *
sizeof(float);
4266 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
4270 gammaAccumulatorLocalMemorySize =
sizeof(float);
4272 gammaAccumulatorLocalMemorySize = chunkSize *
sizeof(float);
4273 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
4275 if (totalLocalMemorySize <= deviceLocalMemorySize)
4279 pixelPerWorkgroup = pixelPerWorkgroup/2;
4280 chunkSize = chunkSize/2;
4281 if (pixelPerWorkgroup == 0
4290 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
4291 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
4293 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"ResizeVerticalFilter");
4294 if (horizontalKernel == NULL)
4296 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
4301 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(cl_mem), (
void*)&imageBuffer);
4302 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&imageColumns);
4303 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&imageRows);
4304 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&matte);
4305 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&yFactor);
4306 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(cl_mem), (
void*)&resizedImage);
4308 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&resizedColumns);
4309 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), (
void*)&resizedRows);
4311 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
int), (
void*)&resizeFilterType);
4312 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
int), (
void*)&resizeWindowType);
4313 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(cl_mem), (
void*)&resizeFilterCubicCoefficients);
4315 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
4316 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterScale);
4318 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
4319 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterSupport);
4321 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
4322 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterWindowSupport);
4324 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
4325 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
float), (
void*)&resizeFilterBlur);
4328 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
4329 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
int), &numCachedPixels);
4330 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), &pixelPerWorkgroup);
4331 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++,
sizeof(
unsigned int), &chunkSize);
4334 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
4335 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
4336 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
4338 if (clStatus != CL_SUCCESS)
4340 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
4344 global_work_size[0] = resizedColumns;
4345 global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
4347 local_work_size[0] = 1;
4348 local_work_size[1] = workgroupSize;
4349 events=GetOpenCLEvents(image,&event_count);
4350 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, event_count, events, &event);
4351 events=(cl_event *) RelinquishMagickMemory(events);
4352 if (clStatus != CL_SUCCESS)
4354 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
4357 if (RecordProfileData(clEnv,ResizeVerticalKernel,event) == MagickFalse)
4359 AddOpenCLEvent(image,event);
4360 AddOpenCLEvent(filteredImage,event);
4362 clEnv->library->clReleaseEvent(event);
4363 status = MagickTrue;
4367 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4369 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
4374 static Image *ComputeResizeImage(
const Image* image,
4375 const size_t resizedColumns,
const size_t resizedRows,
4388 cubicCoefficientsBuffer,
4389 filteredImageBuffer,
4393 const MagickRealType
4394 *resizeFilterCoefficient;
4397 coefficientBuffer[7],
4417 outputReady = MagickFalse;
4418 filteredImage = NULL;
4422 tempImageBuffer = NULL;
4423 filteredImageBuffer = NULL;
4424 cubicCoefficientsBuffer = NULL;
4427 clEnv = GetDefaultOpenCLEnv();
4428 context = GetOpenCLContext(clEnv);
4429 queue = AcquireOpenCLCommandQueue(clEnv);
4431 filteredImage=CloneImage(image,resizedColumns,resizedRows,MagickTrue,exception);
4432 if (filteredImage == (
Image *) NULL)
4435 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
4436 if (imageBuffer == (cl_mem) NULL)
4438 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4439 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
4442 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,exception);
4443 if (filteredImageBuffer == (cl_mem) NULL)
4445 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4446 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
4450 resizeFilterCoefficient=GetResizeFilterCoefficient(resizeFilter);
4451 for (i = 0; i < 7; i++)
4452 coefficientBuffer[i]=(
float) resizeFilterCoefficient[i];
4454 cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(coefficientBuffer), coefficientBuffer, &clStatus);
4455 if (clStatus != CL_SUCCESS)
4457 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
4461 xFactor=(float) resizedColumns/(
float) image->columns;
4462 yFactor=(float) resizedRows/(
float) image->rows;
4463 if (xFactor > yFactor)
4466 length = resizedColumns*image->rows;
4467 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*
sizeof(CLPixelPacket), NULL, &clStatus);
4468 if (clStatus != CL_SUCCESS)
4470 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
4474 status = resizeHorizontalFilter(image,filteredImage,imageBuffer, (
unsigned int) image->columns, (
unsigned int) image->rows, (image->matte != MagickFalse)?1:0
4475 , tempImageBuffer, (
unsigned int) resizedColumns, (
unsigned int) image->rows
4476 , resizeFilter, cubicCoefficientsBuffer
4477 , xFactor, clEnv, queue, exception);
4478 if (status != MagickTrue)
4481 status = resizeVerticalFilter(image,filteredImage,tempImageBuffer, (
unsigned int) resizedColumns, (
unsigned int) image->rows, (image->matte != MagickFalse)?1:0
4482 , filteredImageBuffer, (
unsigned int) resizedColumns, (
unsigned int) resizedRows
4483 , resizeFilter, cubicCoefficientsBuffer
4484 , yFactor, clEnv, queue, exception);
4485 if (status != MagickTrue)
4490 length = image->columns*resizedRows;
4491 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*
sizeof(CLPixelPacket), NULL, &clStatus);
4492 if (clStatus != CL_SUCCESS)
4494 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
4498 status = resizeVerticalFilter(image,filteredImage,imageBuffer, (
unsigned int) image->columns, (
unsigned int) image->rows, (image->matte != MagickFalse)?1:0
4499 , tempImageBuffer, (
unsigned int) image->columns, (
unsigned int) resizedRows
4500 , resizeFilter, cubicCoefficientsBuffer
4501 , yFactor, clEnv, queue, exception);
4502 if (status != MagickTrue)
4505 status = resizeHorizontalFilter(image,filteredImage,tempImageBuffer, (
unsigned int) image->columns, (
unsigned int) resizedRows, (image->matte != MagickFalse)?1:0
4506 , filteredImageBuffer, (
unsigned int) resizedColumns, (
unsigned int) resizedRows
4507 , resizeFilter, cubicCoefficientsBuffer
4508 , xFactor, clEnv, queue, exception);
4509 if (status != MagickTrue)
4512 outputReady=MagickTrue;
4515 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4517 if (imageBuffer != (cl_mem) NULL)
4518 clEnv->library->clReleaseMemObject(imageBuffer);
4519 if (filteredImageBuffer != (cl_mem) NULL)
4520 clEnv->library->clReleaseMemObject(filteredImageBuffer);
4521 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
4522 if (cubicCoefficientsBuffer!=NULL) clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer);
4523 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4524 if ((outputReady == MagickFalse) && (filteredImage != NULL))
4525 filteredImage=(
Image *) DestroyImage(filteredImage);
4526 return(filteredImage);
4529 static MagickBooleanType gpuSupportedResizeWeighting(
4530 ResizeWeightingFunctionType f)
4537 if (supportedResizeWeighting[i] == LastWeightingFunction)
4539 if (supportedResizeWeighting[i] == f)
4542 return(MagickFalse);
4545 MagickPrivate
Image *AccelerateResizeImage(
const Image *image,
4546 const size_t resizedColumns,
const size_t resizedRows,
4552 assert(image != NULL);
4555 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
4556 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
4559 if (gpuSupportedResizeWeighting(GetResizeFilterWeightingType(resizeFilter)) == MagickFalse ||
4560 gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(resizeFilter)) == MagickFalse)
4563 filteredImage=ComputeResizeImage(image,resizedColumns,resizedRows,resizeFilter,exception);
4564 return(filteredImage);
4579 static Image *ComputeUnsharpMaskImage(
const Image *image,
4580 const ChannelType channel,
const double radius,
const double sigma,
4581 const double gain,
const double threshold,
ExceptionInfo *exception)
4584 geometry[MaxTextExtent];
4600 unsharpMaskBlurColumnKernel;
4603 filteredImageBuffer,
4646 filteredImage = NULL;
4650 filteredImageBuffer = NULL;
4651 tempImageBuffer = NULL;
4652 imageKernelBuffer = NULL;
4653 blurRowKernel = NULL;
4654 unsharpMaskBlurColumnKernel = NULL;
4656 outputReady = MagickFalse;
4658 clEnv = GetDefaultOpenCLEnv();
4659 context = GetOpenCLContext(clEnv);
4660 queue = AcquireOpenCLCommandQueue(clEnv);
4662 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
4663 if (filteredImage == (
Image *) NULL)
4666 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
4667 if (imageBuffer == (cl_mem) NULL)
4669 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4670 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
4673 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,exception);
4674 if (filteredImageBuffer == (cl_mem) NULL)
4676 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4677 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
4683 (void) FormatLocaleString(geometry,MaxTextExtent,
"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
4684 kernel=AcquireKernelInfo(geometry);
4687 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireKernelInfo failed.",
".");
4691 kernelBufferPtr=AcquireQuantumMemory(kernel->width,
sizeof(
float));
4692 if (kernelBufferPtr == (
float *) NULL)
4694 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"Memory allocation failed.",
".");
4697 for (i = 0; i < kernel->width; i++)
4698 kernelBufferPtr[i]=(
float) kernel->values[i];
4700 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, kernel->width *
sizeof(
float), kernelBufferPtr, &clStatus);
4701 kernelBufferPtr=RelinquishMagickMemory(kernelBufferPtr);
4702 if (clStatus != CL_SUCCESS)
4704 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
4712 length = image->columns * image->rows;
4713 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 *
sizeof(
float), NULL, &clStatus);
4714 if (clStatus != CL_SUCCESS)
4716 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
4723 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"BlurRow");
4724 if (blurRowKernel == NULL)
4726 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
4730 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"UnsharpMaskBlurColumn");
4731 if (unsharpMaskBlurColumnKernel == NULL)
4733 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
4741 imageColumns = (
unsigned int) image->columns;
4742 imageRows = (
unsigned int) image->rows;
4744 kernelWidth = (
unsigned int) kernel->width;
4748 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4749 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
4750 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(ChannelType),&channel);
4751 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4752 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&kernelWidth);
4753 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
4754 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
4755 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,
sizeof(CLPixelPacket)*(chunkSize+kernel->width),(
void *) NULL);
4756 if (clStatus != CL_SUCCESS)
4758 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
4768 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
4769 gsize[1] = image->rows;
4770 wsize[0] = chunkSize;
4773 events=GetOpenCLEvents(image,&event_count);
4774 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, event_count, events, NULL);
4775 events=(cl_event *) RelinquishMagickMemory(events);
4776 if (clStatus != CL_SUCCESS)
4778 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
4786 imageColumns = (
unsigned int) image->columns;
4787 imageRows = (
unsigned int) image->rows;
4788 kernelWidth = (
unsigned int) kernel->width;
4789 fGain = (
float) gain;
4790 fThreshold = (float) threshold;
4793 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4794 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&tempImageBuffer);
4795 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4796 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
4797 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
4798 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*
sizeof(cl_float4),NULL);
4799 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*
sizeof(
float),NULL);
4800 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(ChannelType),&channel);
4801 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4802 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
unsigned int),(
void *)&kernelWidth);
4803 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
float),(
void *)&fGain);
4804 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,
sizeof(
float),(
void *)&fThreshold);
4806 if (clStatus != CL_SUCCESS)
4808 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
4818 gsize[0] = image->columns;
4819 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
4821 wsize[1] = chunkSize;
4823 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, event_count, events, &event);
4824 if (clStatus != CL_SUCCESS)
4826 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
4829 if (RecordProfileData(clEnv,UnsharpMaskBlurColumnKernel,event) == MagickFalse)
4831 AddOpenCLEvent(image,event);
4832 AddOpenCLEvent(filteredImage,event);
4834 clEnv->library->clReleaseEvent(event);
4839 outputReady=MagickTrue;
4842 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4844 if (imageBuffer != (cl_mem) NULL)
4845 clEnv->library->clReleaseMemObject(imageBuffer);
4846 if (filteredImageBuffer != (cl_mem) NULL)
4847 clEnv->library->clReleaseMemObject(filteredImageBuffer);
4848 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
4849 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
4850 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
4851 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
4852 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
4853 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4854 if ((outputReady == MagickFalse) && (filteredImage != NULL))
4855 filteredImage=(
Image *) DestroyImage(filteredImage);
4856 return(filteredImage);
4859 static Image *ComputeUnsharpMaskImageSingle(
const Image *image,
4860 const double radius,
const double sigma,
const double gain,
4861 const double threshold,
int blurOnly,
ExceptionInfo *exception)
4864 geometry[MaxTextExtent];
4883 filteredImageBuffer,
4914 filteredImage = NULL;
4918 filteredImageBuffer = NULL;
4919 imageKernelBuffer = NULL;
4920 unsharpMaskKernel = NULL;
4922 outputReady = MagickFalse;
4924 clEnv = GetDefaultOpenCLEnv();
4925 context = GetOpenCLContext(clEnv);
4926 queue = AcquireOpenCLCommandQueue(clEnv);
4928 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
4929 if (filteredImage == (
Image *) NULL)
4932 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
4933 if (imageBuffer == (cl_mem) NULL)
4935 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4936 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
4939 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
4940 if (filteredImageBuffer == (cl_mem) NULL)
4942 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4943 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
4949 (void) FormatLocaleString(geometry,MaxTextExtent,
"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
4950 kernel=AcquireKernelInfo(geometry);
4953 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireKernelInfo failed.",
".");
4958 float *kernelBufferPtr = (
float *) AcquireQuantumMemory(kernel->width,
sizeof(
float));
4959 for (i = 0; i < kernel->width; i++)
4960 kernelBufferPtr[i] = (
float)kernel->values[i];
4962 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, kernel->width *
sizeof(
float), kernelBufferPtr, &clStatus);
4963 RelinquishMagickMemory(kernelBufferPtr);
4964 if (clStatus != CL_SUCCESS)
4966 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clCreateBuffer failed.",
".");
4975 unsharpMaskKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"UnsharpMask");
4976 if (unsharpMaskKernel == NULL)
4978 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
4984 imageColumns = (
unsigned int) image->columns;
4985 imageRows = (
unsigned int) image->rows;
4986 kernelWidth = (
unsigned int) kernel->width;
4987 fGain = (
float) gain;
4988 fThreshold = (float) threshold;
4989 justBlur = blurOnly;
4993 clStatus=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&imageBuffer);
4994 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&filteredImageBuffer);
4995 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(cl_mem),(
void *)&imageKernelBuffer);
4996 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(
unsigned int),(
void *)&kernelWidth);
4997 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(
unsigned int),(
void *)&imageColumns);
4998 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(
unsigned int),(
void *)&imageRows);
4999 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(cl_float4)*(8 * (32 + kernel->width)),(
void *) NULL);
5000 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(
float),(
void *)&fGain);
5001 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(
float),(
void *)&fThreshold);
5002 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,
sizeof(cl_uint),(
void *)&justBlur);
5003 if (clStatus != CL_SUCCESS)
5005 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clSetKernelArg failed.",
"'%s'",
".");
5015 gsize[0] = ((image->columns + 7) / 8) * 8;
5016 gsize[1] = ((image->rows + 31) / 32) * 32;
5020 events=GetOpenCLEvents(image,&event_count);
5021 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, event_count, events, &event);
5022 events=(cl_event *) RelinquishMagickMemory(events);
5023 if (clStatus != CL_SUCCESS)
5025 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
5028 if (RecordProfileData(clEnv,UnsharpMaskKernel,event) == MagickFalse)
5030 AddOpenCLEvent(image,event);
5031 AddOpenCLEvent(filteredImage, event);
5033 clEnv->library->clReleaseEvent(event);
5037 outputReady=MagickTrue;
5040 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5042 if (imageBuffer != (cl_mem) NULL)
5043 clEnv->library->clReleaseMemObject(imageBuffer);
5044 if (filteredImageBuffer != (cl_mem) NULL)
5045 clEnv->library->clReleaseMemObject(filteredImageBuffer);
5046 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
5047 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
5048 if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskKernel);
5049 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
5050 if ((outputReady == MagickFalse) && (filteredImage != NULL))
5051 filteredImage=(
Image *) DestroyImage(filteredImage);
5052 return(filteredImage);
5055 MagickPrivate
Image *AccelerateUnsharpMaskImage(
const Image *image,
5056 const ChannelType channel,
const double radius,
const double sigma,
5057 const double gain,
const double threshold,
ExceptionInfo *exception)
5062 assert(image != NULL);
5065 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
5066 (checkAccelerateCondition(image, channel) == MagickFalse))
5070 filteredImage = ComputeUnsharpMaskImageSingle(image,radius,sigma,gain,threshold, 0, exception);
5072 filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
5074 return(filteredImage);
5077 static Image *ComputeWaveletDenoiseImage(
const Image *image,
5096 filteredImageBuffer,
5117 filteredImage = NULL;
5120 filteredImageBuffer = NULL;
5121 denoiseKernel = NULL;
5123 outputReady = MagickFalse;
5125 clEnv = GetDefaultOpenCLEnv();
5128 if (paramMatchesValue(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE_NAME,
5129 "Intel(R) HD Graphics",exception) != MagickFalse)
5132 context = GetOpenCLContext(clEnv);
5133 queue = AcquireOpenCLCommandQueue(clEnv);
5135 filteredImage = CloneImage(image,0,0,MagickTrue, exception);
5136 if (filteredImage == (
Image *) NULL)
5139 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
5140 if (imageBuffer == (cl_mem) NULL)
5142 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
5143 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
5146 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
5147 if (filteredImageBuffer == (cl_mem) NULL)
5149 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
5150 ResourceLimitWarning,
"GetAuthenticOpenCLBuffer failed.",
".");
5155 denoiseKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
"WaveletDenoise");
5156 if (denoiseKernel == NULL)
5158 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"AcquireOpenCLKernel failed.",
"'%s'",
".");
5167 const int PASSES = 5;
5168 cl_int width = (cl_int)image->columns;
5169 cl_int height = (cl_int)image->rows;
5170 cl_float thresh = threshold;
5172 passes = (((1.0f * image->columns) * image->rows) + 1999999.0f) / 2000000.0f;
5173 passes = (passes < 1) ? 1 : passes;
5177 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++,
sizeof(cl_mem), (
void *)&imageBuffer);
5178 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++,
sizeof(cl_mem), (
void *)&filteredImageBuffer);
5179 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++,
sizeof(cl_float), (
void *)&thresh);
5180 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++,
sizeof(cl_int), (
void *)&PASSES);
5181 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++,
sizeof(cl_int), (
void *)&width);
5182 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++,
sizeof(cl_int), (
void *)&height);
5184 for (x = 0; x < passes; ++x)
5186 const int TILESIZE = 64;
5187 const int PAD = 1 << (PASSES - 1);
5188 const int SIZE = TILESIZE - 2 * PAD;
5194 gsize[0] = ((width + (SIZE - 1)) / SIZE) * TILESIZE;
5195 gsize[1] = ((((height + (SIZE - 1)) / SIZE) + passes - 1) / passes) * 4;
5196 wsize[0] = TILESIZE;
5199 goffset[1] = x * gsize[1];
5201 events=GetOpenCLEvents(image,&event_count);
5202 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, denoiseKernel, 2, goffset, gsize, wsize, event_count, events, &event);
5203 events=(cl_event *) RelinquishMagickMemory(events);
5204 if (clStatus != CL_SUCCESS)
5206 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning,
"clEnv->library->clEnqueueNDRangeKernel failed.",
"'%s'",
".");
5209 clEnv->library->clFlush(queue);
5210 if (RecordProfileData(clEnv, WaveletDenoiseKernel, event) == MagickFalse)
5212 AddOpenCLEvent(image, event);
5213 AddOpenCLEvent(filteredImage, event);
5215 clEnv->library->clReleaseEvent(event);
5219 outputReady=MagickTrue;
5222 OpenCLLogException(__FUNCTION__, __LINE__, exception);
5224 if (imageBuffer != (cl_mem) NULL)
5225 clEnv->library->clReleaseMemObject(imageBuffer);
5226 if (filteredImageBuffer != (cl_mem) NULL)
5227 clEnv->library->clReleaseMemObject(filteredImageBuffer);
5228 if (denoiseKernel != NULL)
5229 RelinquishOpenCLKernel(clEnv, denoiseKernel);
5231 RelinquishOpenCLCommandQueue(clEnv, queue);
5232 if ((outputReady == MagickFalse) && (filteredImage != NULL))
5233 filteredImage=(
Image *) DestroyImage(filteredImage);
5234 return(filteredImage);
5237 MagickPrivate
Image *AccelerateWaveletDenoiseImage(
const Image *image,
5243 assert(image != NULL);
5246 if ((checkAccelerateCondition(image,DefaultChannels) == MagickFalse) ||
5247 (checkOpenCLEnvironment(exception) == MagickFalse))
5248 return (
Image *) NULL;
5250 filteredImage=ComputeWaveletDenoiseImage(image,threshold,exception);
5252 return(filteredImage);