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