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