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