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