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