MagickCore 6.9.13
Loading...
Searching...
No Matches
accelerate.c
1/*
2%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3% %
4% %
5% %
6% AAA CCCC CCCC EEEEE L EEEEE RRRR AAA TTTTT EEEEE %
7% A A C C E L E R R A A T E %
8% AAAAA C C EEE L EEE RRRR AAAAA T EEE %
9% A A C C E L E R R A A T E %
10% A A CCCC CCCC EEEEE LLLLL EEEEE R R A A T EEEEE %
11% %
12% %
13% MagickCore Acceleration Methods %
14% %
15% Software Design %
16% Cristy %
17% SiuChi Chan %
18% Guansong Zhang %
19% January 2010 %
20% Dirk Lemstra %
21% May 2016 %
22% %
23% %
24% Copyright 1999 ImageMagick Studio LLC, a non-profit organization %
25% dedicated to making software imaging solutions freely available. %
26% %
27% You may not use this file except in compliance with the License. You may %
28% obtain a copy of the License at %
29% %
30% https://imagemagick.org/script/license.php %
31% %
32% Unless required by applicable law or agreed to in writing, software %
33% distributed under the License is distributed on an "AS IS" BASIS, %
34% WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. %
35% See the License for the specific language governing permissions and %
36% limitations under the License. %
37% %
38%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
39*/
40
41/*
42Include declarations.
43*/
44#include "magick/studio.h"
45#include "magick/accelerate-private.h"
46#include "magick/artifact.h"
47#include "magick/cache.h"
48#include "magick/cache-private.h"
49#include "magick/cache-view.h"
50#include "magick/color-private.h"
51#include "magick/delegate-private.h"
52#include "magick/enhance.h"
53#include "magick/exception.h"
54#include "magick/exception-private.h"
55#include "magick/gem.h"
56#include "magick/hashmap.h"
57#include "magick/image.h"
58#include "magick/image-private.h"
59#include "magick/list.h"
60#include "magick/memory_.h"
61#include "magick/monitor-private.h"
62#include "magick/opencl.h"
63#include "magick/opencl-private.h"
64#include "magick/option.h"
65#include "magick/pixel-private.h"
66#include "magick/prepress.h"
67#include "magick/quantize.h"
68#include "magick/random_.h"
69#include "magick/random-private.h"
70#include "magick/registry.h"
71#include "magick/resize.h"
72#include "magick/resize-private.h"
73#include "magick/semaphore.h"
74#include "magick/splay-tree.h"
75#include "magick/statistic.h"
76#include "magick/string_.h"
77#include "magick/string-private.h"
78#include "magick/token.h"
79
80#ifdef MAGICKCORE_CLPERFMARKER
81#include "CLPerfMarker.h"
82#endif
83
84#define MAGICK_MAX(x,y) (((x) >= (y))?(x):(y))
85#define MAGICK_MIN(x,y) (((x) <= (y))?(x):(y))
86
87#if defined(MAGICKCORE_OPENCL_SUPPORT)
88
89/*
90 Define declarations.
91*/
92#define ALIGNED(pointer,type) ((((size_t)(pointer)) & (sizeof(type)-1)) == 0)
93
94/*
95 Static declarations.
96*/
97static const ResizeWeightingFunctionType supportedResizeWeighting[] =
98{
99 BoxWeightingFunction,
100 TriangleWeightingFunction,
101 HanningWeightingFunction,
102 HammingWeightingFunction,
103 BlackmanWeightingFunction,
104 CubicBCWeightingFunction,
105 SincWeightingFunction,
106 SincFastWeightingFunction,
107 LastWeightingFunction
108};
109
110/*
111 Forward declarations.
112*/
113static Image *ComputeUnsharpMaskImageSingle(const Image *image,
114 const double radius,const double sigma,const double gain,
115 const double threshold,int blurOnly, ExceptionInfo *exception);
116
117/*
118 Helper functions.
119*/
120
121static MagickBooleanType checkAccelerateCondition(const Image* image,
122 const ChannelType channel)
123{
124 /* only direct class images are supported */
125 if (image->storage_class != DirectClass)
126 return(MagickFalse);
127
128 /* check if the image's colorspace is supported */
129 if (image->colorspace != RGBColorspace &&
130 image->colorspace != sRGBColorspace &&
131 image->colorspace != LinearGRAYColorspace &&
132 image->colorspace != GRAYColorspace)
133 return(MagickFalse);
134
135 /* check if the channel is supported */
136 if (((channel & RedChannel) == 0) ||
137 ((channel & GreenChannel) == 0) ||
138 ((channel & BlueChannel) == 0))
139 return(MagickFalse);
140
141 /* check if the virtual pixel method is compatible with the OpenCL implementation */
142 if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
143 (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod))
144 return(MagickFalse);
145
146 /* check if the image has clip_mask / mask */
147 if ((image->clip_mask != (Image *) NULL) || (image->mask != (Image *) NULL))
148 return(MagickFalse);
149
150 return(MagickTrue);
151}
152
153static MagickBooleanType checkHistogramCondition(Image *image,
154 const ChannelType channel)
155{
156 /* ensure this is the only pass get in for now. */
157 if ((channel & SyncChannels) == 0)
158 return MagickFalse;
159
160 if (image->intensity == Rec601LuminancePixelIntensityMethod ||
161 image->intensity == Rec709LuminancePixelIntensityMethod)
162 return MagickFalse;
163
164 if (image->colorspace != sRGBColorspace)
165 return MagickFalse;
166
167 return MagickTrue;
168}
169
170static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception)
171{
172 MagickBooleanType
173 flag;
174
175 MagickCLEnv
176 clEnv;
177
178 clEnv=GetDefaultOpenCLEnv();
179
180 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,
181 sizeof(MagickBooleanType),&flag,exception);
182 if (flag != MagickFalse)
183 return(MagickFalse);
184
185 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED,
186 sizeof(MagickBooleanType),&flag,exception);
187 if (flag == MagickFalse)
188 {
189 if (InitOpenCLEnv(clEnv,exception) == MagickFalse)
190 return(MagickFalse);
191
192 GetMagickOpenCLEnvParam(clEnv,MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,
193 sizeof(MagickBooleanType),&flag,exception);
194 if (flag != MagickFalse)
195 return(MagickFalse);
196 }
197
198 return(MagickTrue);
199}
200
201/* pad the global workgroup size to the next multiple of
202 the local workgroup size */
203inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
204 const unsigned int orgGlobalSize,const unsigned int localGroupSize)
205{
206 return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
207}
208
209static MagickBooleanType paramMatchesValue(MagickCLEnv clEnv,
210 MagickOpenCLEnvParam param,const char *value,ExceptionInfo *exception)
211{
212 char
213 *val;
214
215 MagickBooleanType
216 status;
217
218 status=GetMagickOpenCLEnvParam(clEnv,param,sizeof(val),&val,exception);
219 if (status != MagickFalse)
220 {
221 status=strcmp(value,val) == 0 ? MagickTrue : MagickFalse;
222 RelinquishMagickMemory(val);
223 }
224 return(status);
225}
226
227/*
228%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
229% %
230% %
231% %
232% A c c e l e r a t e A d d N o i s e I m a g e %
233% %
234% %
235% %
236%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
237*/
238
239static Image *ComputeAddNoiseImage(const Image *image,
240 const ChannelType channel,const NoiseType noise_type,
241 ExceptionInfo *exception)
242{
243 cl_command_queue
244 queue;
245
246 cl_context
247 context;
248
249 cl_int
250 inputPixelCount,
251 pixelsPerWorkitem,
252 clStatus;
253
254 cl_uint
255 event_count,
256 seed0,
257 seed1;
258
259 cl_kernel
260 addNoiseKernel;
261
262 cl_event
263 event;
264
265 cl_mem
266 filteredImageBuffer,
267 imageBuffer;
268
269 const char
270 *option;
271
272 cl_event
273 *events;
274
275 float
276 attenuate;
277
278 MagickBooleanType
279 outputReady;
280
281 MagickCLEnv
282 clEnv;
283
284 Image
285 *filteredImage;
286
287 RandomInfo
288 **magick_restrict random_info;
289
290 size_t
291 global_work_size[1],
292 local_work_size[1];
293
294 unsigned int
295 k,
296 numRandomNumberPerPixel;
297
298#if defined(MAGICKCORE_OPENMP_SUPPORT)
299 unsigned long
300 key;
301#endif
302
303 outputReady = MagickFalse;
304 clEnv = NULL;
305 filteredImage = NULL;
306 context = NULL;
307 imageBuffer = NULL;
308 filteredImageBuffer = NULL;
309 queue = NULL;
310 addNoiseKernel = NULL;
311
312 clEnv = GetDefaultOpenCLEnv();
313 context = GetOpenCLContext(clEnv);
314 queue = AcquireOpenCLCommandQueue(clEnv);
315
316 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
317 if (filteredImage == (Image *) NULL)
318 goto cleanup;
319
320 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
321 if (imageBuffer == (cl_mem) NULL)
322 {
323 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
324 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
325 goto cleanup;
326 }
327 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
328 if (filteredImageBuffer == (cl_mem) NULL)
329 {
330 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
331 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
332 goto cleanup;
333 }
334
335 /* find out how many random numbers needed by pixel */
336 numRandomNumberPerPixel = 0;
337 {
338 unsigned int numRandPerChannel = 0;
339 switch (noise_type)
340 {
341 case UniformNoise:
342 case ImpulseNoise:
343 case LaplacianNoise:
344 case RandomNoise:
345 default:
346 numRandPerChannel = 1;
347 break;
348 case GaussianNoise:
349 case MultiplicativeGaussianNoise:
350 case PoissonNoise:
351 numRandPerChannel = 2;
352 break;
353 };
354
355 if ((channel & RedChannel) != 0)
356 numRandomNumberPerPixel+=numRandPerChannel;
357 if ((channel & GreenChannel) != 0)
358 numRandomNumberPerPixel+=numRandPerChannel;
359 if ((channel & BlueChannel) != 0)
360 numRandomNumberPerPixel+=numRandPerChannel;
361 if ((channel & OpacityChannel) != 0)
362 numRandomNumberPerPixel+=numRandPerChannel;
363 }
364
365 /* set up the random number generators */
366 attenuate=1.0;
367 option=GetImageArtifact(image,"attenuate");
368 if (option != (char *) NULL)
369 attenuate=StringToDouble(option,(char **) NULL);
370 random_info=AcquireRandomInfoTLS();
371#if defined(MAGICKCORE_OPENMP_SUPPORT)
372 key=GetRandomSecretKey(random_info[0]);
373 (void) key;
374#endif
375
376 addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoise");
377
378 {
379 cl_uint computeUnitCount;
380 cl_uint workItemCount;
381 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &computeUnitCount, NULL);
382 workItemCount = computeUnitCount * 2 * 256; // 256 work items per group, 2 groups per CU
383 inputPixelCount = (cl_int) (image->columns * image->rows);
384 pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount;
385 pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4;
386
387 local_work_size[0] = 256;
388 global_work_size[0] = workItemCount;
389 }
390 {
391 RandomInfo* randomInfo = AcquireRandomInfo();
392 const unsigned long* s = GetRandomInfoSeed(randomInfo);
393 seed0 = s[0];
394 GetPseudoRandomValue(randomInfo);
395 seed1 = s[0];
396 randomInfo = DestroyRandomInfo(randomInfo);
397 }
398
399 k = 0;
400 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&imageBuffer);
401 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
402 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&inputPixelCount);
403 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&pixelsPerWorkitem);
404 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
405 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
406 attenuate=1.0f;
407 option=GetImageArtifact(image,"attenuate");
408 if (option != (char *) NULL)
409 attenuate=(float)StringToDouble(option,(char **) NULL);
410 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
411 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed0);
412 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed1);
413 clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
414
415 events=GetOpenCLEvents(image,&event_count);
416 clStatus=clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,NULL,event_count,events,&event);
417 events=(cl_event *) RelinquishMagickMemory(events);
418 if (clStatus != CL_SUCCESS)
419 {
420 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
421 goto cleanup;
422 }
423 if (RecordProfileData(clEnv,AddNoiseKernel,event) == MagickFalse)
424 {
425 AddOpenCLEvent(image,event);
426 AddOpenCLEvent(filteredImage,event);
427 }
428 clEnv->library->clReleaseEvent(event);
429 outputReady=MagickTrue;
430
431cleanup:
432 OpenCLLogException(__FUNCTION__,__LINE__,exception);
433
434 if (imageBuffer != (cl_mem) NULL)
435 clEnv->library->clReleaseMemObject(imageBuffer);
436 if (filteredImageBuffer != (cl_mem) NULL)
437 clEnv->library->clReleaseMemObject(filteredImageBuffer);
438 if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
439 if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
440 if ((outputReady == MagickFalse) && (filteredImage != NULL))
441 filteredImage=(Image *) DestroyImage(filteredImage);
442
443 return(filteredImage);
444}
445
446MagickPrivate Image *AccelerateAddNoiseImage(const Image *image,
447 const ChannelType channel,const NoiseType noise_type,
448 ExceptionInfo *exception)
449{
450 /* Temporary disabled because of repetition.
451
452 Image
453 *filteredImage;
454
455 assert(image != NULL);
456 assert(exception != (ExceptionInfo *) NULL);
457
458 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
459 (checkAccelerateCondition(image, channel) == MagickFalse))
460 return NULL;
461
462 filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
463
464 return(filteredImage);
465 */
466 magick_unreferenced(image);
467 magick_unreferenced(channel);
468 magick_unreferenced(noise_type);
469 magick_unreferenced(exception);
470 return((Image *)NULL);
471}
472
473/*
474%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
475% %
476% %
477% %
478% A c c e l e r a t e B l u r I m a g e %
479% %
480% %
481% %
482%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
483*/
484
485static Image *ComputeBlurImage(const Image* image,const ChannelType channel,
486 const double radius,const double sigma,ExceptionInfo *exception)
487{
488 char
489 geometry[MaxTextExtent];
490
491 cl_command_queue
492 queue;
493
494 cl_context
495 context;
496
497 cl_int
498 clStatus;
499
500 cl_kernel
501 blurColumnKernel,
502 blurRowKernel;
503
504 cl_event
505 event;
506
507 cl_mem
508 filteredImageBuffer,
509 imageBuffer,
510 imageKernelBuffer,
511 tempImageBuffer;
512
513 cl_uint
514 event_count;
515
516 cl_event
517 *events;
518
519 float
520 *kernelBufferPtr;
521
522 Image
523 *filteredImage;
524
525 MagickBooleanType
526 outputReady;
527
528 MagickCLEnv
529 clEnv;
530
531 MagickSizeType
532 length;
533
534 KernelInfo
535 *kernel;
536
537 unsigned int
538 i,
539 imageColumns,
540 imageRows,
541 kernelWidth;
542
543 context = NULL;
544 filteredImage = NULL;
545 imageBuffer = NULL;
546 tempImageBuffer = NULL;
547 filteredImageBuffer = NULL;
548 imageKernelBuffer = NULL;
549 blurRowKernel = NULL;
550 blurColumnKernel = NULL;
551 queue = NULL;
552 kernel = NULL;
553
554 outputReady = MagickFalse;
555
556 clEnv = GetDefaultOpenCLEnv();
557 context = GetOpenCLContext(clEnv);
558 queue = AcquireOpenCLCommandQueue(clEnv);
559
560 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
561 if (filteredImage == (Image *) NULL)
562 goto cleanup;
563
564 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
565 if (imageBuffer == (cl_mem) NULL)
566 {
567 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
568 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
569 goto cleanup;
570 }
571 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
572 if (filteredImageBuffer == (cl_mem) NULL)
573 {
574 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
575 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
576 goto cleanup;
577 }
578
579 /* create processing kernel */
580 {
581 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
582 kernel=AcquireKernelInfo(geometry);
583 if (kernel == (KernelInfo *) NULL)
584 {
585 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.",".");
586 goto cleanup;
587 }
588
589 {
590 kernelBufferPtr = (float *)AcquireMagickMemory(kernel->width * sizeof(float));
591 if (kernelBufferPtr == (float *) NULL)
592 {
593 (void)OpenCLThrowMagickException(exception,GetMagickModule(),
594 ResourceLimitWarning,"AcquireMagickMemory failed.", "'%s'", ".");
595 goto cleanup;
596 }
597 for (i = 0; i < kernel->width; i++)
598 kernelBufferPtr[i] = (float)kernel->values[i];
599
600 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, kernel->width * sizeof(float), kernelBufferPtr, &clStatus);
601 RelinquishMagickMemory(kernelBufferPtr);
602 if (clStatus != CL_SUCCESS)
603 {
604 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.", ".");
605 goto cleanup;
606 }
607 }
608 }
609
610 {
611
612 /* create temp buffer */
613 {
614 length = image->columns * image->rows;
615 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
616 if (clStatus != CL_SUCCESS)
617 {
618 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
619 goto cleanup;
620 }
621 }
622
623 /* get the OpenCL kernels */
624 {
625 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
626 if (blurRowKernel == NULL)
627 {
628 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
629 goto cleanup;
630 };
631
632 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn");
633 if (blurColumnKernel == NULL)
634 {
635 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
636 goto cleanup;
637 };
638 }
639
640 {
641 /* need logic to decide this value */
642 int chunkSize = 256;
643
644 {
645 imageColumns = (unsigned int) image->columns;
646 imageRows = (unsigned int) image->rows;
647
648 /* set the kernel arguments */
649 i = 0;
650 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
651 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
652 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
653 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
654 kernelWidth = (unsigned int) kernel->width;
655 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
656 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
657 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
658 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
659 if (clStatus != CL_SUCCESS)
660 {
661 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
662 goto cleanup;
663 }
664 }
665
666 /* launch the kernel */
667 {
668 size_t gsize[2];
669 size_t wsize[2];
670
671 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
672 gsize[1] = image->rows;
673 wsize[0] = chunkSize;
674 wsize[1] = 1;
675
676 events=GetOpenCLEvents(image,&event_count);
677 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, event_count, events, &event);
678 events=(cl_event *) RelinquishMagickMemory(events);
679 if (clStatus != CL_SUCCESS)
680 {
681 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
682 goto cleanup;
683 }
684 if (RecordProfileData(clEnv,BlurRowKernel,event) == MagickFalse)
685 {
686 AddOpenCLEvent(image,event);
687 AddOpenCLEvent(filteredImage,event);
688 }
689 clEnv->library->clReleaseEvent(event);
690 }
691 }
692
693 {
694 /* need logic to decide this value */
695 int chunkSize = 256;
696
697 {
698 imageColumns = (unsigned int) image->columns;
699 imageRows = (unsigned int) image->rows;
700
701 /* set the kernel arguments */
702 i = 0;
703 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
704 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
705 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&channel);
706 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
707 kernelWidth = (unsigned int) kernel->width;
708 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
709 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
710 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
711 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *) NULL);
712 if (clStatus != CL_SUCCESS)
713 {
714 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
715 goto cleanup;
716 }
717 }
718
719 /* launch the kernel */
720 {
721 size_t gsize[2];
722 size_t wsize[2];
723
724 gsize[0] = image->columns;
725 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
726 wsize[0] = 1;
727 wsize[1] = chunkSize;
728
729 events=GetOpenCLEvents(image,&event_count);
730 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, event_count, events, &event);
731 events=(cl_event *) RelinquishMagickMemory(events);
732 if (clStatus != CL_SUCCESS)
733 {
734 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
735 goto cleanup;
736 }
737 if (RecordProfileData(clEnv,BlurColumnKernel,event) == MagickFalse)
738 {
739 AddOpenCLEvent(image,event);
740 AddOpenCLEvent(filteredImage,event);
741 }
742 clEnv->library->clReleaseEvent(event);
743 }
744 }
745
746 }
747
748 outputReady=MagickTrue;
749
750cleanup:
751 OpenCLLogException(__FUNCTION__,__LINE__,exception);
752
753 if (imageBuffer != (cl_mem) NULL)
754 clEnv->library->clReleaseMemObject(imageBuffer);
755 if (filteredImageBuffer != (cl_mem) NULL)
756 clEnv->library->clReleaseMemObject(filteredImageBuffer);
757 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
758 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
759 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
760 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
761 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
762 if (kernel!=NULL) DestroyKernelInfo(kernel);
763 if ((outputReady == MagickFalse) && (filteredImage != NULL))
764 filteredImage=(Image *) DestroyImage(filteredImage);
765 return(filteredImage);
766}
767
768MagickPrivate Image* AccelerateBlurImage(const Image *image,
769 const ChannelType channel,const double radius,const double sigma,
770 ExceptionInfo *exception)
771{
772 Image
773 *filteredImage;
774
775 assert(image != NULL);
776 assert(exception != (ExceptionInfo *) NULL);
777
778 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
779 (checkAccelerateCondition(image, channel) == MagickFalse))
780 return NULL;
781
782 filteredImage=ComputeBlurImage(image, channel, radius, sigma, exception);
783 return(filteredImage);
784}
785
786/*
787%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
788% %
789% %
790% %
791% A c c e l e r a t e C o m p o s i t e I m a g e %
792% %
793% %
794% %
795%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
796*/
797
798static MagickBooleanType LaunchCompositeKernel(const Image *image,
799 MagickCLEnv clEnv,cl_command_queue queue,cl_mem imageBuffer,
800 const unsigned int inputWidth,const unsigned int inputHeight,
801 const unsigned int inputMatte,const ChannelType channel,
802 const CompositeOperator compose,const cl_mem compositeImageBuffer,
803 const unsigned int compositeWidth,const unsigned int compositeHeight,
804 const unsigned int compositeMatte,const float destination_dissolve,
805 const float source_dissolve)
806{
807 cl_int
808 clStatus;
809
810 cl_kernel
811 compositeKernel;
812
813 cl_event
814 event;
815
816 cl_uint
817 event_count;
818
819 cl_event
820 *events;
821
822 int
823 k;
824
825 size_t
826 global_work_size[2],
827 local_work_size[2];
828
829 unsigned int
830 composeOp;
831
832 compositeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
833 "Composite");
834
835 k = 0;
836 clStatus = clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(cl_mem), (void*)&imageBuffer);
837 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&inputWidth);
838 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&inputHeight);
839 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&inputMatte);
840 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(cl_mem), (void*)&compositeImageBuffer);
841 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&compositeWidth);
842 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&compositeHeight);
843 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&compositeMatte);
844 composeOp = (unsigned int)compose;
845 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(unsigned int), (void*)&composeOp);
846 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(ChannelType), (void*)&channel);
847 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(float), (void*)&destination_dissolve);
848 clStatus |= clEnv->library->clSetKernelArg(compositeKernel, k++, sizeof(float), (void*)&source_dissolve);
849
850 if (clStatus != CL_SUCCESS)
851 return MagickFalse;
852
853 local_work_size[0] = 64;
854 local_work_size[1] = 1;
855
856 global_work_size[0] = padGlobalWorkgroupSizeToLocalWorkgroupSize(inputWidth,
857 (unsigned int)local_work_size[0]);
858 global_work_size[1] = inputHeight;
859 events=GetOpenCLEvents(image,&event_count);
860 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL,
861 global_work_size, local_work_size, event_count, events, &event);
862 events=(cl_event *) RelinquishMagickMemory(events);
863 if (clStatus == CL_SUCCESS)
864 AddOpenCLEvent(image,event);
865 clEnv->library->clReleaseEvent(event);
866
867 RelinquishOpenCLKernel(clEnv, compositeKernel);
868
869 return((clStatus == CL_SUCCESS) ? MagickTrue : MagickFalse);
870}
871
872static MagickBooleanType ComputeCompositeImage(Image *image,
873 const ChannelType channel, const CompositeOperator compose,
874 const Image *compositeImage, const ssize_t magick_unused(x_offset),
875 const ssize_t magick_unused(y_offset), const float destination_dissolve,
876 const float source_dissolve, ExceptionInfo *exception)
877{
878 cl_command_queue
879 queue;
880
881 cl_context
882 context;
883
884 cl_mem
885 compositeImageBuffer,
886 imageBuffer;
887
888 MagickBooleanType
889 outputReady,
890 status;
891
892 MagickCLEnv
893 clEnv;
894
895 magick_unreferenced(x_offset);
896 magick_unreferenced(y_offset);
897
898 status = MagickFalse;
899 outputReady = MagickFalse;
900 imageBuffer = NULL;
901 compositeImageBuffer = NULL;
902
903 clEnv = GetDefaultOpenCLEnv();
904 context = GetOpenCLContext(clEnv);
905 queue = AcquireOpenCLCommandQueue(clEnv);
906
907 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
908 if (imageBuffer == (cl_mem) NULL)
909 {
910 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
911 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
912 goto cleanup;
913 }
914
915 compositeImageBuffer = GetAuthenticOpenCLBuffer(compositeImage,exception);
916 if (compositeImageBuffer == (cl_mem) NULL)
917 {
918 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
919 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
920 goto cleanup;
921 }
922
923 status = LaunchCompositeKernel(image,clEnv, queue, imageBuffer,
924 (unsigned int)image->columns,
925 (unsigned int)image->rows,
926 (unsigned int)image->matte,
927 channel, compose, compositeImageBuffer,
928 (unsigned int)compositeImage->columns,
929 (unsigned int)compositeImage->rows,
930 (unsigned int)compositeImage->matte,
931 destination_dissolve, source_dissolve);
932
933 if (status == MagickFalse)
934 goto cleanup;
935
936 outputReady = MagickTrue;
937
938cleanup:
939
940 if (imageBuffer != (cl_mem) NULL)
941 clEnv->library->clReleaseMemObject(imageBuffer);
942 if (compositeImageBuffer != (cl_mem) NULL)
943 clEnv->library->clReleaseMemObject(compositeImageBuffer);
944 if (queue != NULL)
945 RelinquishOpenCLCommandQueue(clEnv, queue);
946
947 return(outputReady);
948}
949
950MagickPrivate MagickBooleanType AccelerateCompositeImage(Image *image,
951 const ChannelType channel, const CompositeOperator compose,
952 const Image *composite, const ssize_t x_offset, const ssize_t y_offset,
953 const float destination_dissolve, const float source_dissolve,
954 ExceptionInfo *exception)
955{
956 MagickBooleanType
957 status;
958
959 assert(image != NULL);
960 assert(exception != (ExceptionInfo *)NULL);
961
962 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
963 (checkAccelerateCondition(image, channel) == MagickFalse))
964 return(MagickFalse);
965
966 /* only support zero offset and
967 images with the size for now */
968 if (x_offset != 0
969 || y_offset != 0
970 || image->columns != composite->columns
971 || image->rows != composite->rows)
972 return MagickFalse;
973
974 switch (compose) {
975 case ColorDodgeCompositeOp:
976 case BlendCompositeOp:
977 break;
978 default:
979 /* unsupported compose operator, quit */
980 return MagickFalse;
981 };
982
983 status = ComputeCompositeImage(image, channel, compose, composite,
984 x_offset, y_offset, destination_dissolve, source_dissolve, exception);
985
986 return(status);
987}
988
989/*
990%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
991% %
992% %
993% %
994% A c c e l e r a t e C o n t r a s t I m a g e %
995% %
996% %
997% %
998%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
999*/
1000
1001static MagickBooleanType ComputeContrastImage(Image *image,
1002 const MagickBooleanType sharpen,ExceptionInfo *exception)
1003{
1004 cl_command_queue
1005 queue;
1006
1007 cl_context
1008 context;
1009
1010 cl_int
1011 clStatus;
1012
1013 cl_kernel
1014 filterKernel;
1015
1016 cl_event
1017 event;
1018
1019 cl_mem
1020 imageBuffer;
1021
1022 cl_uint
1023 event_count;
1024
1025 cl_event
1026 *events;
1027
1028 MagickBooleanType
1029 outputReady;
1030
1031 MagickCLEnv
1032 clEnv;
1033
1034 size_t
1035 global_work_size[2];
1036
1037 unsigned int
1038 i,
1039 uSharpen;
1040
1041 outputReady = MagickFalse;
1042 clEnv = NULL;
1043 context = NULL;
1044 imageBuffer = NULL;
1045 filterKernel = NULL;
1046 queue = NULL;
1047
1048 clEnv = GetDefaultOpenCLEnv();
1049 context = GetOpenCLContext(clEnv);
1050
1051 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
1052 if (imageBuffer == (cl_mem) NULL)
1053 {
1054 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1055 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
1056 goto cleanup;
1057 }
1058
1059 filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast");
1060 if (filterKernel == NULL)
1061 {
1062 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1063 goto cleanup;
1064 }
1065
1066 i = 0;
1067 clStatus=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1068
1069 uSharpen = (sharpen == MagickFalse)?0:1;
1070 clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen);
1071 if (clStatus != CL_SUCCESS)
1072 {
1073 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1074 goto cleanup;
1075 }
1076
1077 global_work_size[0] = image->columns;
1078 global_work_size[1] = image->rows;
1079 /* launch the kernel */
1080 queue = AcquireOpenCLCommandQueue(clEnv);
1081 events=GetOpenCLEvents(image,&event_count);
1082 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
1083 events=(cl_event *) RelinquishMagickMemory(events);
1084 if (clStatus != CL_SUCCESS)
1085 {
1086 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1087 goto cleanup;
1088 }
1089 if (RecordProfileData(clEnv,ContrastKernel,event) == MagickFalse)
1090 AddOpenCLEvent(image,event);
1091 clEnv->library->clReleaseEvent(event);
1092 outputReady=MagickTrue;
1093
1094cleanup:
1095 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1096
1097
1098 if (imageBuffer != (cl_mem) NULL)
1099 clEnv->library->clReleaseMemObject(imageBuffer);
1100 if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel);
1101 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
1102 return(outputReady);
1103}
1104
1105MagickPrivate MagickBooleanType AccelerateContrastImage(Image *image,
1106 const MagickBooleanType sharpen,ExceptionInfo *exception)
1107{
1108 MagickBooleanType
1109 status;
1110
1111 assert(image != NULL);
1112 assert(exception != (ExceptionInfo *) NULL);
1113
1114 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1115 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
1116 return(MagickFalse);
1117
1118 status = ComputeContrastImage(image,sharpen,exception);
1119 return(status);
1120}
1121
1122/*
1123%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1124% %
1125% %
1126% %
1127% A c c e l e r a t e C o n t r a s t S t r e t c h I m a g e %
1128% %
1129% %
1130% %
1131%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1132*/
1133
1134static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
1135 cl_command_queue queue,cl_mem imageBuffer,cl_mem histogramBuffer,
1136 Image *image,const ChannelType channel,ExceptionInfo *exception)
1137{
1138 MagickBooleanType
1139 outputReady;
1140
1141 cl_event
1142 event;
1143
1144 cl_int
1145 clStatus,
1146 colorspace,
1147 method;
1148
1149 cl_kernel
1150 histogramKernel;
1151
1152 cl_uint
1153 event_count;
1154
1155 cl_event
1156 *events;
1157
1158 ssize_t
1159 i;
1160
1161 size_t
1162 global_work_size[2];
1163
1164 histogramKernel = NULL;
1165
1166 outputReady = MagickFalse;
1167 method = image->intensity;
1168 colorspace = image->colorspace;
1169
1170 /* get the OpenCL kernel */
1171 histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram");
1172 if (histogramKernel == NULL)
1173 {
1174 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1175 goto cleanup;
1176 }
1177
1178 /* set the kernel arguments */
1179 i = 0;
1180 clStatus=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1181 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(ChannelType),&channel);
1182 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&method);
1183 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_int),&colorspace);
1184 clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer);
1185 if (clStatus != CL_SUCCESS)
1186 {
1187 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1188 goto cleanup;
1189 }
1190
1191 /* launch the kernel */
1192 global_work_size[0] = image->columns;
1193 global_work_size[1] = image->rows;
1194
1195 events=GetOpenCLEvents(image,&event_count);
1196 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, histogramKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
1197 events=(cl_event *) RelinquishMagickMemory(events);
1198
1199 if (clStatus != CL_SUCCESS)
1200 {
1201 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1202 goto cleanup;
1203 }
1204 if (RecordProfileData(clEnv,HistogramKernel,event) == MagickFalse)
1205 AddOpenCLEvent(image,event);
1206 clEnv->library->clReleaseEvent(event);
1207
1208 outputReady = MagickTrue;
1209
1210cleanup:
1211 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1212
1213 if (histogramKernel!=NULL)
1214 RelinquishOpenCLKernel(clEnv, histogramKernel);
1215
1216 return(outputReady);
1217}
1218
1219MagickPrivate MagickBooleanType ComputeContrastStretchImageChannel(Image *image,
1220 const ChannelType channel,const double black_point,const double white_point,
1221 ExceptionInfo *exception)
1222{
1223#define ContrastStretchImageTag "ContrastStretch/Image"
1224#define MaxRange(color) ((MagickRealType) ScaleQuantumToMap((Quantum) (color)))
1225 cl_command_queue
1226 queue;
1227
1228 cl_context
1229 context;
1230
1231 cl_int
1232 clStatus;
1233
1234 cl_mem
1235 histogramBuffer,
1236 imageBuffer,
1237 stretchMapBuffer;
1238
1239 cl_kernel
1240 histogramKernel,
1241 stretchKernel;
1242
1243 cl_event
1244 event;
1245
1246 cl_uint
1247 event_count;
1248
1249 cl_uint4
1250 *histogram;
1251
1252 cl_event
1253 *events;
1254
1255 double
1256 intensity;
1257
1258 cl_float4
1259 black,
1260 white;
1261
1262 MagickBooleanType
1263 outputReady,
1264 status;
1265
1266 MagickCLEnv
1267 clEnv;
1268
1269 MagickSizeType
1270 length;
1271
1272 PixelPacket
1273 *stretch_map;
1274
1275 ssize_t
1276 i;
1277
1278 size_t
1279 global_work_size[2];
1280
1281 histogram=NULL;
1282 stretch_map=NULL;
1283 imageBuffer = NULL;
1284 histogramBuffer = NULL;
1285 stretchMapBuffer = NULL;
1286 histogramKernel = NULL;
1287 stretchKernel = NULL;
1288 context = NULL;
1289 queue = NULL;
1290 outputReady = MagickFalse;
1291
1292
1293 assert(image != (Image *) NULL);
1294 assert(image->signature == MagickCoreSignature);
1295 if (IsEventLogging() != MagickFalse)
1296 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
1297
1298 /* exception=(&image->exception); */
1299
1300 /*
1301 * initialize opencl env
1302 */
1303 clEnv = GetDefaultOpenCLEnv();
1304 context = GetOpenCLContext(clEnv);
1305 queue = AcquireOpenCLCommandQueue(clEnv);
1306
1307 /*
1308 Allocate and initialize histogram arrays.
1309 */
1310 length = (MaxMap+1);
1311 histogram=(cl_uint4 *) AcquireQuantumMemory(length, sizeof(*histogram));
1312
1313 if (histogram == (cl_uint4 *) NULL)
1314 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed", image->filename);
1315
1316 /* reset histogram */
1317 (void) memset(histogram,0,length*sizeof(*histogram));
1318
1319 /*
1320 if (SetImageGray(image,exception) != MagickFalse)
1321 (void) SetImageColorspace(image,GRAYColorspace);
1322 */
1323
1324 status=MagickTrue;
1325
1326 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
1327 if (imageBuffer == (cl_mem) NULL)
1328 {
1329 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1330 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
1331 goto cleanup;
1332 }
1333
1334 /* create a CL buffer for histogram */
1335 histogramBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, length * sizeof(cl_uint4), histogram, &clStatus);
1336 if (clStatus != CL_SUCCESS)
1337 {
1338 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1339 goto cleanup;
1340 }
1341
1342 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, channel, exception);
1343 if (status == MagickFalse)
1344 goto cleanup;
1345
1346 /* this blocks, should be fixed it in the future */
1347 events=GetOpenCLEvents(image,&event_count);
1348 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), event_count, events, NULL, &clStatus);
1349 events=(cl_event *) RelinquishMagickMemory(events);
1350 if (clStatus != CL_SUCCESS)
1351 {
1352 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
1353 goto cleanup;
1354 }
1355
1356 /* unmap, don't block gpu to use this buffer again. */
1357 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
1358 if (clStatus != CL_SUCCESS)
1359 {
1360 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1361 goto cleanup;
1362 }
1363
1364 /* CPU stuff */
1365 /*
1366 Find the histogram boundaries by locating the black/white levels.
1367 */
1368 black.z=0.0;
1369 white.z=MaxRange(QuantumRange);
1370 if ((channel & RedChannel) != 0)
1371 {
1372 intensity=0.0;
1373 for (i=0; i <= (ssize_t) MaxMap; i++)
1374 {
1375 intensity+=histogram[i].s[2];
1376 if (intensity > black_point)
1377 break;
1378 }
1379 black.z=(MagickRealType) i;
1380 intensity=0.0;
1381 for (i=(ssize_t) MaxMap; i != 0; i--)
1382 {
1383 intensity+=histogram[i].s[2];
1384 if (intensity > ((double) image->columns*image->rows-white_point))
1385 break;
1386 }
1387 white.z=(MagickRealType) i;
1388 }
1389 black.y=0.0;
1390 white.y=MaxRange(QuantumRange);
1391 if ((channel & GreenChannel) != 0)
1392 {
1393 intensity=0.0;
1394 for (i=0; i <= (ssize_t) MaxMap; i++)
1395 {
1396 intensity+=histogram[i].s[2];
1397 if (intensity > black_point)
1398 break;
1399 }
1400 black.y=(MagickRealType) i;
1401 intensity=0.0;
1402 for (i=(ssize_t) MaxMap; i != 0; i--)
1403 {
1404 intensity+=histogram[i].s[2];
1405 if (intensity > ((double) image->columns*image->rows-white_point))
1406 break;
1407 }
1408 white.y=(MagickRealType) i;
1409 }
1410 black.x=0.0;
1411 white.x=MaxRange(QuantumRange);
1412 if ((channel & BlueChannel) != 0)
1413 {
1414 intensity=0.0;
1415 for (i=0; i <= (ssize_t) MaxMap; i++)
1416 {
1417 intensity+=histogram[i].s[2];
1418 if (intensity > black_point)
1419 break;
1420 }
1421 black.x=(MagickRealType) i;
1422 intensity=0.0;
1423 for (i=(ssize_t) MaxMap; i != 0; i--)
1424 {
1425 intensity+=histogram[i].s[2];
1426 if (intensity > ((double) image->columns*image->rows-white_point))
1427 break;
1428 }
1429 white.x=(MagickRealType) i;
1430 }
1431 black.w=0.0;
1432 white.w=MaxRange(QuantumRange);
1433 if ((channel & OpacityChannel) != 0)
1434 {
1435 intensity=0.0;
1436 for (i=0; i <= (ssize_t) MaxMap; i++)
1437 {
1438 intensity+=histogram[i].s[2];
1439 if (intensity > black_point)
1440 break;
1441 }
1442 black.w=(MagickRealType) i;
1443 intensity=0.0;
1444 for (i=(ssize_t) MaxMap; i != 0; i--)
1445 {
1446 intensity+=histogram[i].s[2];
1447 if (intensity > ((double) image->columns*image->rows-white_point))
1448 break;
1449 }
1450 white.w=(MagickRealType) i;
1451 }
1452 /*
1453 black.index=0.0;
1454 white.index=MaxRange(QuantumRange);
1455 if (((channel & IndexChannel) != 0) && (image->colorspace == CMYKColorspace))
1456 {
1457 intensity=0.0;
1458 for (i=0; i <= (ssize_t) MaxMap; i++)
1459 {
1460 intensity+=histogram[i].index;
1461 if (intensity > black_point)
1462 break;
1463 }
1464 black.index=(MagickRealType) i;
1465 intensity=0.0;
1466 for (i=(ssize_t) MaxMap; i != 0; i--)
1467 {
1468 intensity+=histogram[i].index;
1469 if (intensity > ((double) image->columns*image->rows-white_point))
1470 break;
1471 }
1472 white.index=(MagickRealType) i;
1473 }
1474 */
1475
1476
1477 stretch_map=(PixelPacket *) AcquireQuantumMemory(length,
1478 sizeof(*stretch_map));
1479
1480 if (stretch_map == (PixelPacket *) NULL)
1481 ThrowBinaryException(ResourceLimitError,"MemoryAllocationFailed",
1482 image->filename);
1483
1484 /*
1485 Stretch the histogram to create the stretched image mapping.
1486 */
1487 (void) memset(stretch_map,0,length*sizeof(*stretch_map));
1488 for (i=0; i <= (ssize_t) MaxMap; i++)
1489 {
1490 if ((channel & RedChannel) != 0)
1491 {
1492 if (i < (ssize_t) black.z)
1493 stretch_map[i].red=(Quantum) 0;
1494 else
1495 if (i > (ssize_t) white.z)
1496 stretch_map[i].red=QuantumRange;
1497 else
1498 if (black.z != white.z)
1499 stretch_map[i].red=ScaleMapToQuantum((MagickRealType) (MaxMap*
1500 (i-black.z)/(white.z-black.z)));
1501 }
1502 if ((channel & GreenChannel) != 0)
1503 {
1504 if (i < (ssize_t) black.y)
1505 stretch_map[i].green=0;
1506 else
1507 if (i > (ssize_t) white.y)
1508 stretch_map[i].green=QuantumRange;
1509 else
1510 if (black.y != white.y)
1511 stretch_map[i].green=ScaleMapToQuantum((MagickRealType) (MaxMap*
1512 (i-black.y)/(white.y-black.y)));
1513 }
1514 if ((channel & BlueChannel) != 0)
1515 {
1516 if (i < (ssize_t) black.x)
1517 stretch_map[i].blue=0;
1518 else
1519 if (i > (ssize_t) white.x)
1520 stretch_map[i].blue= QuantumRange;
1521 else
1522 if (black.x != white.x)
1523 stretch_map[i].blue=ScaleMapToQuantum((MagickRealType) (MaxMap*
1524 (i-black.x)/(white.x-black.x)));
1525 }
1526 if ((channel & OpacityChannel) != 0)
1527 {
1528 if (i < (ssize_t) black.w)
1529 stretch_map[i].opacity=0;
1530 else
1531 if (i > (ssize_t) white.w)
1532 stretch_map[i].opacity=QuantumRange;
1533 else
1534 if (black.w != white.w)
1535 stretch_map[i].opacity=ScaleMapToQuantum((MagickRealType) (MaxMap*
1536 (i-black.w)/(white.w-black.w)));
1537 }
1538 /*
1539 if (((channel & IndexChannel) != 0) &&
1540 (image->colorspace == CMYKColorspace))
1541 {
1542 if (i < (ssize_t) black.index)
1543 stretch_map[i].index=0;
1544 else
1545 if (i > (ssize_t) white.index)
1546 stretch_map[i].index=QuantumRange;
1547 else
1548 if (black.index != white.index)
1549 stretch_map[i].index=ScaleMapToQuantum((MagickRealType) (MaxMap*
1550 (i-black.index)/(white.index-black.index)));
1551 }
1552 */
1553 }
1554
1555 /*
1556 Stretch the image.
1557 */
1558 if (((channel & OpacityChannel) != 0) || (((channel & IndexChannel) != 0) &&
1559 (image->colorspace == CMYKColorspace)))
1560 image->storage_class=DirectClass;
1561 if (image->storage_class == PseudoClass)
1562 {
1563 /*
1564 Stretch colormap.
1565 */
1566 for (i=0; i < (ssize_t) image->colors; i++)
1567 {
1568 if ((channel & RedChannel) != 0)
1569 {
1570 if (black.z != white.z)
1571 image->colormap[i].red=stretch_map[
1572 ScaleQuantumToMap(image->colormap[i].red)].red;
1573 }
1574 if ((channel & GreenChannel) != 0)
1575 {
1576 if (black.y != white.y)
1577 image->colormap[i].green=stretch_map[
1578 ScaleQuantumToMap(image->colormap[i].green)].green;
1579 }
1580 if ((channel & BlueChannel) != 0)
1581 {
1582 if (black.x != white.x)
1583 image->colormap[i].blue=stretch_map[
1584 ScaleQuantumToMap(image->colormap[i].blue)].blue;
1585 }
1586 if ((channel & OpacityChannel) != 0)
1587 {
1588 if (black.w != white.w)
1589 image->colormap[i].opacity=stretch_map[
1590 ScaleQuantumToMap(image->colormap[i].opacity)].opacity;
1591 }
1592 }
1593 }
1594
1595
1596 /* create a CL buffer for stretch_map */
1597 stretchMapBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, length, stretch_map, &clStatus);
1598 if (clStatus != CL_SUCCESS)
1599 {
1600 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1601 goto cleanup;
1602 }
1603
1604 /* get the OpenCL kernel */
1605 stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ContrastStretch");
1606 if (stretchKernel == NULL)
1607 {
1608 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1609 goto cleanup;
1610 }
1611
1612 /* set the kernel arguments */
1613 i = 0;
1614 clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1615 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(ChannelType),&channel);
1616 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&stretchMapBuffer);
1617 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_float4),&white);
1618 clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_float4),&black);
1619 if (clStatus != CL_SUCCESS)
1620 {
1621 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1622 goto cleanup;
1623 }
1624
1625 /* launch the kernel */
1626 global_work_size[0] = image->columns;
1627 global_work_size[1] = image->rows;
1628
1629 events=GetOpenCLEvents(image,&event_count);
1630 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
1631 events=(cl_event *) RelinquishMagickMemory(events);
1632
1633 if (clStatus != CL_SUCCESS)
1634 {
1635 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1636 goto cleanup;
1637 }
1638
1639 if (RecordProfileData(clEnv,ContrastStretchKernel,event) == MagickFalse)
1640 AddOpenCLEvent(image, event);
1641 clEnv->library->clReleaseEvent(event);
1642
1643 outputReady=MagickTrue;
1644
1645cleanup:
1646 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1647
1648 if (imageBuffer != (cl_mem) NULL)
1649 clEnv->library->clReleaseMemObject(imageBuffer);
1650
1651 if (stretchMapBuffer!=NULL)
1652 clEnv->library->clReleaseMemObject(stretchMapBuffer);
1653 if (stretch_map!=NULL)
1654 stretch_map=(PixelPacket *) RelinquishMagickMemory(stretch_map);
1655
1656
1657 if (histogramBuffer!=NULL)
1658 clEnv->library->clReleaseMemObject(histogramBuffer);
1659 if (histogram!=NULL)
1660 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
1661
1662
1663 if (histogramKernel!=NULL)
1664 RelinquishOpenCLKernel(clEnv, histogramKernel);
1665 if (stretchKernel!=NULL)
1666 RelinquishOpenCLKernel(clEnv, stretchKernel);
1667
1668 if (queue != NULL)
1669 RelinquishOpenCLCommandQueue(clEnv, queue);
1670
1671 return(outputReady);
1672}
1673
1674MagickPrivate MagickBooleanType AccelerateContrastStretchImageChannel(
1675 Image *image,const ChannelType channel,const double black_point,
1676 const double white_point,ExceptionInfo *exception)
1677{
1678 MagickBooleanType
1679 status;
1680
1681 assert(image != NULL);
1682 assert(exception != (ExceptionInfo *) NULL);
1683
1684 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1685 (checkAccelerateCondition(image, channel) == MagickFalse) ||
1686 (checkHistogramCondition(image, channel) == MagickFalse))
1687 return(MagickFalse);
1688
1689 status=ComputeContrastStretchImageChannel(image,channel, black_point, white_point, exception);
1690 return(status);
1691}
1692
1693/*
1694%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1695% %
1696% %
1697% %
1698% A c c e l e r a t e C o n v o l v e I m a g e %
1699% %
1700% %
1701% %
1702%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1703*/
1704
1705static Image *ComputeConvolveImage(const Image* image,
1706 const ChannelType channel,const KernelInfo *kernel,ExceptionInfo *exception)
1707{
1708 cl_command_queue
1709 queue;
1710
1711 cl_context
1712 context;
1713
1714 cl_kernel
1715 clkernel;
1716
1717 cl_event
1718 event;
1719
1720 cl_int
1721 clStatus;
1722
1723 cl_mem
1724 convolutionKernel,
1725 filteredImageBuffer,
1726 imageBuffer;
1727
1728 cl_uint
1729 event_count;
1730
1731 cl_ulong
1732 deviceLocalMemorySize;
1733
1734 cl_event
1735 *events;
1736
1737 float
1738 *kernelBufferPtr;
1739
1740 Image
1741 *filteredImage;
1742
1743 MagickBooleanType
1744 outputReady;
1745
1746 MagickCLEnv
1747 clEnv;
1748
1749 size_t
1750 global_work_size[3],
1751 localGroupSize[3],
1752 localMemoryRequirement;
1753
1754 unsigned
1755 kernelSize;
1756
1757 unsigned int
1758 filterHeight,
1759 filterWidth,
1760 i,
1761 imageHeight,
1762 imageWidth,
1763 matte;
1764
1765 /* intialize all CL objects to NULL */
1766 context = NULL;
1767 imageBuffer = NULL;
1768 filteredImageBuffer = NULL;
1769 convolutionKernel = NULL;
1770 clkernel = NULL;
1771 queue = NULL;
1772
1773 filteredImage = NULL;
1774 outputReady = MagickFalse;
1775
1776 clEnv = GetDefaultOpenCLEnv();
1777
1778 context = GetOpenCLContext(clEnv);
1779
1780 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
1781 if (filteredImage == (Image *) NULL)
1782 goto cleanup;
1783
1784 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
1785 if (imageBuffer == (cl_mem) NULL)
1786 {
1787 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1788 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
1789 goto cleanup;
1790 }
1791 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,exception);
1792 if (filteredImageBuffer == (cl_mem) NULL)
1793 {
1794 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
1795 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
1796 goto cleanup;
1797 }
1798
1799 kernelSize = (unsigned int) (kernel->width * kernel->height);
1800 convolutionKernel = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus);
1801 if (clStatus != CL_SUCCESS)
1802 {
1803 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
1804 goto cleanup;
1805 }
1806
1807 queue = AcquireOpenCLCommandQueue(clEnv);
1808
1809 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, convolutionKernel, CL_TRUE, CL_MAP_WRITE, 0, kernelSize * sizeof(float)
1810 , 0, NULL, NULL, &clStatus);
1811 if (clStatus != CL_SUCCESS)
1812 {
1813 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
1814 goto cleanup;
1815 }
1816 for (i = 0; i < kernelSize; i++)
1817 {
1818 kernelBufferPtr[i] = (float) kernel->values[i];
1819 }
1820 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
1821 if (clStatus != CL_SUCCESS)
1822 {
1823 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
1824 goto cleanup;
1825 }
1826
1827 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
1828
1829 /* Compute the local memory requirement for a 16x16 workgroup.
1830 If it's larger than 16k, reduce the workgroup size to 8x8 */
1831 localGroupSize[0] = 16;
1832 localGroupSize[1] = 16;
1833 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
1834 + kernel->width*kernel->height*sizeof(float);
1835
1836 if (localMemoryRequirement > deviceLocalMemorySize)
1837 {
1838 localGroupSize[0] = 8;
1839 localGroupSize[1] = 8;
1840 localMemoryRequirement = (localGroupSize[0]+kernel->width-1) * (localGroupSize[1]+kernel->height-1) * sizeof(CLPixelPacket)
1841 + kernel->width*kernel->height*sizeof(float);
1842 }
1843 if (localMemoryRequirement <= deviceLocalMemorySize)
1844 {
1845 /* get the OpenCL kernel */
1846 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized");
1847 if (clkernel == NULL)
1848 {
1849 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1850 goto cleanup;
1851 }
1852
1853 /* set the kernel arguments */
1854 i = 0;
1855 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1856 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1857 imageWidth = (unsigned int) image->columns;
1858 imageHeight = (unsigned int) image->rows;
1859 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
1860 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
1861 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
1862 filterWidth = (unsigned int) kernel->width;
1863 filterHeight = (unsigned int) kernel->height;
1864 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
1865 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
1866 matte = (image->matte==MagickTrue)?1:0;
1867 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
1868 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
1869 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, (localGroupSize[0] + kernel->width-1)*(localGroupSize[1] + kernel->height-1)*sizeof(CLPixelPacket),NULL);
1870 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL);
1871 if (clStatus != CL_SUCCESS)
1872 {
1873 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1874 goto cleanup;
1875 }
1876
1877 /* pad the global size to a multiple of the local work size dimension */
1878 global_work_size[0] = ((image->columns + localGroupSize[0] - 1)/localGroupSize[0] ) * localGroupSize[0] ;
1879 global_work_size[1] = ((image->rows + localGroupSize[1] - 1)/localGroupSize[1]) * localGroupSize[1];
1880
1881 /* launch the kernel */
1882 events = GetOpenCLEvents(image, &event_count);
1883 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, event_count, events, &event);
1884 events=(cl_event *) RelinquishMagickMemory(events);
1885 if (clStatus != CL_SUCCESS)
1886 {
1887 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1888 goto cleanup;
1889 }
1890 if (RecordProfileData(clEnv,ConvolveKernel,event) == MagickFalse)
1891 {
1892 AddOpenCLEvent(image, event);
1893 AddOpenCLEvent(filteredImage, event);
1894 }
1895 clEnv->library->clReleaseEvent(event);
1896 }
1897 else
1898 {
1899 /* get the OpenCL kernel */
1900 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve");
1901 if (clkernel == NULL)
1902 {
1903 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
1904 goto cleanup;
1905 }
1906
1907 /* set the kernel arguments */
1908 i = 0;
1909 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
1910 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
1911 imageWidth = (unsigned int) image->columns;
1912 imageHeight = (unsigned int) image->rows;
1913 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageWidth);
1914 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&imageHeight);
1915 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&convolutionKernel);
1916 filterWidth = (unsigned int) kernel->width;
1917 filterHeight = (unsigned int) kernel->height;
1918 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterWidth);
1919 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&filterHeight);
1920 matte = (image->matte==MagickTrue)?1:0;
1921 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&matte);
1922 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
1923 if (clStatus != CL_SUCCESS)
1924 {
1925 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
1926 goto cleanup;
1927 }
1928
1929 localGroupSize[0] = 8;
1930 localGroupSize[1] = 8;
1931 global_work_size[0] = (image->columns + (localGroupSize[0]-1))/localGroupSize[0] * localGroupSize[0];
1932 global_work_size[1] = (image->rows + (localGroupSize[1]-1))/localGroupSize[1] * localGroupSize[1];
1933 events=GetOpenCLEvents(image,&event_count);
1934 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, event_count, events, &event);
1935 events=(cl_event *) RelinquishMagickMemory(events);
1936
1937 if (clStatus != CL_SUCCESS)
1938 {
1939 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
1940 goto cleanup;
1941 }
1942 if (RecordProfileData(clEnv,ConvolveKernel,event) == MagickFalse)
1943 {
1944 AddOpenCLEvent(image,event);
1945 AddOpenCLEvent(filteredImage,event);
1946 }
1947 clEnv->library->clReleaseEvent(event);
1948 }
1949
1950 outputReady = MagickTrue;
1951
1952cleanup:
1953 OpenCLLogException(__FUNCTION__,__LINE__,exception);
1954
1955 if (imageBuffer != (cl_mem) NULL)
1956 clEnv->library->clReleaseMemObject(imageBuffer);
1957
1958 if (filteredImageBuffer != (cl_mem) NULL)
1959 clEnv->library->clReleaseMemObject(filteredImageBuffer);
1960
1961 if (convolutionKernel != NULL)
1962 clEnv->library->clReleaseMemObject(convolutionKernel);
1963
1964 if (clkernel != NULL)
1965 RelinquishOpenCLKernel(clEnv, clkernel);
1966
1967 if (queue != NULL)
1968 RelinquishOpenCLCommandQueue(clEnv, queue);
1969
1970 if ((outputReady == MagickFalse) && (filteredImage != NULL))
1971 filteredImage=(Image *) DestroyImage(filteredImage);
1972
1973 return(filteredImage);
1974}
1975
1976MagickPrivate Image *AccelerateConvolveImageChannel(const Image *image,
1977 const ChannelType channel,const KernelInfo *kernel,ExceptionInfo *exception)
1978{
1979 Image
1980 *filteredImage;
1981
1982 assert(image != NULL);
1983 assert(kernel != (KernelInfo *) NULL);
1984 assert(exception != (ExceptionInfo *) NULL);
1985
1986 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
1987 (checkAccelerateCondition(image, channel) == MagickFalse))
1988 return NULL;
1989
1990 filteredImage=ComputeConvolveImage(image, channel, kernel, exception);
1991 return(filteredImage);
1992}
1993
1994/*
1995%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1996% %
1997% %
1998% %
1999% A c c e l e r a t e D e s p e c k l e I m a g e %
2000% %
2001% %
2002% %
2003%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2004*/
2005
2006static Image *ComputeDespeckleImage(const Image *image,
2007 ExceptionInfo*exception)
2008{
2009 static const int
2010 X[4] = {0, 1, 1,-1},
2011 Y[4] = {1, 0, 1, 1};
2012
2013 cl_command_queue
2014 queue;
2015
2016 cl_context
2017 context;
2018
2019 cl_int
2020 clStatus;
2021
2022 cl_kernel
2023 hullPass1,
2024 hullPass2;
2025
2026 cl_event
2027 event;
2028
2029 cl_mem
2030 filteredImageBuffer,
2031 imageBuffer,
2032 tempImageBuffer[2];
2033
2034 cl_uint
2035 event_count;
2036
2037 cl_event
2038 *events;
2039
2040 Image
2041 *filteredImage;
2042
2043 int
2044 k,
2045 matte;
2046
2047 MagickBooleanType
2048 outputReady;
2049
2050 MagickCLEnv
2051 clEnv;
2052
2053 size_t
2054 global_work_size[2];
2055
2056 unsigned int
2057 imageHeight,
2058 imageWidth;
2059
2060 outputReady = MagickFalse;
2061 clEnv = NULL;
2062 filteredImage = NULL;
2063 context = NULL;
2064 imageBuffer = NULL;
2065 filteredImageBuffer = NULL;
2066 hullPass1 = NULL;
2067 hullPass2 = NULL;
2068 queue = NULL;
2069 tempImageBuffer[0] = tempImageBuffer[1] = NULL;
2070 clEnv = GetDefaultOpenCLEnv();
2071 context = GetOpenCLContext(clEnv);
2072 queue = AcquireOpenCLCommandQueue(clEnv);
2073 events = NULL;
2074
2075 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
2076 if (filteredImage == (Image *) NULL)
2077 goto cleanup;
2078
2079 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
2080 if (imageBuffer == (cl_mem) NULL)
2081 {
2082 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2083 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
2084 goto cleanup;
2085 }
2086 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage, exception);
2087 if (filteredImageBuffer == (cl_mem) NULL)
2088 {
2089 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2090 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
2091 goto cleanup;
2092 }
2093
2094 hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1");
2095 hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2");
2096
2097 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&imageBuffer);
2098 clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1));
2099 imageWidth = (unsigned int) image->columns;
2100 clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth);
2101 imageHeight = (unsigned int) image->rows;
2102 clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight);
2103 matte = (image->matte==MagickFalse)?0:1;
2104 clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
2105 if (clStatus != CL_SUCCESS)
2106 {
2107 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2108 goto cleanup;
2109 }
2110
2111 clStatus = clEnv->library->clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1));
2112 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer);
2113 imageWidth = (unsigned int) image->columns;
2114 clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth);
2115 imageHeight = (unsigned int) image->rows;
2116 clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight);
2117 matte = (image->matte==MagickFalse)?0:1;
2118 clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
2119 if (clStatus != CL_SUCCESS)
2120 {
2121 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2122 goto cleanup;
2123 }
2124
2125
2126 global_work_size[0] = image->columns;
2127 global_work_size[1] = image->rows;
2128
2129 events=GetOpenCLEvents(image,&event_count);
2130 for (k = 0; k < 4; k++)
2131 {
2132 cl_int2 offset;
2133 int polarity;
2134
2135
2136 offset.s[0] = X[k];
2137 offset.s[1] = Y[k];
2138 polarity = 1;
2139 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
2140 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
2141 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
2142 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
2143 if (clStatus != CL_SUCCESS)
2144 {
2145 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2146 goto cleanup;
2147 }
2148 /* launch the kernel */
2149 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2150 if (clStatus != CL_SUCCESS)
2151 {
2152 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2153 goto cleanup;
2154 }
2155 RecordProfileData(clEnv,HullPass1Kernel,event);
2156 clEnv->library->clReleaseEvent(event);
2157 /* launch the kernel */
2158 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2159 if (clStatus != CL_SUCCESS)
2160 {
2161 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2162 goto cleanup;
2163 }
2164 RecordProfileData(clEnv,HullPass2Kernel,event);
2165 clEnv->library->clReleaseEvent(event);
2166
2167
2168 if (k == 0)
2169 clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
2170 offset.s[0] = -X[k];
2171 offset.s[1] = -Y[k];
2172 polarity = 1;
2173 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
2174 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
2175 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
2176 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
2177 if (clStatus != CL_SUCCESS)
2178 {
2179 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2180 goto cleanup;
2181 }
2182 /* launch the kernel */
2183 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2184 if (clStatus != CL_SUCCESS)
2185 {
2186 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2187 goto cleanup;
2188 }
2189 RecordProfileData(clEnv,HullPass1Kernel,event);
2190 clEnv->library->clReleaseEvent(event);
2191 /* launch the kernel */
2192 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2193 if (clStatus != CL_SUCCESS)
2194 {
2195 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2196 goto cleanup;
2197 }
2198 RecordProfileData(clEnv,HullPass2Kernel,event);
2199 clEnv->library->clReleaseEvent(event);
2200
2201 offset.s[0] = -X[k];
2202 offset.s[1] = -Y[k];
2203 polarity = -1;
2204 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
2205 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
2206 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
2207 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
2208 if (clStatus != CL_SUCCESS)
2209 {
2210 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2211 goto cleanup;
2212 }
2213 /* launch the kernel */
2214 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2215 if (clStatus != CL_SUCCESS)
2216 {
2217 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2218 goto cleanup;
2219 }
2220 RecordProfileData(clEnv,HullPass1Kernel,event);
2221 clEnv->library->clReleaseEvent(event);
2222 /* launch the kernel */
2223 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2224 if (clStatus != CL_SUCCESS)
2225 {
2226 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2227 goto cleanup;
2228 }
2229 RecordProfileData(clEnv,HullPass2Kernel,event);
2230 clEnv->library->clReleaseEvent(event);
2231
2232 offset.s[0] = X[k];
2233 offset.s[1] = Y[k];
2234 polarity = -1;
2235 clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
2236 clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
2237 clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
2238 clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
2239
2240 if (k == 3)
2241 clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer);
2242
2243 if (clStatus != CL_SUCCESS)
2244 {
2245 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2246 goto cleanup;
2247 }
2248 /* launch the kernel */
2249 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, event_count, events, &event);
2250 if (clStatus != CL_SUCCESS)
2251 {
2252 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2253 goto cleanup;
2254 }
2255 RecordProfileData(clEnv,HullPass1Kernel,event);
2256 clEnv->library->clReleaseEvent(event);
2257 /* launch the kernel */
2258 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, event_count, events, &event);
2259 if (clStatus != CL_SUCCESS)
2260 {
2261 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2262 goto cleanup;
2263 }
2264 if ((k == 3) && (RecordProfileData(clEnv,HullPass2Kernel,event) == MagickFalse))
2265 {
2266 AddOpenCLEvent(image,event);
2267 AddOpenCLEvent(filteredImage,event);
2268 }
2269 clEnv->library->clReleaseEvent(event);
2270 }
2271
2272 outputReady=MagickTrue;
2273
2274cleanup:
2275 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2276
2277 if (imageBuffer != (cl_mem) NULL)
2278 clEnv->library->clReleaseMemObject(imageBuffer);
2279 if (filteredImageBuffer != (cl_mem) NULL)
2280 clEnv->library->clReleaseMemObject(filteredImageBuffer);
2281 events=(cl_event *) RelinquishMagickMemory(events);
2282 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2283 for (k = 0; k < 2; k++)
2284 {
2285 if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
2286 }
2287 if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
2288 if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
2289 if ((outputReady == MagickFalse) && (filteredImage != NULL))
2290 filteredImage=(Image *) DestroyImage(filteredImage);
2291 return(filteredImage);
2292}
2293
2294MagickPrivate Image *AccelerateDespeckleImage(const Image* image,
2295 ExceptionInfo* exception)
2296{
2297 Image
2298 *filteredImage;
2299
2300 assert(image != NULL);
2301 assert(exception != (ExceptionInfo *) NULL);
2302
2303 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2304 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
2305 return NULL;
2306
2307 filteredImage=ComputeDespeckleImage(image,exception);
2308 return(filteredImage);
2309}
2310
2311/*
2312%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2313% %
2314% %
2315% %
2316% A c c e l e r a t e E q u a l i z e I m a g e %
2317% %
2318% %
2319% %
2320%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2321*/
2322
2323MagickPrivate MagickBooleanType ComputeEqualizeImage(Image *image,
2324 const ChannelType channel,ExceptionInfo *exception)
2325{
2326#define EqualizeImageTag "Equalize/Image"
2327
2328 cl_command_queue
2329 queue;
2330
2331 cl_context
2332 context;
2333
2334 cl_int
2335 clStatus;
2336
2337 cl_mem
2338 equalizeMapBuffer,
2339 histogramBuffer,
2340 imageBuffer;
2341
2342 cl_kernel
2343 equalizeKernel,
2344 histogramKernel;
2345
2346 cl_event
2347 event;
2348
2349 cl_uint
2350 event_count;
2351
2352 cl_uint4
2353 *histogram;
2354
2355 cl_event
2356 *events;
2357
2358 cl_float4
2359 white,
2360 black,
2361 intensity,
2362 *map;
2363
2364 MagickBooleanType
2365 outputReady,
2366 status;
2367
2368 MagickCLEnv
2369 clEnv;
2370
2371 MagickSizeType
2372 length;
2373
2374 PixelPacket
2375 *equalize_map;
2376
2377 ssize_t
2378 i;
2379
2380 size_t
2381 global_work_size[2];
2382
2383 map=NULL;
2384 histogram=NULL;
2385 equalize_map=NULL;
2386 imageBuffer = NULL;
2387 histogramBuffer = NULL;
2388 equalizeMapBuffer = NULL;
2389 histogramKernel = NULL;
2390 equalizeKernel = NULL;
2391 context = NULL;
2392 queue = NULL;
2393 outputReady = MagickFalse;
2394
2395 assert(image != (Image *) NULL);
2396 assert(image->signature == MagickCoreSignature);
2397 if (IsEventLogging() != MagickFalse)
2398 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
2399
2400 /*
2401 * initialize opencl env
2402 */
2403 clEnv = GetDefaultOpenCLEnv();
2404 context = GetOpenCLContext(clEnv);
2405 queue = AcquireOpenCLCommandQueue(clEnv);
2406
2407 /*
2408 Allocate and initialize histogram arrays.
2409 */
2410 length=MaxMap+1UL;
2411 histogram=(cl_uint4 *) AcquireQuantumMemory(length, sizeof(*histogram));
2412 if (histogram == (cl_uint4 *) NULL)
2413 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
2414
2415 /* reset histogram */
2416 (void) memset(histogram,0,length*sizeof(*histogram));
2417
2418 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
2419 if (imageBuffer == (cl_mem) NULL)
2420 {
2421 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2422 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
2423 goto cleanup;
2424 }
2425
2426 /* create a CL buffer for histogram */
2427 histogramBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, length * sizeof(cl_uint4), histogram, &clStatus);
2428 if (clStatus != CL_SUCCESS)
2429 {
2430 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2431 goto cleanup;
2432 }
2433
2434 status = LaunchHistogramKernel(clEnv, queue, imageBuffer, histogramBuffer, image, channel, exception);
2435 if (status == MagickFalse)
2436 goto cleanup;
2437
2438 /* this blocks, should be fixed it in the future */
2439 events=GetOpenCLEvents(image,&event_count);
2440 clEnv->library->clEnqueueMapBuffer(queue, histogramBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(cl_uint4), event_count, events, NULL, &clStatus);
2441 events=(cl_event *) RelinquishMagickMemory(events);
2442 if (clStatus != CL_SUCCESS)
2443 {
2444 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
2445 goto cleanup;
2446 }
2447
2448 /* unmap, don't block gpu to use this buffer again. */
2449 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
2450 if (clStatus != CL_SUCCESS)
2451 {
2452 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
2453 goto cleanup;
2454 }
2455
2456 /* CPU stuff */
2457 equalize_map=(PixelPacket *) AcquireQuantumMemory(length, sizeof(*equalize_map));
2458 if (equalize_map == (PixelPacket *) NULL)
2459 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
2460
2461 map=(cl_float4 *) AcquireQuantumMemory(length,sizeof(*map));
2462 if (map == (cl_float4 *) NULL)
2463 ThrowBinaryException(ResourceLimitWarning,"MemoryAllocationFailed", image->filename);
2464
2465 /*
2466 Integrate the histogram to get the equalization map.
2467 */
2468 (void) memset(&intensity,0,sizeof(intensity));
2469 for (i=0; i <= (ssize_t) MaxMap; i++)
2470 {
2471 if ((channel & SyncChannels) != 0)
2472 {
2473 intensity.z+=histogram[i].s[2];
2474 map[i]=intensity;
2475 continue;
2476 }
2477 if ((channel & RedChannel) != 0)
2478 intensity.z+=histogram[i].s[2];
2479 if ((channel & GreenChannel) != 0)
2480 intensity.y+=histogram[i].s[1];
2481 if ((channel & BlueChannel) != 0)
2482 intensity.x+=histogram[i].s[0];
2483 if ((channel & OpacityChannel) != 0)
2484 intensity.w+=histogram[i].s[3];
2485 /*
2486 if (((channel & IndexChannel) != 0) &&
2487 (image->colorspace == CMYKColorspace))
2488 {
2489 intensity.index+=histogram[i].index;
2490 }
2491 */
2492 map[i]=intensity;
2493 }
2494 black=map[0];
2495 white=map[(int) MaxMap];
2496 (void) memset(equalize_map,0,length*sizeof(*equalize_map));
2497 for (i=0; i <= (ssize_t) MaxMap; i++)
2498 {
2499 if ((channel & SyncChannels) != 0)
2500 {
2501 if (white.z != black.z)
2502 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2503 (map[i].z-black.z))/(white.z-black.z)));
2504 continue;
2505 }
2506 if (((channel & RedChannel) != 0) && (white.z != black.z))
2507 equalize_map[i].red=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2508 (map[i].z-black.z))/(white.z-black.z)));
2509 if (((channel & GreenChannel) != 0) && (white.y != black.y))
2510 equalize_map[i].green=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2511 (map[i].y-black.y))/(white.y-black.y)));
2512 if (((channel & BlueChannel) != 0) && (white.x != black.x))
2513 equalize_map[i].blue=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2514 (map[i].x-black.x))/(white.x-black.x)));
2515 if (((channel & OpacityChannel) != 0) && (white.w != black.w))
2516 equalize_map[i].opacity=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2517 (map[i].w-black.w))/(white.w-black.w)));
2518 /*
2519 if ((((channel & IndexChannel) != 0) &&
2520 (image->colorspace == CMYKColorspace)) &&
2521 (white.index != black.index))
2522 equalize_map[i].index=ScaleMapToQuantum((MagickRealType) ((MaxMap*
2523 (map[i].index-black.index))/(white.index-black.index)));
2524 */
2525 }
2526
2527 if (image->storage_class == PseudoClass)
2528 {
2529 /*
2530 Equalize colormap.
2531 */
2532 for (i=0; i < (ssize_t) image->colors; i++)
2533 {
2534 if ((channel & SyncChannels) != 0)
2535 {
2536 if (white.z != black.z)
2537 {
2538 image->colormap[i].red=equalize_map[
2539 ScaleQuantumToMap(image->colormap[i].red)].red;
2540 image->colormap[i].green=equalize_map[
2541 ScaleQuantumToMap(image->colormap[i].green)].red;
2542 image->colormap[i].blue=equalize_map[
2543 ScaleQuantumToMap(image->colormap[i].blue)].red;
2544 image->colormap[i].opacity=equalize_map[
2545 ScaleQuantumToMap(image->colormap[i].opacity)].red;
2546 }
2547 continue;
2548 }
2549 if (((channel & RedChannel) != 0) && (white.z != black.z))
2550 image->colormap[i].red=equalize_map[
2551 ScaleQuantumToMap(image->colormap[i].red)].red;
2552 if (((channel & GreenChannel) != 0) && (white.y != black.y))
2553 image->colormap[i].green=equalize_map[
2554 ScaleQuantumToMap(image->colormap[i].green)].green;
2555 if (((channel & BlueChannel) != 0) && (white.x != black.x))
2556 image->colormap[i].blue=equalize_map[
2557 ScaleQuantumToMap(image->colormap[i].blue)].blue;
2558 if (((channel & OpacityChannel) != 0) &&
2559 (white.w != black.w))
2560 image->colormap[i].opacity=equalize_map[
2561 ScaleQuantumToMap(image->colormap[i].opacity)].opacity;
2562 }
2563 }
2564
2565 /* create a CL buffer for eqaulize_map */
2566 equalizeMapBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, length * sizeof(PixelPacket), equalize_map, &clStatus);
2567 if (clStatus != CL_SUCCESS)
2568 {
2569 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
2570 goto cleanup;
2571 }
2572
2573 /* get the OpenCL kernel */
2574 equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize");
2575 if (equalizeKernel == NULL)
2576 {
2577 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2578 goto cleanup;
2579 }
2580
2581 /* set the kernel arguments */
2582 i = 0;
2583 clStatus=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2584 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(ChannelType),&channel);
2585 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_mem),(void *)&equalizeMapBuffer);
2586 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_float4),&white);
2587 clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(cl_float4),&black);
2588 if (clStatus != CL_SUCCESS)
2589 {
2590 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2591 goto cleanup;
2592 }
2593
2594 /* launch the kernel */
2595 global_work_size[0] = image->columns;
2596 global_work_size[1] = image->rows;
2597
2598 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, equalizeKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
2599
2600 if (clStatus != CL_SUCCESS)
2601 {
2602 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2603 goto cleanup;
2604 }
2605 if (RecordProfileData(clEnv,EqualizeKernel,event) == MagickFalse)
2606 AddOpenCLEvent(image,event);
2607 clEnv->library->clReleaseEvent(event);
2608
2609cleanup:
2610 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2611
2612 if (imageBuffer != (cl_mem) NULL)
2613 clEnv->library->clReleaseMemObject(imageBuffer);
2614
2615 if (map!=NULL)
2616 map=(cl_float4 *) RelinquishMagickMemory(map);
2617
2618 if (equalizeMapBuffer!=NULL)
2619 clEnv->library->clReleaseMemObject(equalizeMapBuffer);
2620 if (equalize_map!=NULL)
2621 equalize_map=(PixelPacket *) RelinquishMagickMemory(equalize_map);
2622
2623 if (histogramBuffer!=NULL)
2624 clEnv->library->clReleaseMemObject(histogramBuffer);
2625 if (histogram!=NULL)
2626 histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
2627
2628 if (histogramKernel!=NULL)
2629 RelinquishOpenCLKernel(clEnv, histogramKernel);
2630 if (equalizeKernel!=NULL)
2631 RelinquishOpenCLKernel(clEnv, equalizeKernel);
2632
2633 if (queue != NULL)
2634 RelinquishOpenCLCommandQueue(clEnv, queue);
2635
2636 return(outputReady);
2637}
2638
2639MagickPrivate MagickBooleanType AccelerateEqualizeImage(Image *image,
2640 const ChannelType channel,ExceptionInfo *exception)
2641{
2642 MagickBooleanType
2643 status;
2644
2645 assert(image != NULL);
2646 assert(exception != (ExceptionInfo *) NULL);
2647
2648 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2649 (checkAccelerateCondition(image, channel) == MagickFalse) ||
2650 (checkHistogramCondition(image, channel) == MagickFalse))
2651 return(MagickFalse);
2652
2653 status=ComputeEqualizeImage(image,channel,exception);
2654 return(status);
2655}
2656
2657/*
2658%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2659% %
2660% %
2661% %
2662% A c c e l e r a t e F u n c t i o n I m a g e %
2663% %
2664% %
2665% %
2666%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2667*/
2668
2669static MagickBooleanType ComputeFunctionImage(Image *image,
2670 const ChannelType channel,const MagickFunction function,
2671 const size_t number_parameters,const double *parameters,
2672 ExceptionInfo *exception)
2673{
2674 cl_command_queue
2675 queue;
2676
2677 cl_context
2678 context;
2679
2680 cl_int
2681 clStatus;
2682
2683 cl_kernel
2684 clkernel;
2685
2686 cl_event
2687 event;
2688
2689 cl_mem
2690 imageBuffer,
2691 parametersBuffer;
2692
2693 cl_event
2694 *events;
2695
2696 float
2697 *parametersBufferPtr;
2698
2699 MagickBooleanType
2700 status;
2701
2702 MagickCLEnv
2703 clEnv;
2704
2705 size_t
2706 globalWorkSize[2];
2707
2708 unsigned int
2709 event_count,
2710 i;
2711
2712 status = MagickFalse;
2713
2714 context = NULL;
2715 clkernel = NULL;
2716 queue = NULL;
2717 imageBuffer = NULL;
2718 parametersBuffer = NULL;
2719
2720 clEnv = GetDefaultOpenCLEnv();
2721 context = GetOpenCLContext(clEnv);
2722
2723 queue = AcquireOpenCLCommandQueue(clEnv);
2724
2725 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
2726 if (imageBuffer == (cl_mem) NULL)
2727 {
2728 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2729 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
2730 goto cleanup;
2731 }
2732
2733
2734 {
2735 parametersBufferPtr = (float*)AcquireMagickMemory(number_parameters * sizeof(float));
2736
2737 for (i = 0; i < number_parameters; i++)
2738 parametersBufferPtr[i] = (float)parameters[i];
2739
2740 parametersBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, number_parameters * sizeof(float), parametersBufferPtr, &clStatus);
2741 parametersBufferPtr=(float *) RelinquishMagickMemory(parametersBufferPtr);
2742 }
2743
2744 clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ComputeFunction");
2745 if (clkernel == NULL)
2746 {
2747 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2748 goto cleanup;
2749 }
2750
2751 /* set the kernel arguments */
2752 i = 0;
2753 clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2754 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel);
2755 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(MagickFunction),(void *)&function);
2756 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&number_parameters);
2757 clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&parametersBuffer);
2758 if (clStatus != CL_SUCCESS)
2759 {
2760 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2761 goto cleanup;
2762 }
2763
2764 globalWorkSize[0] = image->columns;
2765 globalWorkSize[1] = image->rows;
2766 /* launch the kernel */
2767 events=GetOpenCLEvents(image,&event_count);
2768 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, event_count, events, &event);
2769 events=(cl_event *) RelinquishMagickMemory(events);
2770 if (clStatus != CL_SUCCESS)
2771 {
2772 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2773 goto cleanup;
2774 }
2775 if (RecordProfileData(clEnv,ComputeFunctionKernel,event) == MagickFalse)
2776 AddOpenCLEvent(image,event);
2777 clEnv->library->clReleaseEvent(event);
2778 status = MagickTrue;
2779
2780cleanup:
2781 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2782
2783 if (imageBuffer != (cl_mem) NULL)
2784 clEnv->library->clReleaseMemObject(imageBuffer);
2785 if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel);
2786 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
2787 if (parametersBuffer != NULL) clEnv->library->clReleaseMemObject(parametersBuffer);
2788
2789 return(status);
2790}
2791
2792MagickPrivate MagickBooleanType AccelerateFunctionImage(Image *image,
2793 const ChannelType channel,const MagickFunction function,
2794 const size_t number_parameters,const double *parameters,
2795 ExceptionInfo *exception)
2796{
2797 MagickBooleanType
2798 status;
2799
2800 assert(image != NULL);
2801 assert(exception != (ExceptionInfo *) NULL);
2802
2803 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2804 (checkAccelerateCondition(image, channel) == MagickFalse))
2805 return(MagickFalse);
2806
2807 status=ComputeFunctionImage(image, channel, function, number_parameters, parameters, exception);
2808 return(status);
2809}
2810
2811/*
2812%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2813% %
2814% %
2815% %
2816% A c c e l e r a t e G r a y s c a l e I m a g e %
2817% %
2818% %
2819% %
2820%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2821*/
2822
2823MagickBooleanType ComputeGrayscaleImage(Image *image,
2824 const PixelIntensityMethod method,ExceptionInfo *exception)
2825{
2826 cl_command_queue
2827 queue;
2828
2829 cl_context
2830 context;
2831
2832 cl_int
2833 clStatus,
2834 intensityMethod;
2835
2836 cl_int
2837 colorspace;
2838
2839 cl_kernel
2840 grayscaleKernel;
2841
2842 cl_event
2843 event;
2844
2845 cl_mem
2846 imageBuffer;
2847
2848 cl_uint
2849 event_count;
2850
2851 cl_event
2852 *events;
2853
2854 MagickBooleanType
2855 outputReady;
2856
2857 MagickCLEnv
2858 clEnv;
2859
2860 ssize_t
2861 i;
2862
2863 imageBuffer = NULL;
2864 grayscaleKernel = NULL;
2865
2866 assert(image != (Image *) NULL);
2867 assert(image->signature == MagickCoreSignature);
2868 if (IsEventLogging() != MagickFalse)
2869 (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
2870
2871 /*
2872 * initialize opencl env
2873 */
2874 clEnv = GetDefaultOpenCLEnv();
2875 context = GetOpenCLContext(clEnv);
2876 queue = AcquireOpenCLCommandQueue(clEnv);
2877
2878 outputReady = MagickFalse;
2879
2880 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
2881 if (imageBuffer == (cl_mem) NULL)
2882 {
2883 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
2884 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
2885 goto cleanup;
2886 }
2887
2888 intensityMethod = method;
2889 colorspace = image->colorspace;
2890
2891 grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Grayscale");
2892 if (grayscaleKernel == NULL)
2893 {
2894 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
2895 goto cleanup;
2896 }
2897
2898 i = 0;
2899 clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
2900 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&intensityMethod);
2901 clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&colorspace);
2902 if (clStatus != CL_SUCCESS)
2903 {
2904 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
2905 printf("no kernel\n");
2906 goto cleanup;
2907 }
2908
2909 {
2910 size_t global_work_size[2];
2911 global_work_size[0] = image->columns;
2912 global_work_size[1] = image->rows;
2913 /* launch the kernel */
2914 events=GetOpenCLEvents(image,&event_count);
2915 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, grayscaleKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
2916 events=(cl_event *) RelinquishMagickMemory(events);
2917 if (clStatus != CL_SUCCESS)
2918 {
2919 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
2920 goto cleanup;
2921 }
2922 if (RecordProfileData(clEnv,GrayScaleKernel,event) == MagickFalse)
2923 AddOpenCLEvent(image,event);
2924 clEnv->library->clReleaseEvent(event);
2925 }
2926
2927 outputReady=MagickTrue;
2928
2929cleanup:
2930 OpenCLLogException(__FUNCTION__,__LINE__,exception);
2931
2932 if (imageBuffer != (cl_mem) NULL)
2933 clEnv->library->clReleaseMemObject(imageBuffer);
2934 if (grayscaleKernel!=NULL)
2935 RelinquishOpenCLKernel(clEnv, grayscaleKernel);
2936 if (queue != NULL)
2937 RelinquishOpenCLCommandQueue(clEnv, queue);
2938
2939 return(outputReady);
2940}
2941
2942MagickPrivate MagickBooleanType AccelerateGrayscaleImage(Image* image,
2943 const PixelIntensityMethod method,ExceptionInfo *exception)
2944{
2945 MagickBooleanType
2946 status;
2947
2948 assert(image != NULL);
2949 assert(exception != (ExceptionInfo *) NULL);
2950
2951 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
2952 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
2953 return(MagickFalse);
2954
2955 if (method == Rec601LuminancePixelIntensityMethod || method == Rec709LuminancePixelIntensityMethod)
2956 return(MagickFalse);
2957
2958 if (image->colorspace != sRGBColorspace)
2959 return(MagickFalse);
2960
2961 status=ComputeGrayscaleImage(image,method,exception);
2962 return(status);
2963}
2964
2965/*
2966%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2967% %
2968% %
2969% %
2970% A c c e l e r a t e L o c a l C o n t r a s t I m a g e %
2971% %
2972% %
2973% %
2974%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2975*/
2976
2977static Image *ComputeLocalContrastImage(const Image *image,
2978 const double radius,const double strength,ExceptionInfo *exception)
2979{
2980 cl_command_queue
2981 queue;
2982
2983 cl_context
2984 context;
2985
2986 cl_int
2987 clStatus,
2988 iRadius;
2989
2990 cl_kernel
2991 blurRowKernel,
2992 blurColumnKernel;
2993
2994 cl_event
2995 event;
2996
2997 cl_mem
2998 filteredImageBuffer,
2999 imageBuffer,
3000 tempImageBuffer;
3001
3002 cl_event
3003 *events;
3004
3005 Image
3006 *filteredImage;
3007
3008 MagickBooleanType
3009 outputReady;
3010
3011 MagickCLEnv
3012 clEnv;
3013
3014 MagickSizeType
3015 length;
3016
3017 unsigned int
3018 event_count,
3019 i,
3020 imageColumns,
3021 imageRows,
3022 passes;
3023
3024 clEnv = NULL;
3025 filteredImage = NULL;
3026 context = NULL;
3027 imageBuffer = NULL;
3028 filteredImageBuffer = NULL;
3029 tempImageBuffer = NULL;
3030 blurRowKernel = NULL;
3031 blurColumnKernel = NULL;
3032 queue = NULL;
3033 outputReady = MagickFalse;
3034
3035 clEnv = GetDefaultOpenCLEnv();
3036 context = GetOpenCLContext(clEnv);
3037 queue = AcquireOpenCLCommandQueue(clEnv);
3038
3039 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3040 if (filteredImage == (Image *) NULL)
3041 goto cleanup;
3042
3043 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
3044 if (imageBuffer == (cl_mem) NULL)
3045 {
3046 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3047 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3048 goto cleanup;
3049 }
3050 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
3051 if (filteredImageBuffer == (cl_mem) NULL)
3052 {
3053 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3054 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3055 goto cleanup;
3056 }
3057
3058 {
3059 /* create temp buffer */
3060 {
3061 length = image->columns * image->rows;
3062 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(float), NULL, &clStatus);
3063 if (clStatus != CL_SUCCESS)
3064 {
3065 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3066 goto cleanup;
3067 }
3068 }
3069
3070 /* get the opencl kernel */
3071 {
3072 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "LocalContrastBlurRow");
3073 if (blurRowKernel == NULL)
3074 {
3075 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3076 goto cleanup;
3077 };
3078
3079 blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "LocalContrastBlurApplyColumn");
3080 if (blurColumnKernel == NULL)
3081 {
3082 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3083 goto cleanup;
3084 };
3085 }
3086
3087 {
3088 imageColumns = (unsigned int) image->columns;
3089 imageRows = (unsigned int) image->rows;
3090 iRadius = (cl_int) (image->rows > image->columns ? image->rows : image->columns) * 0.002f * fabs(radius); /*Normalized radius, 100% gives blur radius of 20% of the largest dimension */
3091
3092 passes = (((1.0f * imageRows) * imageColumns * iRadius) + 3999999999) / 4000000000.0f;
3093 passes = (passes < 1) ? 1: passes;
3094
3095 /* set the kernel arguments */
3096 i = 0;
3097 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
3098 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
3099 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
3100 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_int),(void *)&iRadius);
3101 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
3102 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
3103
3104 if (clStatus != CL_SUCCESS)
3105 {
3106 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3107 goto cleanup;
3108 }
3109 }
3110
3111 /* launch the kernel */
3112 {
3113 int x;
3114 for (x = 0; x < passes; ++x) {
3115 size_t gsize[2];
3116 size_t wsize[2];
3117 size_t goffset[2];
3118
3119 gsize[0] = 256;
3120 gsize[1] = (image->rows + passes - 1) / passes;
3121 wsize[0] = 256;
3122 wsize[1] = 1;
3123 goffset[0] = 0;
3124 goffset[1] = x * gsize[1];
3125
3126 events=GetOpenCLEvents(image,&event_count);
3127 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, event_count, events, &event);
3128 events=(cl_event *) RelinquishMagickMemory(events);
3129 if (clStatus != CL_SUCCESS)
3130 {
3131 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3132 goto cleanup;
3133 }
3134 clEnv->library->clFlush(queue);
3135 if (RecordProfileData(clEnv,LocalContrastBlurRowKernel,event) == MagickFalse)
3136 {
3137 AddOpenCLEvent(image,event);
3138 AddOpenCLEvent(filteredImage, event);
3139 }
3140 clEnv->library->clReleaseEvent(event);
3141 }
3142 }
3143
3144 {
3145 cl_float FStrength = strength;
3146 i = 0;
3147 clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
3148 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
3149 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
3150 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&iRadius);
3151 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float),(void *)&FStrength);
3152 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
3153 clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
3154
3155 if (clStatus != CL_SUCCESS)
3156 {
3157 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3158 goto cleanup;
3159 }
3160 }
3161
3162 /* launch the kernel */
3163 {
3164 int x;
3165 for (x = 0; x < passes; ++x) {
3166 size_t gsize[2];
3167 size_t wsize[2];
3168 size_t goffset[2];
3169
3170 gsize[0] = ((image->columns + 3) / 4) * 4;
3171 gsize[1] = ((((image->rows + 63) / 64) + (passes + 1)) / passes) * 64;
3172 wsize[0] = 4;
3173 wsize[1] = 64;
3174 goffset[0] = 0;
3175 goffset[1] = x * gsize[1];
3176
3177 events=GetOpenCLEvents(image,&event_count);
3178 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, event_count, events, &event);
3179 events=(cl_event *) RelinquishMagickMemory(events);
3180 if (clStatus != CL_SUCCESS)
3181 {
3182 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3183 goto cleanup;
3184 }
3185 clEnv->library->clFlush(queue);
3186 if (RecordProfileData(clEnv, LocalContrastBlurApplyColumnKernel, event) == MagickFalse)
3187 {
3188 AddOpenCLEvent(image,event);
3189 AddOpenCLEvent(filteredImage,event);
3190 }
3191 clEnv->library->clReleaseEvent(event);
3192 }
3193 }
3194 }
3195
3196 outputReady = MagickTrue;
3197
3198
3199cleanup:
3200 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3201
3202 if (imageBuffer != (cl_mem) NULL)
3203 clEnv->library->clReleaseMemObject(imageBuffer);
3204 if (filteredImageBuffer != (cl_mem) NULL)
3205 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3206 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
3207 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
3208 if (blurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurColumnKernel);
3209 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3210 if ((outputReady == MagickFalse) && (filteredImage != NULL))
3211 filteredImage=(Image *) DestroyImage(filteredImage);
3212 return(filteredImage);
3213}
3214
3215MagickPrivate Image *AccelerateLocalContrastImage(const Image *image,
3216 const double radius,const double strength,ExceptionInfo *exception)
3217{
3218 Image
3219 *filteredImage;
3220
3221 assert(image != NULL);
3222 assert(exception != (ExceptionInfo *) NULL);
3223
3224 if ((checkOpenCLEnvironment(exception) == MagickFalse))
3225 return NULL;
3226
3227 filteredImage=ComputeLocalContrastImage(image,radius,strength,exception);
3228
3229 return(filteredImage);
3230}
3231
3232/*
3233%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3234% %
3235% %
3236% %
3237% A c c e l e r a t e M o d u l a t e I m a g e %
3238% %
3239% %
3240% %
3241%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3242*/
3243
3244MagickBooleanType ComputeModulateImage(Image *image,
3245 double percent_brightness, double percent_hue, double percent_saturation,
3246 ColorspaceType colorspace, ExceptionInfo *exception)
3247{
3248 cl_float
3249 bright,
3250 hue,
3251 saturation;
3252
3253 cl_context
3254 context;
3255
3256 cl_command_queue
3257 queue;
3258
3259 cl_int
3260 color,
3261 clStatus;
3262
3263 cl_kernel
3264 modulateKernel;
3265
3266 cl_event
3267 event;
3268
3269 cl_mem
3270 imageBuffer;
3271
3272 cl_event
3273 *events;
3274
3275 MagickBooleanType
3276 outputReady;
3277
3278 MagickCLEnv
3279 clEnv;
3280
3281 ssize_t
3282 i;
3283
3284 unsigned int
3285 event_count;
3286
3287 imageBuffer = NULL;
3288 modulateKernel = NULL;
3289 event_count = 0;
3290
3291 assert(image != (Image *)NULL);
3292 assert(image->signature == MagickCoreSignature);
3293 if (IsEventLogging() != MagickFalse)
3294 (void) LogMagickEvent(TraceEvent, GetMagickModule(), "%s", image->filename);
3295
3296 /*
3297 * initialize opencl env
3298 */
3299 clEnv = GetDefaultOpenCLEnv();
3300 context = GetOpenCLContext(clEnv);
3301 queue = AcquireOpenCLCommandQueue(clEnv);
3302
3303 outputReady = MagickFalse;
3304
3305 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
3306 if (imageBuffer == (cl_mem) NULL)
3307 {
3308 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3309 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3310 goto cleanup;
3311 }
3312
3313 modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate");
3314 if (modulateKernel == NULL)
3315 {
3316 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3317 goto cleanup;
3318 }
3319
3320 bright = percent_brightness;
3321 hue = percent_hue;
3322 saturation = percent_saturation;
3323 color = colorspace;
3324
3325 i = 0;
3326 clStatus = clEnv->library->clSetKernelArg(modulateKernel, i++, sizeof(cl_mem), (void *)&imageBuffer);
3327 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++, sizeof(cl_float), &bright);
3328 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++, sizeof(cl_float), &hue);
3329 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++, sizeof(cl_float), &saturation);
3330 clStatus |= clEnv->library->clSetKernelArg(modulateKernel, i++, sizeof(cl_float), &color);
3331 if (clStatus != CL_SUCCESS)
3332 {
3333 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3334 printf("no kernel\n");
3335 goto cleanup;
3336 }
3337
3338 {
3339 size_t global_work_size[2];
3340 global_work_size[0] = image->columns;
3341 global_work_size[1] = image->rows;
3342 /* launch the kernel */
3343 events=GetOpenCLEvents(image,&event_count);
3344 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, event_count, events, &event);
3345 events=(cl_event *) RelinquishMagickMemory(events);
3346 if (clStatus != CL_SUCCESS)
3347 {
3348 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3349 goto cleanup;
3350 }
3351 if (RecordProfileData(clEnv, ModulateKernel, event) == MagickFalse)
3352 AddOpenCLEvent(image,event);
3353 clEnv->library->clReleaseEvent(event);
3354 }
3355
3356 outputReady=MagickTrue;
3357
3358cleanup:
3359 OpenCLLogException(__FUNCTION__, __LINE__, exception);
3360
3361 if (imageBuffer != (cl_mem) NULL)
3362 clEnv->library->clReleaseMemObject(imageBuffer);
3363 if (modulateKernel != NULL)
3364 RelinquishOpenCLKernel(clEnv, modulateKernel);
3365 if (queue != NULL)
3366 RelinquishOpenCLCommandQueue(clEnv, queue);
3367
3368 return(outputReady);
3369}
3370
3371MagickPrivate MagickBooleanType AccelerateModulateImage(Image *image,
3372 double percent_brightness, double percent_hue, double percent_saturation,
3373 ColorspaceType colorspace, ExceptionInfo *exception)
3374{
3375 MagickBooleanType
3376 status;
3377
3378 assert(image != NULL);
3379 assert(exception != (ExceptionInfo *)NULL);
3380
3381 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3382 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
3383 return(MagickFalse);
3384
3385 if ((colorspace != HSLColorspace && colorspace != UndefinedColorspace))
3386 return(MagickFalse);
3387
3388 status = ComputeModulateImage(image, percent_brightness, percent_hue, percent_saturation, colorspace, exception);
3389 return(status);
3390}
3391
3392/*
3393%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3394% %
3395% %
3396% %
3397% A c c e l e r a t e M o t i o n B l u r I m a g e %
3398% %
3399% %
3400% %
3401%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3402*/
3403
3404static Image* ComputeMotionBlurImage(const Image *image,
3405 const ChannelType channel,const double *kernel,const size_t width,
3406 const OffsetInfo *offset,ExceptionInfo *exception)
3407{
3408 cl_command_queue
3409 queue;
3410
3411 cl_context
3412 context;
3413
3414 cl_float4
3415 biasPixel;
3416
3417 cl_int
3418 clStatus;
3419
3420 cl_kernel
3421 motionBlurKernel;
3422
3423 cl_event
3424 event;
3425
3426 cl_mem
3427 filteredImageBuffer,
3428 imageBuffer,
3429 imageKernelBuffer,
3430 offsetBuffer;
3431
3432 cl_uint
3433 event_count;
3434
3435 cl_event
3436 *events;
3437
3438 float
3439 *kernelBufferPtr;
3440
3441 Image
3442 *filteredImage;
3443
3444 int
3445 *offsetBufferPtr;
3446
3447 MagickBooleanType
3448 outputReady;
3449
3450 MagickCLEnv
3451 clEnv;
3452
3453 MagickPixelPacket
3454 bias;
3455
3456 size_t
3457 global_work_size[2],
3458 local_work_size[2];
3459
3460 unsigned int
3461 i,
3462 imageHeight,
3463 imageWidth,
3464 matte;
3465
3466 outputReady = MagickFalse;
3467 context = NULL;
3468 filteredImage = NULL;
3469 imageBuffer = NULL;
3470 filteredImageBuffer = NULL;
3471 imageKernelBuffer = NULL;
3472 motionBlurKernel = NULL;
3473 queue = NULL;
3474
3475 clEnv = GetDefaultOpenCLEnv();
3476 context = GetOpenCLContext(clEnv);
3477
3478 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3479 if (filteredImage == (Image *) NULL)
3480 goto cleanup;
3481
3482 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
3483 if (imageBuffer == (cl_mem) NULL)
3484 {
3485 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3486 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3487 goto cleanup;
3488 }
3489 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage, exception);
3490 if (filteredImageBuffer == (cl_mem) NULL)
3491 {
3492 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3493 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3494 goto cleanup;
3495 }
3496
3497 imageKernelBuffer = clEnv->library->clCreateBuffer(context,
3498 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(float), NULL,
3499 &clStatus);
3500 if (clStatus != CL_SUCCESS)
3501 {
3502 (void) ThrowMagickException(exception, GetMagickModule(),
3503 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
3504 goto cleanup;
3505 }
3506
3507 queue = AcquireOpenCLCommandQueue(clEnv);
3508 events=GetOpenCLEvents(image,&event_count);
3509 /* this blocks, should be fixed it in the future */
3510 kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer,
3511 CL_TRUE, CL_MAP_WRITE, 0, width * sizeof(float), event_count, events, NULL, &clStatus);
3512 events=(cl_event *) RelinquishMagickMemory(events);
3513 if (clStatus != CL_SUCCESS)
3514 {
3515 (void) ThrowMagickException(exception, GetMagickModule(),
3516 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
3517 goto cleanup;
3518 }
3519 for (i = 0; i < width; i++)
3520 {
3521 kernelBufferPtr[i] = (float) kernel[i];
3522 }
3523 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr,
3524 0, NULL, NULL);
3525 if (clStatus != CL_SUCCESS)
3526 {
3527 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3528 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
3529 goto cleanup;
3530 }
3531
3532 offsetBuffer = clEnv->library->clCreateBuffer(context,
3533 CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, width * sizeof(cl_int2), NULL,
3534 &clStatus);
3535 if (clStatus != CL_SUCCESS)
3536 {
3537 (void) ThrowMagickException(exception, GetMagickModule(),
3538 ResourceLimitError, "clEnv->library->clCreateBuffer failed.",".");
3539 goto cleanup;
3540 }
3541
3542 offsetBufferPtr = (int*)clEnv->library->clEnqueueMapBuffer(queue, offsetBuffer, CL_TRUE,
3543 CL_MAP_WRITE, 0, width * sizeof(cl_int2), 0, NULL, NULL, &clStatus);
3544 if (clStatus != CL_SUCCESS)
3545 {
3546 (void) ThrowMagickException(exception, GetMagickModule(),
3547 ResourceLimitError, "clEnv->library->clEnqueueMapBuffer failed.",".");
3548 goto cleanup;
3549 }
3550 for (i = 0; i < width; i++)
3551 {
3552 offsetBufferPtr[2*i] = (int)offset[i].x;
3553 offsetBufferPtr[2*i+1] = (int)offset[i].y;
3554 }
3555 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, offsetBuffer, offsetBufferPtr, 0,
3556 NULL, NULL);
3557 if (clStatus != CL_SUCCESS)
3558 {
3559 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3560 "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
3561 goto cleanup;
3562 }
3563
3564
3565 /*
3566 Get the OpenCL kernel.
3567 */
3568 motionBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE,
3569 "MotionBlur");
3570 if (motionBlurKernel == NULL)
3571 {
3572 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3573 "AcquireOpenCLKernel failed.", "'%s'", ".");
3574 goto cleanup;
3575 }
3576
3577 /*
3578 Set the kernel arguments.
3579 */
3580 i = 0;
3581 clStatus=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
3582 (void *)&imageBuffer);
3583 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
3584 (void *)&filteredImageBuffer);
3585 imageWidth = (unsigned int) image->columns;
3586 imageHeight = (unsigned int) image->rows;
3587 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
3588 &imageWidth);
3589 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
3590 &imageHeight);
3591 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
3592 (void *)&imageKernelBuffer);
3593 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int),
3594 &width);
3595 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_mem),
3596 (void *)&offsetBuffer);
3597
3598 GetMagickPixelPacket(image,&bias);
3599 biasPixel.s[0] = bias.red;
3600 biasPixel.s[1] = bias.green;
3601 biasPixel.s[2] = bias.blue;
3602 biasPixel.s[3] = bias.opacity;
3603 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(cl_float4), &biasPixel);
3604
3605 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(ChannelType), &channel);
3606 matte = (image->matte != MagickFalse)?1:0;
3607 clStatus|=clEnv->library->clSetKernelArg(motionBlurKernel,i++,sizeof(unsigned int), &matte);
3608 if (clStatus != CL_SUCCESS)
3609 {
3610 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3611 "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3612 goto cleanup;
3613 }
3614
3615 /*
3616 Launch the kernel.
3617 */
3618 local_work_size[0] = 16;
3619 local_work_size[1] = 16;
3620 global_work_size[0] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3621 (unsigned int) image->columns,(unsigned int) local_work_size[0]);
3622 global_work_size[1] = (size_t)padGlobalWorkgroupSizeToLocalWorkgroupSize(
3623 (unsigned int) image->rows,(unsigned int) local_work_size[1]);
3624 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, motionBlurKernel, 2, NULL,
3625 global_work_size, local_work_size, 0, NULL, &event);
3626
3627 if (clStatus != CL_SUCCESS)
3628 {
3629 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
3630 "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3631 goto cleanup;
3632 }
3633 if (RecordProfileData(clEnv,MotionBlurKernel,event) == MagickFalse)
3634 {
3635 AddOpenCLEvent(image, event);
3636 AddOpenCLEvent(filteredImage, event);
3637 }
3638 clEnv->library->clReleaseEvent(event);
3639
3640 outputReady = MagickTrue;
3641
3642cleanup:
3643
3644 if (imageBuffer != (cl_mem) NULL)
3645 clEnv->library->clReleaseMemObject(imageBuffer);
3646 if (filteredImageBuffer != (cl_mem) NULL)
3647 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3648 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
3649 if (motionBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, motionBlurKernel);
3650 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3651 if ((outputReady == MagickFalse) && (filteredImage != NULL))
3652 filteredImage=(Image *) DestroyImage(filteredImage);
3653
3654 return(filteredImage);
3655}
3656
3657MagickPrivate Image *AccelerateMotionBlurImage(const Image *image,
3658 const ChannelType channel,const double* kernel,const size_t width,
3659 const OffsetInfo *offset,ExceptionInfo *exception)
3660{
3661 Image
3662 *filteredImage;
3663
3664 assert(image != NULL);
3665 assert(kernel != (double *) NULL);
3666 assert(offset != (OffsetInfo *) NULL);
3667 assert(exception != (ExceptionInfo *) NULL);
3668
3669 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3670 (checkAccelerateCondition(image, channel) == MagickFalse))
3671 return NULL;
3672
3673 filteredImage=ComputeMotionBlurImage(image, channel, kernel, width,
3674 offset, exception);
3675 return(filteredImage);
3676}
3677
3678/*
3679%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3680% %
3681% %
3682% %
3683% A c c e l e r a t e R a d i a l B l u r I m a g e %
3684% %
3685% %
3686% %
3687%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3688*/
3689
3690static Image *ComputeRadialBlurImage(const Image *image,
3691 const ChannelType channel,const double angle,ExceptionInfo *exception)
3692{
3693 cl_command_queue
3694 queue;
3695
3696 cl_context
3697 context;
3698
3699 cl_float2
3700 blurCenter;
3701
3702 cl_float4
3703 biasPixel;
3704
3705 cl_int
3706 clStatus;
3707
3708 cl_mem
3709 cosThetaBuffer,
3710 filteredImageBuffer,
3711 imageBuffer,
3712 sinThetaBuffer;
3713
3714 cl_kernel
3715 radialBlurKernel;
3716
3717 cl_event
3718 event;
3719
3720 cl_uint
3721 event_count;
3722
3723 cl_event
3724 *events;
3725
3726 float
3727 blurRadius,
3728 *cosThetaPtr,
3729 offset,
3730 *sinThetaPtr,
3731 theta;
3732
3733 Image
3734 *filteredImage;
3735
3736 MagickBooleanType
3737 outputReady;
3738
3739 MagickCLEnv
3740 clEnv;
3741
3742 MagickPixelPacket
3743 bias;
3744
3745 size_t
3746 global_work_size[2];
3747
3748 unsigned int
3749 cossin_theta_size,
3750 i,
3751 matte;
3752
3753 outputReady = MagickFalse;
3754 context = NULL;
3755 filteredImage = NULL;
3756 imageBuffer = NULL;
3757 filteredImageBuffer = NULL;
3758 sinThetaBuffer = NULL;
3759 cosThetaBuffer = NULL;
3760 queue = NULL;
3761 radialBlurKernel = NULL;
3762
3763
3764 clEnv = GetDefaultOpenCLEnv();
3765 context = GetOpenCLContext(clEnv);
3766
3767 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
3768 if (filteredImage == (Image *) NULL)
3769 goto cleanup;
3770
3771 imageBuffer = GetAuthenticOpenCLBuffer(image, exception);
3772 if (imageBuffer == (cl_mem) NULL)
3773 {
3774 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3775 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3776 goto cleanup;
3777 }
3778 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage, exception);
3779 if (filteredImageBuffer == (cl_mem) NULL)
3780 {
3781 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
3782 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
3783 goto cleanup;
3784 }
3785
3786 blurCenter.s[0] = (float) (image->columns-1)/2.0;
3787 blurCenter.s[1] = (float) (image->rows-1)/2.0;
3788 blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]);
3789 cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL);
3790
3791 /* create a buffer for sin_theta and cos_theta */
3792 sinThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
3793 if (clStatus != CL_SUCCESS)
3794 {
3795 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3796 goto cleanup;
3797 }
3798 cosThetaBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
3799 if (clStatus != CL_SUCCESS)
3800 {
3801 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
3802 goto cleanup;
3803 }
3804
3805 queue = AcquireOpenCLCommandQueue(clEnv);
3806 events=GetOpenCLEvents(image,&event_count);
3807 /* this blocks, should be fixed it in the future */
3808 sinThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), event_count, events, NULL, &clStatus);
3809 events=(cl_event *) RelinquishMagickMemory(events);
3810 if (clStatus != CL_SUCCESS)
3811 {
3812 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
3813 goto cleanup;
3814 }
3815
3816 cosThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
3817 if (clStatus != CL_SUCCESS)
3818 {
3819 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
3820 goto cleanup;
3821 }
3822
3823 theta=DegreesToRadians(angle)/(MagickRealType) (cossin_theta_size-1);
3824 offset=theta*(MagickRealType) (cossin_theta_size-1)/2.0;
3825 for (i=0; i < (ssize_t) cossin_theta_size; i++)
3826 {
3827 cosThetaPtr[i]=(float)cos((double) (theta*i-offset));
3828 sinThetaPtr[i]=(float)sin((double) (theta*i-offset));
3829 }
3830
3831 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL);
3832 clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL);
3833 if (clStatus != CL_SUCCESS)
3834 {
3835 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
3836 goto cleanup;
3837 }
3838
3839 /* get the OpenCL kernel */
3840 radialBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RadialBlur");
3841 if (radialBlurKernel == NULL)
3842 {
3843 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
3844 goto cleanup;
3845 }
3846
3847
3848 /* set the kernel arguments */
3849 i = 0;
3850 clStatus=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
3851 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
3852
3853 GetMagickPixelPacket(image,&bias);
3854 biasPixel.s[0] = bias.red;
3855 biasPixel.s[1] = bias.green;
3856 biasPixel.s[2] = bias.blue;
3857 biasPixel.s[3] = bias.opacity;
3858 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_float4), &biasPixel);
3859 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(ChannelType), &channel);
3860
3861 matte = (image->matte != MagickFalse)?1:0;
3862 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &matte);
3863
3864 clStatus=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_float2), &blurCenter);
3865
3866 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer);
3867 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer);
3868 clStatus|=clEnv->library->clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size);
3869 if (clStatus != CL_SUCCESS)
3870 {
3871 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
3872 goto cleanup;
3873 }
3874
3875
3876 global_work_size[0] = image->columns;
3877 global_work_size[1] = image->rows;
3878 /* launch the kernel */
3879 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, radialBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
3880 if (clStatus != CL_SUCCESS)
3881 {
3882 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
3883 goto cleanup;
3884 }
3885 if (RecordProfileData(clEnv,RadialBlurKernel,event) == MagickFalse)
3886 {
3887 AddOpenCLEvent(image,event);
3888 AddOpenCLEvent(filteredImage,event);
3889 }
3890 clEnv->library->clReleaseEvent(event);
3891
3892 outputReady = MagickTrue;
3893
3894cleanup:
3895 OpenCLLogException(__FUNCTION__,__LINE__,exception);
3896
3897 if (imageBuffer != (cl_mem) NULL)
3898 clEnv->library->clReleaseMemObject(imageBuffer);
3899 if (filteredImageBuffer != (cl_mem) NULL)
3900 clEnv->library->clReleaseMemObject(filteredImageBuffer);
3901 if (sinThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(sinThetaBuffer);
3902 if (cosThetaBuffer!=NULL) clEnv->library->clReleaseMemObject(cosThetaBuffer);
3903 if (radialBlurKernel!=NULL) RelinquishOpenCLKernel(clEnv, radialBlurKernel);
3904 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
3905 if ((outputReady == MagickFalse) && (filteredImage != NULL))
3906 filteredImage=(Image *) DestroyImage(filteredImage);
3907 return filteredImage;
3908}
3909
3910MagickPrivate Image *AccelerateRadialBlurImage(const Image *image,
3911 const ChannelType channel,const double angle,ExceptionInfo *exception)
3912{
3913 Image
3914 *filteredImage;
3915
3916 assert(image != NULL);
3917 assert(exception != (ExceptionInfo *) NULL);
3918
3919 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
3920 (checkAccelerateCondition(image, channel) == MagickFalse))
3921 return NULL;
3922
3923 filteredImage=ComputeRadialBlurImage(image, channel, angle, exception);
3924 return filteredImage;
3925}
3926
3927/*
3928%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3929% %
3930% %
3931% %
3932% A c c e l e r a t e R e s i z e I m a g e %
3933% %
3934% %
3935% %
3936%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3937*/
3938
3939static MagickBooleanType resizeHorizontalFilter(const Image *image,
3940 const Image *filteredImage,cl_mem imageBuffer,const unsigned int imageColumns,
3941 const unsigned int imageRows,const unsigned int matte,cl_mem resizedImage,
3942 const unsigned int resizedColumns,const unsigned int resizedRows,
3943 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
3944 const float xFactor,MagickCLEnv clEnv,cl_command_queue queue,
3945 ExceptionInfo *exception)
3946{
3947 cl_kernel
3948 horizontalKernel;
3949
3950 cl_event
3951 event;
3952
3953 cl_int
3954 clStatus;
3955
3956 cl_uint
3957 event_count;
3958
3959 cl_event
3960 *events;
3961
3962 const unsigned int
3963 workgroupSize = 256;
3964
3965 float
3966 resizeFilterScale,
3967 resizeFilterSupport,
3968 resizeFilterWindowSupport,
3969 resizeFilterBlur,
3970 scale,
3971 support;
3972
3973 int
3974 cacheRangeStart,
3975 cacheRangeEnd,
3976 numCachedPixels,
3977 resizeFilterType,
3978 resizeWindowType;
3979
3980 MagickBooleanType
3981 status = MagickFalse;
3982
3983 size_t
3984 deviceLocalMemorySize,
3985 gammaAccumulatorLocalMemorySize,
3986 global_work_size[2],
3987 imageCacheLocalMemorySize,
3988 pixelAccumulatorLocalMemorySize,
3989 local_work_size[2],
3990 totalLocalMemorySize,
3991 weightAccumulatorLocalMemorySize;
3992
3993 unsigned int
3994 chunkSize,
3995 i,
3996 pixelPerWorkgroup;
3997
3998 horizontalKernel = NULL;
3999 status = MagickFalse;
4000
4001 /*
4002 Apply filter to resize vertically from image to resize image.
4003 */
4004 scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
4005 support=scale*GetResizeFilterSupport(resizeFilter);
4006 if (support < 0.5)
4007 {
4008 /*
4009 Support too small even for nearest neighbour: Reduce to point
4010 sampling.
4011 */
4012 support=(MagickRealType) 0.5;
4013 scale=1.0;
4014 }
4015 scale=PerceptibleReciprocal(scale);
4016
4017 if (resizedColumns < workgroupSize)
4018 {
4019 chunkSize = 32;
4020 pixelPerWorkgroup = 32;
4021 }
4022 else
4023 {
4024 chunkSize = workgroupSize;
4025 pixelPerWorkgroup = workgroupSize;
4026 }
4027
4028 /* get the local memory size supported by the device */
4029 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
4030
4031DisableMSCWarning(4127)
4032 while(1)
4033RestoreMSCWarning
4034 {
4035 /* calculate the local memory size needed per workgroup */
4036 cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
4037 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
4038 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
4039 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
4040 totalLocalMemorySize = imageCacheLocalMemorySize;
4041
4042 /* local size for the pixel accumulator */
4043 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
4044 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
4045
4046 /* local memory size for the weight accumulator */
4047 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
4048 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
4049
4050 /* local memory size for the gamma accumulator */
4051 if (matte == 0)
4052 gammaAccumulatorLocalMemorySize = sizeof(float);
4053 else
4054 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
4055 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
4056
4057 if (totalLocalMemorySize <= deviceLocalMemorySize)
4058 break;
4059 else
4060 {
4061 pixelPerWorkgroup = pixelPerWorkgroup/2;
4062 chunkSize = chunkSize/2;
4063 if (pixelPerWorkgroup == 0
4064 || chunkSize == 0)
4065 {
4066 /* quit, fallback to CPU */
4067 goto cleanup;
4068 }
4069 }
4070 }
4071
4072 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
4073 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
4074
4075 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter");
4076 if (horizontalKernel == NULL)
4077 {
4078 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4079 goto cleanup;
4080 }
4081
4082 i = 0;
4083 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&imageBuffer);
4084 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns);
4085 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows);
4086 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
4087 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
4088 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
4089
4090 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
4091 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
4092
4093 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
4094 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
4095 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
4096
4097 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
4098 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
4099
4100 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
4101 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
4102
4103 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
4104 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
4105
4106 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
4107 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
4108
4109
4110 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
4111 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
4112 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
4113 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
4114
4115
4116 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
4117 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
4118 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
4119
4120 if (clStatus != CL_SUCCESS)
4121 {
4122 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4123 goto cleanup;
4124 }
4125
4126 global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
4127 global_work_size[1] = resizedRows;
4128
4129 local_work_size[0] = workgroupSize;
4130 local_work_size[1] = 1;
4131 events=GetOpenCLEvents(image,&event_count);
4132 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, event_count, events, &event);
4133 events=(cl_event *) RelinquishMagickMemory(events);
4134 if (clStatus != CL_SUCCESS)
4135 {
4136 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4137 goto cleanup;
4138 }
4139 if (RecordProfileData(clEnv,ResizeHorizontalKernel,event) == MagickFalse)
4140 {
4141 AddOpenCLEvent(image,event);
4142 AddOpenCLEvent(filteredImage,event);
4143 }
4144 clEnv->library->clReleaseEvent(event);
4145 status = MagickTrue;
4146
4147
4148cleanup:
4149 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4150
4151 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
4152
4153 return(status);
4154}
4155
4156static MagickBooleanType resizeVerticalFilter(const Image *image,
4157 const Image *filteredImage,cl_mem imageBuffer,const unsigned int imageColumns,
4158 const unsigned int imageRows,const unsigned int matte,cl_mem resizedImage,
4159 const unsigned int resizedColumns,const unsigned int resizedRows,
4160 const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
4161 const float yFactor,MagickCLEnv clEnv,cl_command_queue queue,
4162 ExceptionInfo *exception)
4163{
4164 cl_kernel
4165 horizontalKernel;
4166
4167 cl_event
4168 event;
4169
4170 cl_int
4171 clStatus;
4172
4173 cl_uint
4174 event_count;
4175
4176 cl_event
4177 *events;
4178
4179 const unsigned int
4180 workgroupSize = 256;
4181
4182 float
4183 resizeFilterScale,
4184 resizeFilterSupport,
4185 resizeFilterWindowSupport,
4186 resizeFilterBlur,
4187 scale,
4188 support;
4189
4190 int
4191 cacheRangeStart,
4192 cacheRangeEnd,
4193 numCachedPixels,
4194 resizeFilterType,
4195 resizeWindowType;
4196
4197 MagickBooleanType
4198 status = MagickFalse;
4199
4200 size_t
4201 deviceLocalMemorySize,
4202 gammaAccumulatorLocalMemorySize,
4203 global_work_size[2],
4204 imageCacheLocalMemorySize,
4205 pixelAccumulatorLocalMemorySize,
4206 local_work_size[2],
4207 totalLocalMemorySize,
4208 weightAccumulatorLocalMemorySize;
4209
4210 unsigned int
4211 chunkSize,
4212 i,
4213 pixelPerWorkgroup;
4214
4215 horizontalKernel = NULL;
4216 status = MagickFalse;
4217
4218 /*
4219 Apply filter to resize vertically from image to resize image.
4220 */
4221 scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
4222 support=scale*GetResizeFilterSupport(resizeFilter);
4223 if (support < 0.5)
4224 {
4225 /*
4226 Support too small even for nearest neighbour: Reduce to point
4227 sampling.
4228 */
4229 support=(MagickRealType) 0.5;
4230 scale=1.0;
4231 }
4232 scale=PerceptibleReciprocal(scale);
4233
4234 if (resizedRows < workgroupSize)
4235 {
4236 chunkSize = 32;
4237 pixelPerWorkgroup = 32;
4238 }
4239 else
4240 {
4241 chunkSize = workgroupSize;
4242 pixelPerWorkgroup = workgroupSize;
4243 }
4244
4245 /* get the local memory size supported by the device */
4246 deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
4247
4248DisableMSCWarning(4127)
4249 while(1)
4250RestoreMSCWarning
4251 {
4252 /* calculate the local memory size needed per workgroup */
4253 cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
4254 cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
4255 numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
4256 imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
4257 totalLocalMemorySize = imageCacheLocalMemorySize;
4258
4259 /* local size for the pixel accumulator */
4260 pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
4261 totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
4262
4263 /* local memory size for the weight accumulator */
4264 weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
4265 totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
4266
4267 /* local memory size for the gamma accumulator */
4268 if (matte == 0)
4269 gammaAccumulatorLocalMemorySize = sizeof(float);
4270 else
4271 gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
4272 totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
4273
4274 if (totalLocalMemorySize <= deviceLocalMemorySize)
4275 break;
4276 else
4277 {
4278 pixelPerWorkgroup = pixelPerWorkgroup/2;
4279 chunkSize = chunkSize/2;
4280 if (pixelPerWorkgroup == 0
4281 || chunkSize == 0)
4282 {
4283 /* quit, fallback to CPU */
4284 goto cleanup;
4285 }
4286 }
4287 }
4288
4289 resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
4290 resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
4291
4292 horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter");
4293 if (horizontalKernel == NULL)
4294 {
4295 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4296 goto cleanup;
4297 }
4298
4299 i = 0;
4300 clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&imageBuffer);
4301 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns);
4302 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows);
4303 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
4304 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&yFactor);
4305 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
4306
4307 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
4308 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
4309
4310 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
4311 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
4312 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
4313
4314 resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
4315 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
4316
4317 resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
4318 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
4319
4320 resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
4321 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
4322
4323 resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
4324 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
4325
4326
4327 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
4328 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
4329 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
4330 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
4331
4332
4333 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
4334 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
4335 clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
4336
4337 if (clStatus != CL_SUCCESS)
4338 {
4339 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4340 goto cleanup;
4341 }
4342
4343 global_work_size[0] = resizedColumns;
4344 global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
4345
4346 local_work_size[0] = 1;
4347 local_work_size[1] = workgroupSize;
4348 events=GetOpenCLEvents(image,&event_count);
4349 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, event_count, events, &event);
4350 events=(cl_event *) RelinquishMagickMemory(events);
4351 if (clStatus != CL_SUCCESS)
4352 {
4353 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4354 goto cleanup;
4355 }
4356 if (RecordProfileData(clEnv,ResizeVerticalKernel,event) == MagickFalse)
4357 {
4358 AddOpenCLEvent(image,event);
4359 AddOpenCLEvent(filteredImage,event);
4360 }
4361 clEnv->library->clReleaseEvent(event);
4362 status = MagickTrue;
4363
4364
4365cleanup:
4366 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4367
4368 if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
4369
4370 return(status);
4371}
4372
4373static Image *ComputeResizeImage(const Image* image,
4374 const size_t resizedColumns,const size_t resizedRows,
4375 const ResizeFilter *resizeFilter,ExceptionInfo *exception)
4376{
4377 cl_command_queue
4378 queue;
4379
4380 cl_int
4381 clStatus;
4382
4383 cl_context
4384 context;
4385
4386 cl_mem
4387 cubicCoefficientsBuffer,
4388 filteredImageBuffer,
4389 imageBuffer,
4390 tempImageBuffer;
4391
4392 const MagickRealType
4393 *resizeFilterCoefficient;
4394
4395 float
4396 coefficientBuffer[7],
4397 xFactor,
4398 yFactor;
4399
4400 MagickBooleanType
4401 outputReady,
4402 status;
4403
4404 MagickCLEnv
4405 clEnv;
4406
4407 MagickSizeType
4408 length;
4409
4410 Image
4411 *filteredImage;
4412
4413 size_t
4414 i;
4415
4416 outputReady = MagickFalse;
4417 filteredImage = NULL;
4418 clEnv = NULL;
4419 context = NULL;
4420 imageBuffer = NULL;
4421 tempImageBuffer = NULL;
4422 filteredImageBuffer = NULL;
4423 cubicCoefficientsBuffer = NULL;
4424 queue = NULL;
4425
4426 clEnv = GetDefaultOpenCLEnv();
4427 context = GetOpenCLContext(clEnv);
4428 queue = AcquireOpenCLCommandQueue(clEnv);
4429
4430 filteredImage=CloneImage(image,resizedColumns,resizedRows,MagickTrue,exception);
4431 if (filteredImage == (Image *) NULL)
4432 goto cleanup;
4433
4434 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
4435 if (imageBuffer == (cl_mem) NULL)
4436 {
4437 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4438 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
4439 goto cleanup;
4440 }
4441 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,exception);
4442 if (filteredImageBuffer == (cl_mem) NULL)
4443 {
4444 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4445 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
4446 goto cleanup;
4447 }
4448
4449 resizeFilterCoefficient=GetResizeFilterCoefficient(resizeFilter);
4450 for (i = 0; i < 7; i++)
4451 coefficientBuffer[i]=(float) resizeFilterCoefficient[i];
4452
4453 cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(coefficientBuffer), coefficientBuffer, &clStatus);
4454 if (clStatus != CL_SUCCESS)
4455 {
4456 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4457 goto cleanup;
4458 }
4459
4460 xFactor=(float) resizedColumns/(float) image->columns;
4461 yFactor=(float) resizedRows/(float) image->rows;
4462 if (xFactor > yFactor)
4463 {
4464
4465 length = resizedColumns*image->rows;
4466 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
4467 if (clStatus != CL_SUCCESS)
4468 {
4469 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4470 goto cleanup;
4471 }
4472
4473 status = resizeHorizontalFilter(image,filteredImage,imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, (image->matte != MagickFalse)?1:0
4474 , tempImageBuffer, (unsigned int) resizedColumns, (unsigned int) image->rows
4475 , resizeFilter, cubicCoefficientsBuffer
4476 , xFactor, clEnv, queue, exception);
4477 if (status != MagickTrue)
4478 goto cleanup;
4479
4480 status = resizeVerticalFilter(image,filteredImage,tempImageBuffer, (unsigned int) resizedColumns, (unsigned int) image->rows, (image->matte != MagickFalse)?1:0
4481 , filteredImageBuffer, (unsigned int) resizedColumns, (unsigned int) resizedRows
4482 , resizeFilter, cubicCoefficientsBuffer
4483 , yFactor, clEnv, queue, exception);
4484 if (status != MagickTrue)
4485 goto cleanup;
4486 }
4487 else
4488 {
4489 length = image->columns*resizedRows;
4490 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus);
4491 if (clStatus != CL_SUCCESS)
4492 {
4493 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4494 goto cleanup;
4495 }
4496
4497 status = resizeVerticalFilter(image,filteredImage,imageBuffer, (unsigned int) image->columns, (unsigned int) image->rows, (image->matte != MagickFalse)?1:0
4498 , tempImageBuffer, (unsigned int) image->columns, (unsigned int) resizedRows
4499 , resizeFilter, cubicCoefficientsBuffer
4500 , yFactor, clEnv, queue, exception);
4501 if (status != MagickTrue)
4502 goto cleanup;
4503
4504 status = resizeHorizontalFilter(image,filteredImage,tempImageBuffer, (unsigned int) image->columns, (unsigned int) resizedRows, (image->matte != MagickFalse)?1:0
4505 , filteredImageBuffer, (unsigned int) resizedColumns, (unsigned int) resizedRows
4506 , resizeFilter, cubicCoefficientsBuffer
4507 , xFactor, clEnv, queue, exception);
4508 if (status != MagickTrue)
4509 goto cleanup;
4510 }
4511 outputReady=MagickTrue;
4512
4513cleanup:
4514 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4515
4516 if (imageBuffer != (cl_mem) NULL)
4517 clEnv->library->clReleaseMemObject(imageBuffer);
4518 if (filteredImageBuffer != (cl_mem) NULL)
4519 clEnv->library->clReleaseMemObject(filteredImageBuffer);
4520 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
4521 if (cubicCoefficientsBuffer!=NULL) clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer);
4522 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4523 if ((outputReady == MagickFalse) && (filteredImage != NULL))
4524 filteredImage=(Image *) DestroyImage(filteredImage);
4525 return(filteredImage);
4526}
4527
4528static MagickBooleanType gpuSupportedResizeWeighting(
4529 ResizeWeightingFunctionType f)
4530{
4531 unsigned int
4532 i;
4533
4534 for (i = 0; ;i++)
4535 {
4536 if (supportedResizeWeighting[i] == LastWeightingFunction)
4537 break;
4538 if (supportedResizeWeighting[i] == f)
4539 return(MagickTrue);
4540 }
4541 return(MagickFalse);
4542}
4543
4544MagickPrivate Image *AccelerateResizeImage(const Image *image,
4545 const size_t resizedColumns,const size_t resizedRows,
4546 const ResizeFilter *resizeFilter,ExceptionInfo *exception)
4547{
4548 Image
4549 *filteredImage;
4550
4551 assert(image != NULL);
4552 assert(exception != (ExceptionInfo *) NULL);
4553
4554 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
4555 (checkAccelerateCondition(image, AllChannels) == MagickFalse))
4556 return NULL;
4557
4558 if (gpuSupportedResizeWeighting(GetResizeFilterWeightingType(resizeFilter)) == MagickFalse ||
4559 gpuSupportedResizeWeighting(GetResizeFilterWindowWeightingType(resizeFilter)) == MagickFalse)
4560 return NULL;
4561
4562 filteredImage=ComputeResizeImage(image,resizedColumns,resizedRows,resizeFilter,exception);
4563 return(filteredImage);
4564}
4565
4566/*
4567%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4568% %
4569% %
4570% %
4571% A c c e l e r a t e U n s h a r p M a s k I m a g e %
4572% %
4573% %
4574% %
4575%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
4576*/
4577
4578static Image *ComputeUnsharpMaskImage(const Image *image,
4579 const ChannelType channel,const double radius,const double sigma,
4580 const double gain,const double threshold,ExceptionInfo *exception)
4581{
4582 char
4583 geometry[MaxTextExtent];
4584
4585 cl_command_queue
4586 queue;
4587
4588 cl_context
4589 context;
4590
4591 cl_event
4592 event;
4593
4594 cl_int
4595 clStatus;
4596
4597 cl_kernel
4598 blurRowKernel,
4599 unsharpMaskBlurColumnKernel;
4600
4601 cl_mem
4602 filteredImageBuffer,
4603 imageBuffer,
4604 imageKernelBuffer,
4605 tempImageBuffer;
4606
4607 cl_uint
4608 event_count;
4609
4610 cl_event
4611 *events;
4612
4613 float
4614 fGain,
4615 fThreshold,
4616 *kernelBufferPtr;
4617
4618 Image
4619 *filteredImage;
4620
4621 int
4622 chunkSize;
4623
4624 KernelInfo
4625 *kernel;
4626
4627 MagickBooleanType
4628 outputReady;
4629
4630 MagickCLEnv
4631 clEnv;
4632
4633 MagickSizeType
4634 length;
4635
4636 unsigned int
4637 imageColumns,
4638 imageRows,
4639 kernelWidth;
4640
4641 size_t
4642 i;
4643
4644 clEnv = NULL;
4645 filteredImage = NULL;
4646 kernel = NULL;
4647 context = NULL;
4648 imageBuffer = NULL;
4649 filteredImageBuffer = NULL;
4650 tempImageBuffer = NULL;
4651 imageKernelBuffer = NULL;
4652 blurRowKernel = NULL;
4653 unsharpMaskBlurColumnKernel = NULL;
4654 queue = NULL;
4655 outputReady = MagickFalse;
4656
4657 clEnv = GetDefaultOpenCLEnv();
4658 context = GetOpenCLContext(clEnv);
4659 queue = AcquireOpenCLCommandQueue(clEnv);
4660
4661 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
4662 if (filteredImage == (Image *) NULL)
4663 goto cleanup;
4664
4665 imageBuffer=GetAuthenticOpenCLBuffer(image,exception);
4666 if (imageBuffer == (cl_mem) NULL)
4667 {
4668 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4669 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
4670 goto cleanup;
4671 }
4672 filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,exception);
4673 if (filteredImageBuffer == (cl_mem) NULL)
4674 {
4675 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4676 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
4677 goto cleanup;
4678 }
4679
4680 /* create the blur kernel */
4681 {
4682 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
4683 kernel=AcquireKernelInfo(geometry);
4684 if (kernel == (KernelInfo *) NULL)
4685 {
4686 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
4687 goto cleanup;
4688 }
4689
4690 kernelBufferPtr=(float *) AcquireQuantumMemory(kernel->width,sizeof(float));
4691 if (kernelBufferPtr == (float *) NULL)
4692 {
4693 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Memory allocation failed.",".");
4694 goto cleanup;
4695 }
4696 for (i = 0; i < kernel->width; i++)
4697 kernelBufferPtr[i]=(float) kernel->values[i];
4698
4699 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, kernel->width * sizeof(float), kernelBufferPtr, &clStatus);
4700 kernelBufferPtr=(float *) RelinquishMagickMemory(kernelBufferPtr);
4701 if (clStatus != CL_SUCCESS)
4702 {
4703 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4704 goto cleanup;
4705 }
4706 }
4707
4708 {
4709 /* create temp buffer */
4710 {
4711 length = image->columns * image->rows;
4712 tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
4713 if (clStatus != CL_SUCCESS)
4714 {
4715 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4716 goto cleanup;
4717 }
4718 }
4719
4720 /* get the opencl kernel */
4721 {
4722 blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow");
4723 if (blurRowKernel == NULL)
4724 {
4725 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4726 goto cleanup;
4727 };
4728
4729 unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn");
4730 if (unsharpMaskBlurColumnKernel == NULL)
4731 {
4732 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4733 goto cleanup;
4734 };
4735 }
4736
4737 {
4738 chunkSize = 256;
4739
4740 imageColumns = (unsigned int) image->columns;
4741 imageRows = (unsigned int) image->rows;
4742
4743 kernelWidth = (unsigned int) kernel->width;
4744
4745 /* set the kernel arguments */
4746 i = 0;
4747 clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4748 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
4749 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&channel);
4750 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
4751 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
4752 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
4753 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
4754 clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
4755 if (clStatus != CL_SUCCESS)
4756 {
4757 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4758 goto cleanup;
4759 }
4760 }
4761
4762 /* launch the kernel */
4763 {
4764 size_t gsize[2];
4765 size_t wsize[2];
4766
4767 gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
4768 gsize[1] = image->rows;
4769 wsize[0] = chunkSize;
4770 wsize[1] = 1;
4771
4772 events=GetOpenCLEvents(image,&event_count);
4773 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, event_count, events, NULL);
4774 events=(cl_event *) RelinquishMagickMemory(events);
4775 if (clStatus != CL_SUCCESS)
4776 {
4777 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4778 goto cleanup;
4779 }
4780 }
4781
4782
4783 {
4784 chunkSize = 256;
4785 imageColumns = (unsigned int) image->columns;
4786 imageRows = (unsigned int) image->rows;
4787 kernelWidth = (unsigned int) kernel->width;
4788 fGain = (float) gain;
4789 fThreshold = (float) threshold;
4790
4791 i = 0;
4792 clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4793 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
4794 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
4795 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
4796 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
4797 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
4798 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
4799 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&channel);
4800 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
4801 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
4802 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
4803 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
4804
4805 if (clStatus != CL_SUCCESS)
4806 {
4807 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
4808 goto cleanup;
4809 }
4810 }
4811
4812 /* launch the kernel */
4813 {
4814 size_t gsize[2];
4815 size_t wsize[2];
4816
4817 gsize[0] = image->columns;
4818 gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
4819 wsize[0] = 1;
4820 wsize[1] = chunkSize;
4821
4822 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, event_count, events, &event);
4823 if (clStatus != CL_SUCCESS)
4824 {
4825 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
4826 goto cleanup;
4827 }
4828 if (RecordProfileData(clEnv,UnsharpMaskBlurColumnKernel,event) == MagickFalse)
4829 {
4830 AddOpenCLEvent(image,event);
4831 AddOpenCLEvent(filteredImage,event);
4832 }
4833 clEnv->library->clReleaseEvent(event);
4834 }
4835
4836 }
4837
4838 outputReady=MagickTrue;
4839
4840cleanup:
4841 OpenCLLogException(__FUNCTION__,__LINE__,exception);
4842
4843 if (imageBuffer != (cl_mem) NULL)
4844 clEnv->library->clReleaseMemObject(imageBuffer);
4845 if (filteredImageBuffer != (cl_mem) NULL)
4846 clEnv->library->clReleaseMemObject(filteredImageBuffer);
4847 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
4848 if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
4849 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
4850 if (blurRowKernel!=NULL) RelinquishOpenCLKernel(clEnv, blurRowKernel);
4851 if (unsharpMaskBlurColumnKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskBlurColumnKernel);
4852 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
4853 if ((outputReady == MagickFalse) && (filteredImage != NULL))
4854 filteredImage=(Image *) DestroyImage(filteredImage);
4855 return(filteredImage);
4856}
4857
4858static Image *ComputeUnsharpMaskImageSingle(const Image *image,
4859 const double radius,const double sigma,const double gain,
4860 const double threshold,int blurOnly, ExceptionInfo *exception)
4861{
4862 char
4863 geometry[MaxTextExtent];
4864
4865 cl_command_queue
4866 queue;
4867
4868 cl_context
4869 context;
4870
4871 cl_int
4872 justBlur,
4873 clStatus;
4874
4875 cl_kernel
4876 unsharpMaskKernel;
4877
4878 cl_event
4879 event;
4880
4881 cl_mem
4882 filteredImageBuffer,
4883 imageBuffer,
4884 imageKernelBuffer;
4885
4886 cl_event
4887 *events;
4888
4889 float
4890 fGain,
4891 fThreshold;
4892
4893 Image
4894 *filteredImage;
4895
4896 KernelInfo
4897 *kernel;
4898
4899 MagickBooleanType
4900 outputReady;
4901
4902 MagickCLEnv
4903 clEnv;
4904
4905 unsigned int
4906 event_count,
4907 i,
4908 imageColumns,
4909 imageRows,
4910 kernelWidth;
4911
4912 clEnv = NULL;
4913 filteredImage = NULL;
4914 kernel = NULL;
4915 context = NULL;
4916 imageBuffer = NULL;
4917 filteredImageBuffer = NULL;
4918 imageKernelBuffer = NULL;
4919 unsharpMaskKernel = NULL;
4920 queue = NULL;
4921 outputReady = MagickFalse;
4922
4923 clEnv = GetDefaultOpenCLEnv();
4924 context = GetOpenCLContext(clEnv);
4925 queue = AcquireOpenCLCommandQueue(clEnv);
4926
4927 filteredImage = CloneImage(image,0,0,MagickTrue,exception);
4928 if (filteredImage == (Image *) NULL)
4929 goto cleanup;
4930
4931 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
4932 if (imageBuffer == (cl_mem) NULL)
4933 {
4934 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4935 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
4936 goto cleanup;
4937 }
4938 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
4939 if (filteredImageBuffer == (cl_mem) NULL)
4940 {
4941 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
4942 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
4943 goto cleanup;
4944 }
4945
4946 /* create the blur kernel */
4947 {
4948 (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
4949 kernel=AcquireKernelInfo(geometry);
4950 if (kernel == (KernelInfo *) NULL)
4951 {
4952 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
4953 goto cleanup;
4954 }
4955
4956 {
4957 float *kernelBufferPtr = (float *) AcquireQuantumMemory(kernel->width, sizeof(float));
4958 for (i = 0; i < kernel->width; i++)
4959 kernelBufferPtr[i] = (float)kernel->values[i];
4960
4961 imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, kernel->width * sizeof(float), kernelBufferPtr, &clStatus);
4962 RelinquishMagickMemory(kernelBufferPtr);
4963 if (clStatus != CL_SUCCESS)
4964 {
4965 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
4966 goto cleanup;
4967 }
4968 }
4969 }
4970
4971 {
4972 /* get the opencl kernel */
4973 {
4974 unsharpMaskKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMask");
4975 if (unsharpMaskKernel == NULL)
4976 {
4977 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
4978 goto cleanup;
4979 };
4980 }
4981
4982 {
4983 imageColumns = (unsigned int) image->columns;
4984 imageRows = (unsigned int) image->rows;
4985 kernelWidth = (unsigned int) kernel->width;
4986 fGain = (float) gain;
4987 fThreshold = (float) threshold;
4988 justBlur = blurOnly;
4989
4990 /* set the kernel arguments */
4991 i = 0;
4992 clStatus=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
4993 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
4994 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
4995 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
4996 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
4997 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&imageRows);
4998 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernel->width)),(void *) NULL);
4999 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain);
5000 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold);
5001 clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&justBlur);
5002 if (clStatus != CL_SUCCESS)
5003 {
5004 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
5005 goto cleanup;
5006 }
5007 }
5008
5009 /* launch the kernel */
5010 {
5011 size_t gsize[2];
5012 size_t wsize[2];
5013
5014 gsize[0] = ((image->columns + 7) / 8) * 8;
5015 gsize[1] = ((image->rows + 31) / 32) * 32;
5016 wsize[0] = 8;
5017 wsize[1] = 32;
5018
5019 events=GetOpenCLEvents(image,&event_count);
5020 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, event_count, events, &event);
5021 events=(cl_event *) RelinquishMagickMemory(events);
5022 if (clStatus != CL_SUCCESS)
5023 {
5024 (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5025 goto cleanup;
5026 }
5027 if (RecordProfileData(clEnv,UnsharpMaskKernel,event) == MagickFalse)
5028 {
5029 AddOpenCLEvent(image,event);
5030 AddOpenCLEvent(filteredImage, event);
5031 }
5032 clEnv->library->clReleaseEvent(event);
5033 }
5034 }
5035
5036 outputReady=MagickTrue;
5037
5038cleanup:
5039 OpenCLLogException(__FUNCTION__,__LINE__,exception);
5040
5041 if (imageBuffer != (cl_mem) NULL)
5042 clEnv->library->clReleaseMemObject(imageBuffer);
5043 if (filteredImageBuffer != (cl_mem) NULL)
5044 clEnv->library->clReleaseMemObject(filteredImageBuffer);
5045 if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
5046 if (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
5047 if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskKernel);
5048 if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
5049 if ((outputReady == MagickFalse) && (filteredImage != NULL))
5050 filteredImage=(Image *) DestroyImage(filteredImage);
5051 return(filteredImage);
5052}
5053
5054MagickPrivate Image *AccelerateUnsharpMaskImage(const Image *image,
5055 const ChannelType channel,const double radius,const double sigma,
5056 const double gain,const double threshold,ExceptionInfo *exception)
5057{
5058 Image
5059 *filteredImage;
5060
5061 assert(image != NULL);
5062 assert(exception != (ExceptionInfo *) NULL);
5063
5064 if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
5065 (checkAccelerateCondition(image, channel) == MagickFalse))
5066 return NULL;
5067
5068 if (radius < 12.1)
5069 filteredImage = ComputeUnsharpMaskImageSingle(image,radius,sigma,gain,threshold, 0, exception);
5070 else
5071 filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
5072
5073 return(filteredImage);
5074}
5075
5076static Image *ComputeWaveletDenoiseImage(const Image *image,
5077 const double threshold,ExceptionInfo *exception)
5078{
5079 cl_command_queue
5080 queue;
5081
5082 cl_context
5083 context;
5084
5085 cl_int
5086 clStatus;
5087
5088 cl_kernel
5089 denoiseKernel;
5090
5091 cl_event
5092 event;
5093
5094 cl_mem
5095 filteredImageBuffer,
5096 imageBuffer;
5097
5098 cl_event
5099 *events;
5100
5101 Image
5102 *filteredImage;
5103
5104 MagickBooleanType
5105 outputReady;
5106
5107 MagickCLEnv
5108 clEnv;
5109
5110 unsigned int
5111 event_count,
5112 i,
5113 passes;
5114
5115 clEnv = NULL;
5116 filteredImage = NULL;
5117 context = NULL;
5118 imageBuffer = NULL;
5119 filteredImageBuffer = NULL;
5120 denoiseKernel = NULL;
5121 queue = NULL;
5122 outputReady = MagickFalse;
5123
5124 clEnv = GetDefaultOpenCLEnv();
5125
5126 /* Work around an issue on low end Intel devices */
5127 if (paramMatchesValue(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE_NAME,
5128 "Intel(R) HD Graphics",exception) != MagickFalse)
5129 goto cleanup;
5130
5131 context = GetOpenCLContext(clEnv);
5132 queue = AcquireOpenCLCommandQueue(clEnv);
5133
5134 filteredImage = CloneImage(image,0,0,MagickTrue, exception);
5135 if (filteredImage == (Image *) NULL)
5136 goto cleanup;
5137
5138 imageBuffer = GetAuthenticOpenCLBuffer(image,exception);
5139 if (imageBuffer == (cl_mem) NULL)
5140 {
5141 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
5142 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
5143 goto cleanup;
5144 }
5145 filteredImageBuffer = GetAuthenticOpenCLBuffer(filteredImage,exception);
5146 if (filteredImageBuffer == (cl_mem) NULL)
5147 {
5148 (void) OpenCLThrowMagickException(exception,GetMagickModule(),
5149 ResourceLimitWarning,"GetAuthenticOpenCLBuffer failed.",".");
5150 goto cleanup;
5151 }
5152
5153 /* get the opencl kernel */
5154 denoiseKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "WaveletDenoise");
5155 if (denoiseKernel == NULL)
5156 {
5157 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
5158 goto cleanup;
5159 };
5160
5161 /*
5162 Process image.
5163 */
5164 {
5165 int x;
5166 const int PASSES = 5;
5167 cl_int width = (cl_int)image->columns;
5168 cl_int height = (cl_int)image->rows;
5169 cl_float thresh = threshold;
5170
5171 passes = (((1.0f * image->columns) * image->rows) + 1999999.0f) / 2000000.0f;
5172 passes = (passes < 1) ? 1 : passes;
5173
5174 /* set the kernel arguments */
5175 i = 0;
5176 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_mem), (void *)&imageBuffer);
5177 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_mem), (void *)&filteredImageBuffer);
5178 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_float), (void *)&thresh);
5179 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_int), (void *)&PASSES);
5180 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_int), (void *)&width);
5181 clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_int), (void *)&height);
5182
5183 for (x = 0; x < passes; ++x)
5184 {
5185 const int TILESIZE = 64;
5186 const int PAD = 1 << (PASSES - 1);
5187 const int SIZE = TILESIZE - 2 * PAD;
5188
5189 size_t gsize[2];
5190 size_t wsize[2];
5191 size_t goffset[2];
5192
5193 gsize[0] = ((width + (SIZE - 1)) / SIZE) * TILESIZE;
5194 gsize[1] = ((((height + (SIZE - 1)) / SIZE) + passes - 1) / passes) * 4;
5195 wsize[0] = TILESIZE;
5196 wsize[1] = 4;
5197 goffset[0] = 0;
5198 goffset[1] = x * gsize[1];
5199
5200 events=GetOpenCLEvents(image,&event_count);
5201 clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, denoiseKernel, 2, goffset, gsize, wsize, event_count, events, &event);
5202 events=(cl_event *) RelinquishMagickMemory(events);
5203 if (clStatus != CL_SUCCESS)
5204 {
5205 (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
5206 goto cleanup;
5207 }
5208 clEnv->library->clFlush(queue);
5209 if (RecordProfileData(clEnv, WaveletDenoiseKernel, event) == MagickFalse)
5210 {
5211 AddOpenCLEvent(image, event);
5212 AddOpenCLEvent(filteredImage, event);
5213 }
5214 clEnv->library->clReleaseEvent(event);
5215 }
5216 }
5217
5218 outputReady=MagickTrue;
5219
5220cleanup:
5221 OpenCLLogException(__FUNCTION__, __LINE__, exception);
5222
5223 if (imageBuffer != (cl_mem) NULL)
5224 clEnv->library->clReleaseMemObject(imageBuffer);
5225 if (filteredImageBuffer != (cl_mem) NULL)
5226 clEnv->library->clReleaseMemObject(filteredImageBuffer);
5227 if (denoiseKernel != NULL)
5228 RelinquishOpenCLKernel(clEnv, denoiseKernel);
5229 if (queue != NULL)
5230 RelinquishOpenCLCommandQueue(clEnv, queue);
5231 if ((outputReady == MagickFalse) && (filteredImage != NULL))
5232 filteredImage=(Image *) DestroyImage(filteredImage);
5233 return(filteredImage);
5234}
5235
5236MagickPrivate Image *AccelerateWaveletDenoiseImage(const Image *image,
5237 const double threshold,ExceptionInfo *exception)
5238{
5239 Image
5240 *filteredImage;
5241
5242 assert(image != NULL);
5243 assert(exception != (ExceptionInfo *)NULL);
5244
5245 if ((checkAccelerateCondition(image,DefaultChannels) == MagickFalse) ||
5246 (checkOpenCLEnvironment(exception) == MagickFalse))
5247 return (Image *) NULL;
5248
5249 filteredImage=ComputeWaveletDenoiseImage(image,threshold,exception);
5250
5251 return(filteredImage);
5252}
5253
5254#endif /* MAGICKCORE_OPENCL_SUPPORT */
_ExceptionInfo
Definition exception.h:103
_Image
Definition image.h:134
_MagickCLEnv
Definition opencl.c:3058
_MagickPixelPacket
Definition pixel.h:105
_OffsetInfo
Definition geometry.h:117
_PixelPacket
Definition pixel.h:132
_RandomInfo
Definition random.c:82
_ResizeFilter
Definition resize.c:88