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