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