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