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