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