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