MagickCore  7.0.7
Convert, Edit, Or Compose Bitmap Images
opencl.c
Go to the documentation of this file.
1 /*
2 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3 % %
4 % %
5 % %
6 % OOO PPPP EEEEE N N CCCC L %
7 % O O P P E NN N C L %
8 % O O PPPP EEE N N N C L %
9 % O O P E N NN C L %
10 % OOO P EEEEE N N CCCC LLLLL %
11 % %
12 % %
13 % MagickCore OpenCL Methods %
14 % %
15 % Software Design %
16 % Cristy %
17 % March 2000 %
18 % %
19 % %
20 % Copyright 1999-2018 ImageMagick Studio LLC, a non-profit organization %
21 % dedicated to making software imaging solutions freely available. %
22 % %
23 % You may not use this file except in compliance with the License. You may %
24 % obtain a copy of the License at %
25 % %
26 % https://www.imagemagick.org/script/license.php %
27 % %
28 % Unless required by applicable law or agreed to in writing, software %
29 % distributed under the License is distributed on an "AS IS" BASIS, %
30 % WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. %
31 % See the License for the specific language governing permissions and %
32 % limitations under the License. %
33 % %
34 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
35 %
36 %
37 %
38 */
39 
40 /*
41  Include declarations.
42 */
43 #include "MagickCore/studio.h"
44 #include "MagickCore/artifact.h"
45 #include "MagickCore/cache.h"
47 #include "MagickCore/color.h"
48 #include "MagickCore/compare.h"
49 #include "MagickCore/constitute.h"
50 #include "MagickCore/configure.h"
51 #include "MagickCore/distort.h"
52 #include "MagickCore/draw.h"
53 #include "MagickCore/effect.h"
54 #include "MagickCore/exception.h"
56 #include "MagickCore/fx.h"
57 #include "MagickCore/gem.h"
58 #include "MagickCore/geometry.h"
59 #include "MagickCore/image.h"
61 #include "MagickCore/layer.h"
63 #include "MagickCore/memory_.h"
65 #include "MagickCore/monitor.h"
66 #include "MagickCore/montage.h"
67 #include "MagickCore/morphology.h"
68 #include "MagickCore/nt-base.h"
70 #include "MagickCore/opencl.h"
72 #include "MagickCore/option.h"
73 #include "MagickCore/policy.h"
74 #include "MagickCore/property.h"
75 #include "MagickCore/quantize.h"
76 #include "MagickCore/quantum.h"
77 #include "MagickCore/random_.h"
79 #include "MagickCore/resample.h"
80 #include "MagickCore/resource_.h"
81 #include "MagickCore/splay-tree.h"
82 #include "MagickCore/semaphore.h"
83 #include "MagickCore/statistic.h"
84 #include "MagickCore/string_.h"
86 #include "MagickCore/token.h"
87 #include "MagickCore/utility.h"
89 
90 #if defined(MAGICKCORE_OPENCL_SUPPORT)
91 #if defined(MAGICKCORE_LTDL_DELEGATE)
92 #include "ltdl.h"
93 #endif
94 
95 #ifndef MAGICKCORE_WINDOWS_SUPPORT
96 #include <dlfcn.h>
97 #endif
98 
99 #ifdef MAGICKCORE_HAVE_OPENCL_CL_H
100 #define MAGICKCORE_OPENCL_MACOSX 1
101 #endif
102 
103 /*
104  Define declarations.
105 */
106 #define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"
107 
108 /*
109  Typedef declarations.
110 */
111 typedef struct
112 {
113  long long freq;
114  long long clocks;
115  long long start;
116 } AccelerateTimer;
117 
118 typedef struct
119 {
120  char
121  *name,
122  *platform_name,
123  *vendor_name,
124  *version;
125 
126  cl_uint
127  max_clock_frequency,
128  max_compute_units;
129 
130  double
131  score;
132 } MagickCLDeviceBenchmark;
133 
134 /*
135  Forward declarations.
136 */
137 
138 static MagickBooleanType
139  HasOpenCLDevices(MagickCLEnv,ExceptionInfo *),
140  LoadOpenCLLibrary(void);
141 
142 static MagickCLDevice
143  RelinquishMagickCLDevice(MagickCLDevice);
144 
145 static MagickCLEnv
146  RelinquishMagickCLEnv(MagickCLEnv);
147 
148 static void
149  BenchmarkOpenCLDevices(MagickCLEnv);
150 
151 extern const char
152  *accelerateKernels, *accelerateKernels2;
153 
154 /* OpenCL library */
155 MagickLibrary
156  *openCL_library;
157 
158 /* Default OpenCL environment */
159 MagickCLEnv
160  default_CLEnv;
162  test_thread_id=0;
164  *openCL_lock;
165 
166 /* Cached location of the OpenCL cache files */
167 char
168  *cache_directory;
170  *cache_directory_lock;
171 
172 static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
173  MagickCLDevice b)
174 {
175  if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
176  (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
177  (LocaleCompare(a->name,b->name) == 0) &&
178  (LocaleCompare(a->version,b->version) == 0) &&
179  (a->max_clock_frequency == b->max_clock_frequency) &&
180  (a->max_compute_units == b->max_compute_units))
181  return(MagickTrue);
182 
183  return(MagickFalse);
184 }
185 
186 static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
187  MagickCLDeviceBenchmark *b)
188 {
189  if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
190  (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
191  (LocaleCompare(a->name,b->name) == 0) &&
192  (LocaleCompare(a->version,b->version) == 0) &&
193  (a->max_clock_frequency == b->max_clock_frequency) &&
194  (a->max_compute_units == b->max_compute_units))
195  return(MagickTrue);
196 
197  return(MagickFalse);
198 }
199 
200 static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
201 {
202  size_t
203  i;
204 
205  if (clEnv->devices != (MagickCLDevice *) NULL)
206  {
207  for (i = 0; i < clEnv->number_devices; i++)
208  clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
209  clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
210  }
211  clEnv->number_devices=0;
212 }
213 
214 static inline MagickBooleanType MagickCreateDirectory(const char *path)
215 {
216  int
217  status;
218 
219 #ifdef MAGICKCORE_WINDOWS_SUPPORT
220  status=mkdir(path);
221 #else
222  status=mkdir(path, 0777);
223 #endif
224  return(status == 0 ? MagickTrue : MagickFalse);
225 }
226 
227 static inline void InitAccelerateTimer(AccelerateTimer *timer)
228 {
229 #ifdef _WIN32
230  QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
231 #else
232  timer->freq=(long long)1.0E3;
233 #endif
234  timer->clocks=0;
235  timer->start=0;
236 }
237 
238 static inline double ReadAccelerateTimer(AccelerateTimer *timer)
239 {
240  return (double)timer->clocks/(double)timer->freq;
241 }
242 
243 static inline void StartAccelerateTimer(AccelerateTimer* timer)
244 {
245 #ifdef _WIN32
246  QueryPerformanceCounter((LARGE_INTEGER*)&timer->start);
247 #else
248  struct timeval
249  s;
250  gettimeofday(&s,0);
251  timer->start=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
252  (long long)1.0E3;
253 #endif
254 }
255 
256 static inline void StopAccelerateTimer(AccelerateTimer *timer)
257 {
258  long long
259  n;
260 
261  n=0;
262 #ifdef _WIN32
263  QueryPerformanceCounter((LARGE_INTEGER*)&(n));
264 #else
265  struct timeval
266  s;
267  gettimeofday(&s,0);
268  n=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
269  (long long)1.0E3;
270 #endif
271  n-=timer->start;
272  timer->start=0;
273  timer->clocks+=n;
274 }
275 
276 static const char *GetOpenCLCacheDirectory()
277 {
278  if (cache_directory == (char *) NULL)
279  {
280  if (cache_directory_lock == (SemaphoreInfo *) NULL)
281  ActivateSemaphoreInfo(&cache_directory_lock);
282  LockSemaphoreInfo(cache_directory_lock);
283  if (cache_directory == (char *) NULL)
284  {
285  char
286  *home,
287  path[MagickPathExtent],
288  *temp;
289 
291  status;
292 
293  struct stat
294  attributes;
295 
296  temp=(char *) NULL;
297  home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
298  if (home == (char *) NULL)
299  {
300  home=GetEnvironmentValue("XDG_CACHE_HOME");
301  if (home == (char *) NULL)
302  home=GetEnvironmentValue("LOCALAPPDATA");
303  if (home == (char *) NULL)
304  home=GetEnvironmentValue("APPDATA");
305  if (home == (char *) NULL)
306  home=GetEnvironmentValue("USERPROFILE");
307  }
308 
309  if (home != (char *) NULL)
310  {
311  /* first check if $HOME exists */
312  (void) FormatLocaleString(path,MagickPathExtent,"%s",home);
313  status=GetPathAttributes(path,&attributes);
314  if (status == MagickFalse)
315  status=MagickCreateDirectory(path);
316 
317  /* first check if $HOME/ImageMagick exists */
318  if (status != MagickFalse)
319  {
321  "%s%sImageMagick",home,DirectorySeparator);
322 
323  status=GetPathAttributes(path,&attributes);
324  if (status == MagickFalse)
325  status=MagickCreateDirectory(path);
326  }
327 
328  if (status != MagickFalse)
329  {
330  temp=(char*) AcquireCriticalMemory(strlen(path)+1);
331  CopyMagickString(temp,path,strlen(path)+1);
332  }
333  home=DestroyString(home);
334  }
335  else
336  {
337  home=GetEnvironmentValue("HOME");
338  if (home != (char *) NULL)
339  {
340  /* first check if $HOME/.cache exists */
341  (void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache",
342  home,DirectorySeparator);
343  status=GetPathAttributes(path,&attributes);
344  if (status == MagickFalse)
345  status=MagickCreateDirectory(path);
346 
347  /* first check if $HOME/.cache/ImageMagick exists */
348  if (status != MagickFalse)
349  {
351  "%s%s.cache%sImageMagick",home,DirectorySeparator,
353  status=GetPathAttributes(path,&attributes);
354  if (status == MagickFalse)
355  status=MagickCreateDirectory(path);
356  }
357 
358  if (status != MagickFalse)
359  {
360  temp=(char*) AcquireCriticalMemory(strlen(path)+1);
361  CopyMagickString(temp,path,strlen(path)+1);
362  }
363  home=DestroyString(home);
364  }
365  }
366  if (temp == (char *) NULL)
367  temp=AcquireString("?");
368  cache_directory=temp;
369  }
370  UnlockSemaphoreInfo(cache_directory_lock);
371  }
372  if (*cache_directory == '?')
373  return((const char *) NULL);
374  return(cache_directory);
375 }
376 
377 static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
378 {
380  device;
381 
382  size_t
383  i,
384  j;
385 
386  for (i = 0; i < clEnv->number_devices; i++)
387  clEnv->devices[i]->enabled=MagickFalse;
388 
389  for (i = 0; i < clEnv->number_devices; i++)
390  {
391  device=clEnv->devices[i];
392  if (device->type != type)
393  continue;
394 
395  device->enabled=MagickTrue;
396  for (j = i+1; j < clEnv->number_devices; j++)
397  {
399  other_device;
400 
401  other_device=clEnv->devices[j];
402  if (IsSameOpenCLDevice(device,other_device))
403  other_device->enabled=MagickTrue;
404  }
405  }
406 }
407 
408 static size_t StringSignature(const char* string)
409 {
410  size_t
411  n,
412  i,
413  j,
414  signature,
415  stringLength;
416 
417  union
418  {
419  const char* s;
420  const size_t* u;
421  } p;
422 
423  stringLength=(size_t) strlen(string);
424  signature=stringLength;
425  n=stringLength/sizeof(size_t);
426  p.s=string;
427  for (i = 0; i < n; i++)
428  signature^=p.u[i];
429  if (n * sizeof(size_t) != stringLength)
430  {
431  char
432  padded[4];
433 
434  j=n*sizeof(size_t);
435  for (i = 0; i < 4; i++, j++)
436  {
437  if (j < stringLength)
438  padded[i]=p.s[j];
439  else
440  padded[i]=0;
441  }
442  p.s=padded;
443  signature^=p.u[0];
444  }
445  return(signature);
446 }
447 
448 static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
449 {
450  ssize_t
451  i;
452 
453  for (i=0; i < (ssize_t) info->event_count; i++)
454  openCL_library->clReleaseEvent(info->events[i]);
455  info->events=(cl_event *) RelinquishMagickMemory(info->events);
456  if (info->buffer != (cl_mem) NULL)
457  openCL_library->clReleaseMemObject(info->buffer);
458  RelinquishSemaphoreInfo(&info->events_semaphore);
459  ReleaseOpenCLDevice(info->device);
461 }
462 
463 /*
464  Provide call to OpenCL library methods
465 */
466 
467 MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
468  cl_mem_flags flags,size_t size,void *host_ptr)
469 {
470  return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr,
471  (cl_int *) NULL));
472 }
473 
474 MagickPrivate void ReleaseOpenCLKernel(cl_kernel kernel)
475 {
476  (void) openCL_library->clReleaseKernel(kernel);
477 }
478 
479 MagickPrivate void ReleaseOpenCLMemObject(cl_mem memobj)
480 {
481  (void) openCL_library->clReleaseMemObject(memobj);
482 }
483 
484 MagickPrivate void RetainOpenCLMemObject(cl_mem memobj)
485 {
486  (void) openCL_library->clRetainMemObject(memobj);
487 }
488 
489 MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,size_t arg_index,
490  size_t arg_size,const void *arg_value)
491 {
492  return(openCL_library->clSetKernelArg(kernel,(cl_uint) arg_index,arg_size,
493  arg_value));
494 }
495 
496 /*
497 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
498 % %
499 % %
500 % %
501 + A c q u i r e M a g i c k C L C a c h e I n f o %
502 % %
503 % %
504 % %
505 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
506 %
507 % AcquireMagickCLCacheInfo() acquires an OpenCL cache info structure.
508 %
509 % The format of the AcquireMagickCLCacheInfo method is:
510 %
511 % MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
512 % Quantum *pixels,const MagickSizeType length)
513 %
514 % A description of each parameter follows:
515 %
516 % o device: the OpenCL device.
517 %
518 % o pixels: the pixel buffer of the image.
519 %
520 % o length: the length of the pixel buffer.
521 %
522 */
523 
524 MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
525  Quantum *pixels,const MagickSizeType length)
526 {
527  cl_int
528  status;
529 
531  info;
532 
533  info=(MagickCLCacheInfo) AcquireCriticalMemory(sizeof(*info));
534  (void) ResetMagickMemory(info,0,sizeof(*info));
535  LockSemaphoreInfo(openCL_lock);
536  device->requested++;
537  UnlockSemaphoreInfo(openCL_lock);
538  info->device=device;
539  info->length=length;
540  info->pixels=pixels;
541  info->events_semaphore=AcquireSemaphoreInfo();
542  info->buffer=openCL_library->clCreateBuffer(device->context,
543  CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,
544  &status);
545  if (status == CL_SUCCESS)
546  return(info);
547  DestroyMagickCLCacheInfo(info);
548  return((MagickCLCacheInfo) NULL);
549 }
550 
551 /*
552 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
553 % %
554 % %
555 % %
556 % A c q u i r e M a g i c k C L D e v i c e %
557 % %
558 % %
559 % %
560 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
561 %
562 % AcquireMagickCLDevice() acquires an OpenCL device
563 %
564 % The format of the AcquireMagickCLDevice method is:
565 %
566 % MagickCLDevice AcquireMagickCLDevice()
567 %
568 */
569 
570 static MagickCLDevice AcquireMagickCLDevice()
571 {
573  device;
574 
575  device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device));
576  if (device != NULL)
577  {
578  (void) ResetMagickMemory(device,0,sizeof(*device));
579  ActivateSemaphoreInfo(&device->lock);
580  device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
581  device->command_queues_index=-1;
582  device->enabled=MagickTrue;
583  }
584  return(device);
585 }
586 
587 /*
588 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
589 % %
590 % %
591 % %
592 % A c q u i r e M a g i c k C L E n v %
593 % %
594 % %
595 % %
596 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
597 %
598 % AcquireMagickCLEnv() allocates the MagickCLEnv structure
599 %
600 */
601 
602 static MagickCLEnv AcquireMagickCLEnv(void)
603 {
604  const char
605  *option;
606 
607  MagickCLEnv
608  clEnv;
609 
610  clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv));
611  if (clEnv != (MagickCLEnv) NULL)
612  {
613  (void) ResetMagickMemory(clEnv,0,sizeof(*clEnv));
614  ActivateSemaphoreInfo(&clEnv->lock);
615  clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
616  clEnv->enabled=MagickTrue;
617  option=getenv("MAGICK_OCL_DEVICE");
618  if ((option != (const char *) NULL) && (strcmp(option,"OFF") == 0))
619  clEnv->enabled=MagickFalse;
620  }
621  return clEnv;
622 }
623 
624 /*
625 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
626 % %
627 % %
628 % %
629 + A c q u i r e O p e n C L C o m m a n d Q u e u e %
630 % %
631 % %
632 % %
633 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
634 %
635 % AcquireOpenCLCommandQueue() acquires an OpenCL command queue
636 %
637 % The format of the AcquireOpenCLCommandQueue method is:
638 %
639 % cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
640 %
641 % A description of each parameter follows:
642 %
643 % o device: the OpenCL device.
644 %
645 */
646 
647 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
648 {
649  cl_command_queue
650  queue;
651 
652  cl_command_queue_properties
653  properties;
654 
655  assert(device != (MagickCLDevice) NULL);
656  LockSemaphoreInfo(device->lock);
657  if ((device->profile_kernels == MagickFalse) &&
658  (device->command_queues_index >= 0))
659  {
660  queue=device->command_queues[device->command_queues_index--];
661  UnlockSemaphoreInfo(device->lock);
662  }
663  else
664  {
665  UnlockSemaphoreInfo(device->lock);
666  properties=(cl_command_queue_properties) NULL;
667  if (device->profile_kernels != MagickFalse)
668  properties=CL_QUEUE_PROFILING_ENABLE;
669  queue=openCL_library->clCreateCommandQueue(device->context,
670  device->deviceID,properties,(cl_int *) NULL);
671  }
672  return(queue);
673 }
674 
675 /*
676 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
677 % %
678 % %
679 % %
680 + A c q u i r e O p e n C L K e r n e l %
681 % %
682 % %
683 % %
684 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
685 %
686 % AcquireOpenCLKernel() acquires an OpenCL kernel
687 %
688 % The format of the AcquireOpenCLKernel method is:
689 %
690 % cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
691 % MagickOpenCLProgram program, const char* kernelName)
692 %
693 % A description of each parameter follows:
694 %
695 % o clEnv: the OpenCL environment.
696 %
697 % o program: the OpenCL program module that the kernel belongs to.
698 %
699 % o kernelName: the name of the kernel
700 %
701 */
702 
703 MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
704  const char *kernel_name)
705 {
706  cl_kernel
707  kernel;
708 
709  assert(device != (MagickCLDevice) NULL);
710  kernel=openCL_library->clCreateKernel(device->program,kernel_name,
711  (cl_int *) NULL);
712  return(kernel);
713 }
714 
715 /*
716 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
717 % %
718 % %
719 % %
720 % A u t o S e l e c t O p e n C L D e v i c e s %
721 % %
722 % %
723 % %
724 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
725 %
726 % AutoSelectOpenCLDevices() determines the best device based on the
727 % information from the micro-benchmark.
728 %
729 % The format of the AutoSelectOpenCLDevices method is:
730 %
731 % void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception)
732 %
733 % A description of each parameter follows:
734 %
735 % o clEnv: the OpenCL environment.
736 %
737 % o exception: return any errors or warnings in this structure.
738 %
739 */
740 
741 static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
742 {
743  char
744  keyword[MagickPathExtent],
745  *token;
746 
747  const char
748  *q;
749 
750  MagickCLDeviceBenchmark
751  *device_benchmark;
752 
753  size_t
754  i,
755  extent;
756 
757  if (xml == (char *) NULL)
758  return;
759  device_benchmark=(MagickCLDeviceBenchmark *) NULL;
760  token=AcquireString(xml);
761  extent=strlen(token)+MagickPathExtent;
762  for (q=(char *) xml; *q != '\0'; )
763  {
764  /*
765  Interpret XML.
766  */
767  GetNextToken(q,&q,extent,token);
768  if (*token == '\0')
769  break;
770  (void) CopyMagickString(keyword,token,MagickPathExtent);
771  if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
772  {
773  /*
774  Doctype element.
775  */
776  while ((LocaleNCompare(q,"]>",2) != 0) && (*q != '\0'))
777  GetNextToken(q,&q,extent,token);
778  continue;
779  }
780  if (LocaleNCompare(keyword,"<!--",4) == 0)
781  {
782  /*
783  Comment element.
784  */
785  while ((LocaleNCompare(q,"->",2) != 0) && (*q != '\0'))
786  GetNextToken(q,&q,extent,token);
787  continue;
788  }
789  if (LocaleCompare(keyword,"<device") == 0)
790  {
791  /*
792  Device element.
793  */
794  device_benchmark=(MagickCLDeviceBenchmark *) AcquireMagickMemory(
795  sizeof(*device_benchmark));
796  if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
797  break;
798  (void) ResetMagickMemory(device_benchmark,0,sizeof(*device_benchmark));
799  device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
800  continue;
801  }
802  if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
803  continue;
804  if (LocaleCompare(keyword,"/>") == 0)
805  {
806  if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
807  {
808  if (LocaleCompare(device_benchmark->name, "CPU") == 0)
809  clEnv->cpu_score=device_benchmark->score;
810  else
811  {
813  device;
814 
815  /*
816  Set the score for all devices that match this device.
817  */
818  for (i = 0; i < clEnv->number_devices; i++)
819  {
820  device=clEnv->devices[i];
821  if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
822  device->score=device_benchmark->score;
823  }
824  }
825  }
826 
827  device_benchmark->platform_name=RelinquishMagickMemory(
828  device_benchmark->platform_name);
829  device_benchmark->vendor_name=RelinquishMagickMemory(
830  device_benchmark->vendor_name);
831  device_benchmark->name=RelinquishMagickMemory(device_benchmark->name);
832  device_benchmark->version=RelinquishMagickMemory(
833  device_benchmark->version);
834  device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
835  device_benchmark);
836  continue;
837  }
838  GetNextToken(q,(const char **) NULL,extent,token);
839  if (*token != '=')
840  continue;
841  GetNextToken(q,&q,extent,token);
842  GetNextToken(q,&q,extent,token);
843  switch (*keyword)
844  {
845  case 'M':
846  case 'm':
847  {
848  if (LocaleCompare((char *) keyword,"maxClockFrequency") == 0)
849  {
850  device_benchmark->max_clock_frequency=StringToInteger(token);
851  break;
852  }
853  if (LocaleCompare((char *) keyword,"maxComputeUnits") == 0)
854  {
855  device_benchmark->max_compute_units=StringToInteger(token);
856  break;
857  }
858  break;
859  }
860  case 'N':
861  case 'n':
862  {
863  if (LocaleCompare((char *) keyword,"name") == 0)
864  device_benchmark->name=ConstantString(token);
865  break;
866  }
867  case 'P':
868  case 'p':
869  {
870  if (LocaleCompare((char *) keyword,"platform") == 0)
871  device_benchmark->platform_name=ConstantString(token);
872  break;
873  }
874  case 'S':
875  case 's':
876  {
877  if (LocaleCompare((char *) keyword,"score") == 0)
878  device_benchmark->score=StringToDouble(token,(char **) NULL);
879  break;
880  }
881  case 'V':
882  case 'v':
883  {
884  if (LocaleCompare((char *) keyword,"vendor") == 0)
885  device_benchmark->vendor_name=ConstantString(token);
886  if (LocaleCompare((char *) keyword,"version") == 0)
887  device_benchmark->version=ConstantString(token);
888  break;
889  }
890  default:
891  break;
892  }
893  }
894  token=(char *) RelinquishMagickMemory(token);
895  device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
896  device_benchmark);
897 }
898 
899 static MagickBooleanType CanWriteProfileToFile(const char *filename)
900 {
901  FILE
902  *profileFile;
903 
904  profileFile=fopen(filename,"ab");
905 
906  if (profileFile == (FILE *)NULL)
907  return(MagickFalse);
908 
909  fclose(profileFile);
910  return(MagickTrue);
911 }
912 
913 static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
914 {
915  char
916  filename[MagickPathExtent];
917 
918  StringInfo
919  *option;
920 
921  size_t
922  i;
923 
924  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
925  GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
926 
927  /*
928  We don't run the benchmark when we can not write out a device profile. The
929  first GPU device will be used.
930  */
931 #if !defined(MAGICKCORE_ZERO_CONFIGURATION_SUPPORT)
932  if (CanWriteProfileToFile(filename) == MagickFalse)
933 #endif
934  {
935  for (i = 0; i < clEnv->number_devices; i++)
936  clEnv->devices[i]->score=1.0;
937 
938  SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
939  return(MagickFalse);
940  }
941 
942  option=ConfigureFileToStringInfo(filename);
943  LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(option));
944  option=DestroyStringInfo(option);
945  return(MagickTrue);
946 }
947 
948 static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
949 {
950  const char
951  *option;
952 
953  double
954  best_score;
955 
957  benchmark;
958 
959  size_t
960  i;
961 
962  option=getenv("MAGICK_OCL_DEVICE");
963  if (option != (const char *) NULL)
964  {
965  if (strcmp(option,"GPU") == 0)
966  SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
967  else if (strcmp(option,"CPU") == 0)
968  SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
969  else if (strcmp(option,"OFF") == 0)
970  {
971  for (i = 0; i < clEnv->number_devices; i++)
972  clEnv->devices[i]->enabled=MagickFalse;
973  clEnv->enabled=MagickFalse;
974  }
975  }
976 
977  if (LoadOpenCLBenchmarks(clEnv) == MagickFalse)
978  return;
979 
980  benchmark=MagickFalse;
981  if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
982  benchmark=MagickTrue;
983  else
984  {
985  for (i = 0; i < clEnv->number_devices; i++)
986  {
987  if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
988  {
989  benchmark=MagickTrue;
990  break;
991  }
992  }
993  }
994 
995  if (benchmark != MagickFalse)
996  BenchmarkOpenCLDevices(clEnv);
997 
998  best_score=clEnv->cpu_score;
999  for (i = 0; i < clEnv->number_devices; i++)
1000  best_score=MagickMin(clEnv->devices[i]->score,best_score);
1001 
1002  for (i = 0; i < clEnv->number_devices; i++)
1003  {
1004  if (clEnv->devices[i]->score != best_score)
1005  clEnv->devices[i]->enabled=MagickFalse;
1006  }
1007 }
1008 
1009 /*
1010 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1011 % %
1012 % %
1013 % %
1014 % B e n c h m a r k O p e n C L D e v i c e s %
1015 % %
1016 % %
1017 % %
1018 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1019 %
1020 % BenchmarkOpenCLDevices() benchmarks the OpenCL devices and the CPU to help
1021 % the automatic selection of the best device.
1022 %
1023 % The format of the BenchmarkOpenCLDevices method is:
1024 %
1025 % void BenchmarkOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
1026 %
1027 % A description of each parameter follows:
1028 %
1029 % o clEnv: the OpenCL environment.
1030 %
1031 % o exception: return any errors or warnings
1032 */
1033 
1034 static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
1035 {
1036  AccelerateTimer
1037  timer;
1038 
1040  *exception;
1041 
1042  Image
1043  *inputImage;
1044 
1045  ImageInfo
1046  *imageInfo;
1047 
1048  size_t
1049  i;
1050 
1051  exception=AcquireExceptionInfo();
1052  imageInfo=AcquireImageInfo();
1053  CloneString(&imageInfo->size,"2048x1536");
1054  CopyMagickString(imageInfo->filename,"xc:none",MagickPathExtent);
1055  inputImage=ReadImage(imageInfo,exception);
1056 
1057  InitAccelerateTimer(&timer);
1058 
1059  for (i=0; i<=2; i++)
1060  {
1061  Image
1062  *bluredImage,
1063  *resizedImage,
1064  *unsharpedImage;
1065 
1066  if (i > 0)
1067  StartAccelerateTimer(&timer);
1068 
1069  bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
1070  unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
1071  exception);
1072  resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
1073  exception);
1074 
1075  /*
1076  We need this to get a proper performance benchmark, the operations
1077  are executed asynchronous.
1078  */
1079  if (is_cpu == MagickFalse)
1080  {
1081  CacheInfo
1082  *cache_info;
1083 
1084  cache_info=(CacheInfo *) resizedImage->cache;
1085  if (cache_info->opencl != (MagickCLCacheInfo) NULL)
1086  openCL_library->clWaitForEvents(cache_info->opencl->event_count,
1087  cache_info->opencl->events);
1088  }
1089 
1090  if (i > 0)
1091  StopAccelerateTimer(&timer);
1092 
1093  if (bluredImage != (Image *) NULL)
1094  DestroyImage(bluredImage);
1095  if (unsharpedImage != (Image *) NULL)
1096  DestroyImage(unsharpedImage);
1097  if (resizedImage != (Image *) NULL)
1098  DestroyImage(resizedImage);
1099  }
1100  DestroyImage(inputImage);
1101  return(ReadAccelerateTimer(&timer));
1102 }
1103 
1104 static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
1105  MagickCLDevice device)
1106 {
1107  testEnv->devices[0]=device;
1108  default_CLEnv=testEnv;
1109  device->score=RunOpenCLBenchmark(MagickFalse);
1110  default_CLEnv=clEnv;
1111  testEnv->devices[0]=(MagickCLDevice) NULL;
1112 }
1113 
1114 static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
1115 {
1116  char
1117  filename[MagickPathExtent];
1118 
1119  FILE
1120  *cache_file;
1121 
1123  device;
1124 
1125  size_t
1126  i,
1127  j;
1128 
1129  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1130  GetOpenCLCacheDirectory(),DirectorySeparator,
1131  IMAGEMAGICK_PROFILE_FILE);
1132 
1133  cache_file=fopen_utf8(filename,"wb");
1134  if (cache_file == (FILE *) NULL)
1135  return;
1136  fwrite("<devices>\n",sizeof(char),10,cache_file);
1137  fprintf(cache_file," <device name=\"CPU\" score=\"%.4g\"/>\n",
1138  clEnv->cpu_score);
1139  for (i = 0; i < clEnv->number_devices; i++)
1140  {
1142  duplicate;
1143 
1144  device=clEnv->devices[i];
1145  duplicate=MagickFalse;
1146  for (j = 0; j < i; j++)
1147  {
1148  if (IsSameOpenCLDevice(clEnv->devices[j],device))
1149  {
1150  duplicate=MagickTrue;
1151  break;
1152  }
1153  }
1154 
1155  if (duplicate)
1156  continue;
1157 
1158  if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1159  fprintf(cache_file," <device platform=\"%s\" vendor=\"%s\" name=\"%s\"\
1160  version=\"%s\" maxClockFrequency=\"%d\" maxComputeUnits=\"%d\"\
1161  score=\"%.4g\"/>\n",
1162  device->platform_name,device->vendor_name,device->name,device->version,
1163  (int)device->max_clock_frequency,(int)device->max_compute_units,
1164  device->score);
1165  }
1166  fwrite("</devices>",sizeof(char),10,cache_file);
1167 
1168  fclose(cache_file);
1169 }
1170 
1171 static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
1172 {
1174  device;
1175 
1176  MagickCLEnv
1177  testEnv;
1178 
1179  size_t
1180  i,
1181  j;
1182 
1183  testEnv=AcquireMagickCLEnv();
1184  testEnv->library=openCL_library;
1185  testEnv->devices=(MagickCLDevice *) AcquireCriticalMemory(
1186  sizeof(MagickCLDevice));
1187  testEnv->number_devices=1;
1188  testEnv->benchmark_thread_id=GetMagickThreadId();
1189  testEnv->initialized=MagickTrue;
1190 
1191  for (i = 0; i < clEnv->number_devices; i++)
1192  clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
1193 
1194  for (i = 0; i < clEnv->number_devices; i++)
1195  {
1196  device=clEnv->devices[i];
1197  if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1198  RunDeviceBenckmark(clEnv,testEnv,device);
1199 
1200  /* Set the score on all the other devices that are the same */
1201  for (j = i+1; j < clEnv->number_devices; j++)
1202  {
1204  other_device;
1205 
1206  other_device=clEnv->devices[j];
1207  if (IsSameOpenCLDevice(device,other_device))
1208  other_device->score=device->score;
1209  }
1210  }
1211 
1212  testEnv->enabled=MagickFalse;
1213  default_CLEnv=testEnv;
1214  clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
1215  default_CLEnv=clEnv;
1216 
1217  testEnv=RelinquishMagickCLEnv(testEnv);
1218  CacheOpenCLBenchmarks(clEnv);
1219 }
1220 
1221 /*
1222 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1223 % %
1224 % %
1225 % %
1226 % C o m p i l e O p e n C L K e r n e l %
1227 % %
1228 % %
1229 % %
1230 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1231 %
1232 % CompileOpenCLKernel() compiles the kernel for the specified device. The
1233 % kernel will be cached on disk to reduce the compilation time.
1234 %
1235 % The format of the CompileOpenCLKernel method is:
1236 %
1237 % MagickBooleanType AcquireOpenCLKernel(MagickCLDevice clEnv,
1238 % unsigned int signature,const char *kernel,const char *options,
1239 % ExceptionInfo *exception)
1240 %
1241 % A description of each parameter follows:
1242 %
1243 % o device: the OpenCL device.
1244 %
1245 % o kernel: the source code of the kernel.
1246 %
1247 % o options: options for the compiler.
1248 %
1249 % o signature: a number to uniquely identify the kernel
1250 %
1251 % o exception: return any errors or warnings in this structure.
1252 %
1253 */
1254 
1255 static void CacheOpenCLKernel(MagickCLDevice device,char *filename,
1256  ExceptionInfo *exception)
1257 {
1258  cl_uint
1259  status;
1260 
1261  size_t
1262  binaryProgramSize;
1263 
1264  unsigned char
1265  *binaryProgram;
1266 
1267  status=openCL_library->clGetProgramInfo(device->program,
1268  CL_PROGRAM_BINARY_SIZES,sizeof(size_t),&binaryProgramSize,NULL);
1269  if (status != CL_SUCCESS)
1270  return;
1271  binaryProgram=(unsigned char*) AcquireMagickMemory(binaryProgramSize);
1272  if (binaryProgram == (unsigned char *) NULL)
1273  {
1274  (void) ThrowMagickException(exception,GetMagickModule(),
1275  ResourceLimitError,"MemoryAllocationFailed","`%s'",filename);
1276  return;
1277  }
1278  status=openCL_library->clGetProgramInfo(device->program,
1279  CL_PROGRAM_BINARIES,sizeof(unsigned char*),&binaryProgram,NULL);
1280  if (status == CL_SUCCESS)
1281  (void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
1282  binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1283 }
1284 
1285 static MagickBooleanType LoadCachedOpenCLKernel(MagickCLDevice device,
1286  const char *filename)
1287 {
1288  cl_int
1289  binaryStatus,
1290  status;
1291 
1293  *exception;
1294 
1295  size_t
1296  length;
1297 
1298  unsigned char
1299  *binaryProgram;
1300 
1301  exception=AcquireExceptionInfo();
1302  binaryProgram=(unsigned char *) FileToBlob(filename,~0UL,&length,exception);
1303  exception=DestroyExceptionInfo(exception);
1304  if (binaryProgram == (unsigned char *) NULL)
1305  return(MagickFalse);
1306  device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
1307  &device->deviceID,&length,(const unsigned char**)&binaryProgram,
1308  &binaryStatus,&status);
1309  binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1310  return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
1311  MagickTrue);
1312 }
1313 
1314 static void LogOpenCLBuildFailure(MagickCLDevice device,const char *kernel,
1315  ExceptionInfo *exception)
1316 {
1317  char
1318  filename[MagickPathExtent],
1319  *log;
1320 
1321  size_t
1322  log_size;
1323 
1324  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1325  GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.cl");
1326 
1327  (void) remove_utf8(filename);
1328  (void) BlobToFile(filename,kernel,strlen(kernel),exception);
1329 
1330  openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1331  CL_PROGRAM_BUILD_LOG,0,NULL,&log_size);
1332  log=(char*)AcquireCriticalMemory(log_size);
1333  openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1334  CL_PROGRAM_BUILD_LOG,log_size,log,&log_size);
1335 
1336  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1337  GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.log");
1338 
1339  (void) remove_utf8(filename);
1340  (void) BlobToFile(filename,log,log_size,exception);
1341  log=(char*)RelinquishMagickMemory(log);
1342 }
1343 
1344 static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
1345  const char *kernel,const char *options,size_t signature,
1346  ExceptionInfo *exception)
1347 {
1348  char
1349  deviceName[MagickPathExtent],
1350  filename[MagickPathExtent],
1351  *ptr;
1352 
1353  cl_int
1354  status;
1355 
1357  loaded;
1358 
1359  size_t
1360  length;
1361 
1362  (void) CopyMagickString(deviceName,device->name,MagickPathExtent);
1363  ptr=deviceName;
1364  /* Strip out illegal characters for file names */
1365  while (*ptr != '\0')
1366  {
1367  if ((*ptr == ' ') || (*ptr == '\\') || (*ptr == '/') || (*ptr == ':') ||
1368  (*ptr == '*') || (*ptr == '?') || (*ptr == '"') || (*ptr == '<') ||
1369  (*ptr == '>' || *ptr == '|'))
1370  *ptr = '_';
1371  ptr++;
1372  }
1373  (void) FormatLocaleString(filename,MagickPathExtent,
1374  "%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
1375  DirectorySeparator,"magick_opencl",deviceName,(unsigned int) signature,
1376  (double) sizeof(char*)*8);
1377  loaded=LoadCachedOpenCLKernel(device,filename);
1378  if (loaded == MagickFalse)
1379  {
1380  /* Binary CL program unavailable, compile the program from source */
1381  length=strlen(kernel);
1382  device->program=openCL_library->clCreateProgramWithSource(
1383  device->context,1,&kernel,&length,&status);
1384  if (status != CL_SUCCESS)
1385  return(MagickFalse);
1386  }
1387 
1388  status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
1389  options,NULL,NULL);
1390  if (status != CL_SUCCESS)
1391  {
1393  "clBuildProgram failed.","(%d)",(int)status);
1394  LogOpenCLBuildFailure(device,kernel,exception);
1395  return(MagickFalse);
1396  }
1397 
1398  /* Save the binary to a file to avoid re-compilation of the kernels */
1399  if (loaded == MagickFalse)
1400  CacheOpenCLKernel(device,filename,exception);
1401 
1402  return(MagickTrue);
1403 }
1404 
1405 static cl_event* CopyOpenCLEvents(MagickCLCacheInfo first,
1406  MagickCLCacheInfo second,cl_uint *event_count)
1407 {
1408  cl_event
1409  *events;
1410 
1411  register size_t
1412  i;
1413 
1414  size_t
1415  j;
1416 
1417  assert(first != (MagickCLCacheInfo) NULL);
1418  assert(event_count != (cl_uint *) NULL);
1419  events=(cl_event *) NULL;
1420  LockSemaphoreInfo(first->events_semaphore);
1421  if (second != (MagickCLCacheInfo) NULL)
1422  LockSemaphoreInfo(second->events_semaphore);
1423  *event_count=first->event_count;
1424  if (second != (MagickCLCacheInfo) NULL)
1425  *event_count+=second->event_count;
1426  if (*event_count > 0)
1427  {
1428  events=AcquireQuantumMemory(*event_count,sizeof(*events));
1429  if (events == (cl_event *) NULL)
1430  *event_count=0;
1431  else
1432  {
1433  j=0;
1434  for (i=0; i < first->event_count; i++, j++)
1435  events[j]=first->events[i];
1436  if (second != (MagickCLCacheInfo) NULL)
1437  {
1438  for (i=0; i < second->event_count; i++, j++)
1439  events[j]=second->events[i];
1440  }
1441  }
1442  }
1443  UnlockSemaphoreInfo(first->events_semaphore);
1444  if (second != (MagickCLCacheInfo) NULL)
1445  UnlockSemaphoreInfo(second->events_semaphore);
1446  return(events);
1447 }
1448 
1449 /*
1450 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1451 % %
1452 % %
1453 % %
1454 + C o p y M a g i c k C L C a c h e I n f o %
1455 % %
1456 % %
1457 % %
1458 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1459 %
1460 % CopyMagickCLCacheInfo() copies the memory from the device into host memory.
1461 %
1462 % The format of the CopyMagickCLCacheInfo method is:
1463 %
1464 % void CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1465 %
1466 % A description of each parameter follows:
1467 %
1468 % o info: the OpenCL cache info.
1469 %
1470 */
1471 MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1472 {
1473  cl_command_queue
1474  queue;
1475 
1476  cl_event
1477  *events;
1478 
1479  cl_uint
1480  event_count;
1481 
1482  Quantum
1483  *pixels;
1484 
1485  if (info == (MagickCLCacheInfo) NULL)
1486  return((MagickCLCacheInfo) NULL);
1487  events=CopyOpenCLEvents(info,(MagickCLCacheInfo) NULL,&event_count);
1488  if (events != (cl_event *) NULL)
1489  {
1490  queue=AcquireOpenCLCommandQueue(info->device);
1491  pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
1492  CL_MAP_READ | CL_MAP_WRITE,0,info->length,event_count,events,
1493  (cl_event *) NULL,(cl_int *) NULL);
1494  assert(pixels == info->pixels);
1495  ReleaseOpenCLCommandQueue(info->device,queue);
1496  events=(cl_event *) RelinquishMagickMemory(events);
1497  }
1498  return(RelinquishMagickCLCacheInfo(info,MagickFalse));
1499 }
1500 
1501 /*
1502 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1503 % %
1504 % %
1505 % %
1506 + D u m p O p e n C L P r o f i l e D a t a %
1507 % %
1508 % %
1509 % %
1510 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1511 %
1512 % DumpOpenCLProfileData() dumps the kernel profile data.
1513 %
1514 % The format of the DumpProfileData method is:
1515 %
1516 % void DumpProfileData()
1517 %
1518 */
1519 
1520 MagickPrivate void DumpOpenCLProfileData()
1521 {
1522 #define OpenCLLog(message) \
1523  fwrite(message,sizeof(char),strlen(message),log); \
1524  fwrite("\n",sizeof(char),1,log);
1525 
1526  char
1527  buf[4096],
1528  filename[MagickPathExtent],
1529  indent[160];
1530 
1531  FILE
1532  *log;
1533 
1534  MagickCLEnv
1535  clEnv;
1536 
1537  size_t
1538  i,
1539  j;
1540 
1541  clEnv=GetCurrentOpenCLEnv();
1542  if (clEnv == (MagickCLEnv) NULL)
1543  return;
1544 
1545  for (i = 0; i < clEnv->number_devices; i++)
1546  if (clEnv->devices[i]->profile_kernels != MagickFalse)
1547  break;
1548  if (i == clEnv->number_devices)
1549  return;
1550 
1551  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1552  GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
1553 
1554  log=fopen_utf8(filename,"wb");
1555 
1556  for (i = 0; i < clEnv->number_devices; i++)
1557  {
1559  device;
1560 
1561  device=clEnv->devices[i];
1562  if ((device->profile_kernels == MagickFalse) ||
1563  (device->profile_records == (KernelProfileRecord *) NULL))
1564  continue;
1565 
1566  OpenCLLog("====================================================");
1567  fprintf(log,"Device: %s\n",device->name);
1568  fprintf(log,"Version: %s\n",device->version);
1569  OpenCLLog("====================================================");
1570  OpenCLLog(" average calls min max");
1571  OpenCLLog(" ------- ----- --- ---");
1572  j=0;
1573  while (device->profile_records[j] != (KernelProfileRecord) NULL)
1574  {
1576  profile;
1577 
1578  profile=device->profile_records[j];
1579  strcpy(indent," ");
1580  strncpy(indent,profile->kernel_name,MagickMin(strlen(
1581  profile->kernel_name),strlen(indent)-1));
1582  sprintf(buf,"%s %7d %7d %7d %7d",indent,(int) (profile->total/
1583  profile->count),(int) profile->count,(int) profile->min,
1584  (int) profile->max);
1585  OpenCLLog(buf);
1586  j++;
1587  }
1588  OpenCLLog("====================================================");
1589  fwrite("\n\n",sizeof(char),2,log);
1590  }
1591  fclose(log);
1592 }
1593 /*
1594 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1595 % %
1596 % %
1597 % %
1598 + E n q u e u e O p e n C L K e r n e l %
1599 % %
1600 % %
1601 % %
1602 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1603 %
1604 % EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
1605 % events with the images.
1606 %
1607 % The format of the EnqueueOpenCLKernel method is:
1608 %
1609 % MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
1610 % const size_t *global_work_offset,const size_t *global_work_size,
1611 % const size_t *local_work_size,const Image *input_image,
1612 % const Image *output_image,ExceptionInfo *exception)
1613 %
1614 % A description of each parameter follows:
1615 %
1616 % o kernel: the OpenCL kernel.
1617 %
1618 % o work_dim: the number of dimensions used to specify the global work-items
1619 % and work-items in the work-group.
1620 %
1621 % o offset: can be used to specify an array of work_dim unsigned values
1622 % that describe the offset used to calculate the global ID of a
1623 % work-item.
1624 %
1625 % o gsize: points to an array of work_dim unsigned values that describe the
1626 % number of global work-items in work_dim dimensions that will
1627 % execute the kernel function.
1628 %
1629 % o lsize: points to an array of work_dim unsigned values that describe the
1630 % number of work-items that make up a work-group that will execute
1631 % the kernel specified by kernel.
1632 %
1633 % o input_image: the input image of the operation.
1634 %
1635 % o output_image: the output or secondairy image of the operation.
1636 %
1637 % o exception: return any errors or warnings in this structure.
1638 %
1639 */
1640 
1641 static MagickBooleanType RegisterCacheEvent(MagickCLCacheInfo info,
1642  cl_event event)
1643 {
1644  assert(info != (MagickCLCacheInfo) NULL);
1645  assert(event != (cl_event) NULL);
1646  if (openCL_library->clRetainEvent(event) != CL_SUCCESS)
1647  {
1648  openCL_library->clWaitForEvents(1,&event);
1649  return(MagickFalse);
1650  }
1651  LockSemaphoreInfo(info->events_semaphore);
1652  if (info->events == (cl_event *) NULL)
1653  {
1654  info->events=AcquireMagickMemory(sizeof(*info->events));
1655  info->event_count=1;
1656  }
1657  else
1658  info->events=ResizeQuantumMemory(info->events,++info->event_count,
1659  sizeof(*info->events));
1660  if (info->events == (cl_event *) NULL)
1661  ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
1662  info->events[info->event_count-1]=event;
1663  UnlockSemaphoreInfo(info->events_semaphore);
1664  return(MagickTrue);
1665 }
1666 
1667 MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1668  cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
1669  const size_t *lsize,const Image *input_image,const Image *output_image,
1670  MagickBooleanType flush,ExceptionInfo *exception)
1671 {
1672  CacheInfo
1673  *output_info,
1674  *input_info;
1675 
1676  cl_event
1677  event,
1678  *events;
1679 
1680  cl_int
1681  status;
1682 
1683  cl_uint
1684  event_count;
1685 
1686  assert(input_image != (const Image *) NULL);
1687  input_info=(CacheInfo *) input_image->cache;
1688  assert(input_info != (CacheInfo *) NULL);
1689  assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1690  output_info=(CacheInfo *) NULL;
1691  if (output_image == (const Image *) NULL)
1692  events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
1693  &event_count);
1694  else
1695  {
1696  output_info=(CacheInfo *) output_image->cache;
1697  assert(output_info != (CacheInfo *) NULL);
1698  assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1699  events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
1700  &event_count);
1701  }
1702  status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1703  gsize,lsize,event_count,events,&event);
1704  /* This can fail due to memory issues and calling clFinish might help. */
1705  if ((status != CL_SUCCESS) && (event_count > 0))
1706  {
1707  openCL_library->clFinish(queue);
1708  status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
1709  offset,gsize,lsize,event_count,events,&event);
1710  }
1711  events=(cl_event *) RelinquishMagickMemory(events);
1712  if (status != CL_SUCCESS)
1713  {
1714  (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1716  "clEnqueueNDRangeKernel failed.","'%s'",".");
1717  return(MagickFalse);
1718  }
1719  if (flush != MagickFalse)
1720  openCL_library->clFlush(queue);
1721  if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1722  {
1723  if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
1724  {
1725  if (output_info != (CacheInfo *) NULL)
1726  (void) RegisterCacheEvent(output_info->opencl,event);
1727  }
1728  }
1729  openCL_library->clReleaseEvent(event);
1730  return(MagickTrue);
1731 }
1732 
1733 /*
1734 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1735 % %
1736 % %
1737 % %
1738 + G e t C u r r u n t O p e n C L E n v %
1739 % %
1740 % %
1741 % %
1742 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1743 %
1744 % GetCurrentOpenCLEnv() returns the current OpenCL env
1745 %
1746 % The format of the GetCurrentOpenCLEnv method is:
1747 %
1748 % MagickCLEnv GetCurrentOpenCLEnv()
1749 %
1750 */
1751 
1752 MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
1753 {
1754  if (default_CLEnv != (MagickCLEnv) NULL)
1755  {
1756  if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1757  (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1758  return((MagickCLEnv) NULL);
1759  else
1760  return(default_CLEnv);
1761  }
1762 
1763  if (GetOpenCLCacheDirectory() == (char *) NULL)
1764  return((MagickCLEnv) NULL);
1765 
1766  if (openCL_lock == (SemaphoreInfo *) NULL)
1767  ActivateSemaphoreInfo(&openCL_lock);
1768 
1769  LockSemaphoreInfo(openCL_lock);
1770  if (default_CLEnv == (MagickCLEnv) NULL)
1771  default_CLEnv=AcquireMagickCLEnv();
1772  UnlockSemaphoreInfo(openCL_lock);
1773 
1774  return(default_CLEnv);
1775 }
1776 
1777 /*
1778 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1779 % %
1780 % %
1781 % %
1782 % G e t O p e n C L D e v i c e B e n c h m a r k D u r a t i o n %
1783 % %
1784 % %
1785 % %
1786 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1787 %
1788 % GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
1789 % device. The score is determined by the duration of the micro benchmark so
1790 % that means a lower score is better than a higher score.
1791 %
1792 % The format of the GetOpenCLDeviceBenchmarkScore method is:
1793 %
1794 % double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
1795 %
1796 % A description of each parameter follows:
1797 %
1798 % o device: the OpenCL device.
1799 */
1800 
1802  const MagickCLDevice device)
1803 {
1804  if (device == (MagickCLDevice) NULL)
1805  return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1806  return(device->score);
1807 }
1808 
1809 /*
1810 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1811 % %
1812 % %
1813 % %
1814 % G e t O p e n C L D e v i c e E n a b l e d %
1815 % %
1816 % %
1817 % %
1818 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1819 %
1820 % GetOpenCLDeviceEnabled() returns true if the device is enabled.
1821 %
1822 % The format of the GetOpenCLDeviceEnabled method is:
1823 %
1824 % MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
1825 %
1826 % A description of each parameter follows:
1827 %
1828 % o device: the OpenCL device.
1829 */
1830 
1832  const MagickCLDevice device)
1833 {
1834  if (device == (MagickCLDevice) NULL)
1835  return(MagickFalse);
1836  return(device->enabled);
1837 }
1838 
1839 /*
1840 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1841 % %
1842 % %
1843 % %
1844 % G e t O p e n C L D e v i c e N a m e %
1845 % %
1846 % %
1847 % %
1848 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1849 %
1850 % GetOpenCLDeviceName() returns the name of the device.
1851 %
1852 % The format of the GetOpenCLDeviceName method is:
1853 %
1854 % const char *GetOpenCLDeviceName(const MagickCLDevice device)
1855 %
1856 % A description of each parameter follows:
1857 %
1858 % o device: the OpenCL device.
1859 */
1860 
1861 MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
1862 {
1863  if (device == (MagickCLDevice) NULL)
1864  return((const char *) NULL);
1865  return(device->name);
1866 }
1867 
1868 /*
1869 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1870 % %
1871 % %
1872 % %
1873 % G e t O p e n C L D e v i c e V e n d o r N a m e %
1874 % %
1875 % %
1876 % %
1877 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1878 %
1879 % GetOpenCLDeviceVendorName() returns the vendor name of the device.
1880 %
1881 % The format of the GetOpenCLDeviceVendorName method is:
1882 %
1883 % const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1884 %
1885 % A description of each parameter follows:
1886 %
1887 % o device: the OpenCL device.
1888 */
1889 
1890 MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1891 {
1892  if (device == (MagickCLDevice) NULL)
1893  return((const char *) NULL);
1894  return(device->vendor_name);
1895 }
1896 
1897 /*
1898 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1899 % %
1900 % %
1901 % %
1902 % G e t O p e n C L D e v i c e s %
1903 % %
1904 % %
1905 % %
1906 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1907 %
1908 % GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
1909 % value of length to the number of devices that are available.
1910 %
1911 % The format of the GetOpenCLDevices method is:
1912 %
1913 % const MagickCLDevice *GetOpenCLDevices(size_t *length,
1914 % ExceptionInfo *exception)
1915 %
1916 % A description of each parameter follows:
1917 %
1918 % o length: the number of device.
1919 %
1920 % o exception: return any errors or warnings in this structure.
1921 %
1922 */
1923 
1925  ExceptionInfo *exception)
1926 {
1927  MagickCLEnv
1928  clEnv;
1929 
1930  clEnv=GetCurrentOpenCLEnv();
1931  if (clEnv == (MagickCLEnv) NULL)
1932  {
1933  if (length != (size_t *) NULL)
1934  *length=0;
1935  return((MagickCLDevice *) NULL);
1936  }
1937  InitializeOpenCL(clEnv,exception);
1938  if (length != (size_t *) NULL)
1939  *length=clEnv->number_devices;
1940  return(clEnv->devices);
1941 }
1942 
1943 /*
1944 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1945 % %
1946 % %
1947 % %
1948 % G e t O p e n C L D e v i c e T y p e %
1949 % %
1950 % %
1951 % %
1952 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1953 %
1954 % GetOpenCLDeviceType() returns the type of the device.
1955 %
1956 % The format of the GetOpenCLDeviceType method is:
1957 %
1958 % MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
1959 %
1960 % A description of each parameter follows:
1961 %
1962 % o device: the OpenCL device.
1963 */
1964 
1966  const MagickCLDevice device)
1967 {
1968  if (device == (MagickCLDevice) NULL)
1969  return(UndefinedCLDeviceType);
1970  if (device->type == CL_DEVICE_TYPE_GPU)
1971  return(GpuCLDeviceType);
1972  if (device->type == CL_DEVICE_TYPE_CPU)
1973  return(CpuCLDeviceType);
1974  return(UndefinedCLDeviceType);
1975 }
1976 
1977 /*
1978 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1979 % %
1980 % %
1981 % %
1982 % G e t O p e n C L D e v i c e V e r s i o n %
1983 % %
1984 % %
1985 % %
1986 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1987 %
1988 % GetOpenCLDeviceVersion() returns the version of the device.
1989 %
1990 % The format of the GetOpenCLDeviceName method is:
1991 %
1992 % const char *GetOpenCLDeviceVersion(MagickCLDevice device)
1993 %
1994 % A description of each parameter follows:
1995 %
1996 % o device: the OpenCL device.
1997 */
1998 
1999 MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
2000 {
2001  if (device == (MagickCLDevice) NULL)
2002  return((const char *) NULL);
2003  return(device->version);
2004 }
2005 
2006 /*
2007 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2008 % %
2009 % %
2010 % %
2011 % G e t O p e n C L E n a b l e d %
2012 % %
2013 % %
2014 % %
2015 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2016 %
2017 % GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
2018 %
2019 % The format of the GetOpenCLEnabled method is:
2020 %
2021 % MagickBooleanType GetOpenCLEnabled()
2022 %
2023 */
2024 
2026 {
2027  MagickCLEnv
2028  clEnv;
2029 
2030  clEnv=GetCurrentOpenCLEnv();
2031  if (clEnv == (MagickCLEnv) NULL)
2032  return(MagickFalse);
2033  return(clEnv->enabled);
2034 }
2035 
2036 /*
2037 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2038 % %
2039 % %
2040 % %
2041 % G e t O p e n C L K e r n e l P r o f i l e R e c o r d s %
2042 % %
2043 % %
2044 % %
2045 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2046 %
2047 % GetOpenCLKernelProfileRecords() returns the profile records for the
2048 % specified device and sets length to the number of profile records.
2049 %
2050 % The format of the GetOpenCLKernelProfileRecords method is:
2051 %
2052 % const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
2053 %
2054 % A description of each parameter follows:
2055 %
2056 % o length: the number of profiles records.
2057 */
2058 
2060  const MagickCLDevice device,size_t *length)
2061 {
2062  if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
2063  (KernelProfileRecord *) NULL))
2064  {
2065  if (length != (size_t *) NULL)
2066  *length=0;
2067  return((const KernelProfileRecord *) NULL);
2068  }
2069  if (length != (size_t *) NULL)
2070  {
2071  *length=0;
2072  LockSemaphoreInfo(device->lock);
2073  while (device->profile_records[*length] != (KernelProfileRecord) NULL)
2074  *length=*length+1;
2075  UnlockSemaphoreInfo(device->lock);
2076  }
2077  return(device->profile_records);
2078 }
2079 
2080 /*
2081 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2082 % %
2083 % %
2084 % %
2085 % H a s O p e n C L D e v i c e s %
2086 % %
2087 % %
2088 % %
2089 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2090 %
2091 % HasOpenCLDevices() checks if the OpenCL environment has devices that are
2092 % enabled and compiles the kernel for the device when necessary. False will be
2093 % returned if no enabled devices could be found
2094 %
2095 % The format of the HasOpenCLDevices method is:
2096 %
2097 % MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2098 % ExceptionInfo exception)
2099 %
2100 % A description of each parameter follows:
2101 %
2102 % o clEnv: the OpenCL environment.
2103 %
2104 % o exception: return any errors or warnings in this structure.
2105 %
2106 */
2107 
2108 static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2109  ExceptionInfo *exception)
2110 {
2111  char
2112  *accelerateKernelsBuffer,
2113  options[MagickPathExtent];
2114 
2116  status;
2117 
2118  size_t
2119  i;
2120 
2121  size_t
2122  signature;
2123 
2124  /* Check if there are enabled devices */
2125  for (i = 0; i < clEnv->number_devices; i++)
2126  {
2127  if ((clEnv->devices[i]->enabled != MagickFalse))
2128  break;
2129  }
2130  if (i == clEnv->number_devices)
2131  return(MagickFalse);
2132 
2133  /* Check if we need to compile a kernel for one of the devices */
2134  status=MagickTrue;
2135  for (i = 0; i < clEnv->number_devices; i++)
2136  {
2137  if ((clEnv->devices[i]->enabled != MagickFalse) &&
2138  (clEnv->devices[i]->program == (cl_program) NULL))
2139  {
2140  status=MagickFalse;
2141  break;
2142  }
2143  }
2144  if (status != MagickFalse)
2145  return(MagickTrue);
2146 
2147  /* Get additional options */
2148  (void) FormatLocaleString(options,MaxTextExtent,CLOptions,
2149  (float)QuantumRange,(float)QuantumScale,(float)CLCharQuantumScale,
2150  (float)MagickEpsilon,(float)MagickPI,(unsigned int)MaxMap,
2151  (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2152 
2153  signature=StringSignature(options);
2154  accelerateKernelsBuffer=(char*) AcquireMagickMemory(
2155  strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2156  if (accelerateKernelsBuffer == (char*) NULL)
2157  return(MagickFalse);
2158  sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
2159  signature^=StringSignature(accelerateKernelsBuffer);
2160 
2161  status=MagickTrue;
2162  for (i = 0; i < clEnv->number_devices; i++)
2163  {
2165  device;
2166 
2167  size_t
2168  device_signature;
2169 
2170  device=clEnv->devices[i];
2171  if ((device->enabled == MagickFalse) ||
2172  (device->program != (cl_program) NULL))
2173  continue;
2174 
2175  LockSemaphoreInfo(device->lock);
2176  if (device->program != (cl_program) NULL)
2177  {
2178  UnlockSemaphoreInfo(device->lock);
2179  continue;
2180  }
2181  device_signature=signature;
2182  device_signature^=StringSignature(device->platform_name);
2183  status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2184  device_signature,exception);
2185  UnlockSemaphoreInfo(device->lock);
2186  if (status == MagickFalse)
2187  break;
2188  }
2189  accelerateKernelsBuffer=RelinquishMagickMemory(accelerateKernelsBuffer);
2190  return(status);
2191 }
2192 
2193 /*
2194 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2195 % %
2196 % %
2197 % %
2198 + I n i t i a l i z e O p e n C L %
2199 % %
2200 % %
2201 % %
2202 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2203 %
2204 % InitializeOpenCL() is used to initialize the OpenCL environment. This method
2205 % makes sure the devices are propertly initialized and benchmarked.
2206 %
2207 % The format of the InitializeOpenCL method is:
2208 %
2209 % MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2210 %
2211 % A description of each parameter follows:
2212 %
2213 % o exception: return any errors or warnings in this structure.
2214 %
2215 */
2216 
2217 static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2218 {
2219  char
2220  version[MagickPathExtent];
2221 
2222  cl_uint
2223  num;
2224 
2225  if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2226  MagickPathExtent,version,NULL) != CL_SUCCESS)
2227  return(0);
2228  if (strncmp(version, "OpenCL 1.0 ", 11) == 0)
2229  return(0);
2230  if (clEnv->library->clGetDeviceIDs(platform,
2231  CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2232  return(0);
2233  return(num);
2234 }
2235 
2236 static void LoadOpenCLDevices(MagickCLEnv clEnv)
2237 {
2238  cl_context_properties
2239  properties[3];
2240 
2241  cl_device_id
2242  *devices;
2243 
2244  cl_int
2245  status;
2246 
2247  cl_platform_id
2248  *platforms;
2249 
2250  cl_uint
2251  i,
2252  j,
2253  next,
2254  number_devices,
2255  number_platforms;
2256 
2257  size_t
2258  length;
2259 
2260  number_platforms=0;
2261  if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2262  return;
2263  if (number_platforms == 0)
2264  return;
2265  platforms=(cl_platform_id *) AcquireMagickMemory(number_platforms*
2266  sizeof(cl_platform_id));
2267  if (platforms == (cl_platform_id *) NULL)
2268  return;
2269  if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2270  {
2271  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2272  return;
2273  }
2274  for (i = 0; i < number_platforms; i++)
2275  {
2276  number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2277  if (number_devices == 0)
2278  platforms[i]=(cl_platform_id) NULL;
2279  else
2280  clEnv->number_devices+=number_devices;
2281  }
2282  if (clEnv->number_devices == 0)
2283  {
2284  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2285  return;
2286  }
2287  clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2288  sizeof(MagickCLDevice));
2289  if (clEnv->devices == (MagickCLDevice *) NULL)
2290  {
2291  RelinquishMagickCLDevices(clEnv);
2292  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2293  return;
2294  }
2295  (void) ResetMagickMemory(clEnv->devices,0,clEnv->number_devices*
2296  sizeof(MagickCLDevice));
2297  devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2298  sizeof(cl_device_id));
2299  if (devices == (cl_device_id *) NULL)
2300  {
2301  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2302  RelinquishMagickCLDevices(clEnv);
2303  return;
2304  }
2305  clEnv->number_contexts=(size_t) number_platforms;
2306  clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2307  sizeof(cl_context));
2308  if (clEnv->contexts == (cl_context *) NULL)
2309  {
2310  devices=(cl_device_id *) RelinquishMagickMemory(devices);
2311  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2312  RelinquishMagickCLDevices(clEnv);
2313  return;
2314  }
2315  next=0;
2316  for (i = 0; i < number_platforms; i++)
2317  {
2318  if (platforms[i] == (cl_platform_id) NULL)
2319  continue;
2320 
2321  status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2322  CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2323  if (status != CL_SUCCESS)
2324  continue;
2325 
2326  properties[0]=CL_CONTEXT_PLATFORM;
2327  properties[1]=(cl_context_properties) platforms[i];
2328  properties[2]=0;
2329  clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2330  devices,NULL,NULL,&status);
2331  if (status != CL_SUCCESS)
2332  continue;
2333 
2334  for (j = 0; j < number_devices; j++,next++)
2335  {
2337  device;
2338 
2339  device=AcquireMagickCLDevice();
2340  if (device == (MagickCLDevice) NULL)
2341  break;
2342 
2343  device->context=clEnv->contexts[i];
2344  device->deviceID=devices[j];
2345 
2346  openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,0,NULL,
2347  &length);
2348  device->platform_name=AcquireCriticalMemory(length*
2349  sizeof(*device->platform_name));
2350  openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,length,
2351  device->platform_name,NULL);
2352 
2353  openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,0,NULL,
2354  &length);
2355  device->vendor_name=AcquireCriticalMemory(length*
2356  sizeof(*device->vendor_name));
2357  openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,length,
2358  device->vendor_name,NULL);
2359 
2360  openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_NAME,0,NULL,
2361  &length);
2362  device->name=AcquireCriticalMemory(length*sizeof(*device->name));
2363  openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_NAME,length,
2364  device->name,NULL);
2365 
2366  openCL_library->clGetDeviceInfo(devices[j],CL_DRIVER_VERSION,0,NULL,
2367  &length);
2368  device->version=AcquireCriticalMemory(length*sizeof(*device->version));
2369  openCL_library->clGetDeviceInfo(devices[j],CL_DRIVER_VERSION,length,
2370  device->version,NULL);
2371 
2372  openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2373  sizeof(cl_uint),&device->max_clock_frequency,NULL);
2374 
2375  openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2376  sizeof(cl_uint),&device->max_compute_units,NULL);
2377 
2378  openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2379  sizeof(cl_device_type),&device->type,NULL);
2380 
2381  openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2382  sizeof(cl_ulong),&device->local_memory_size,NULL);
2383 
2384  clEnv->devices[next]=device;
2385  }
2386  }
2387  if (next != clEnv->number_devices)
2388  RelinquishMagickCLDevices(clEnv);
2389  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2390  devices=(cl_device_id *) RelinquishMagickMemory(devices);
2391 }
2392 
2393 MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2394  ExceptionInfo *exception)
2395 {
2396  LockSemaphoreInfo(clEnv->lock);
2397  if (clEnv->initialized != MagickFalse)
2398  {
2399  UnlockSemaphoreInfo(clEnv->lock);
2400  return(HasOpenCLDevices(clEnv,exception));
2401  }
2402  if (LoadOpenCLLibrary() != MagickFalse)
2403  {
2404  clEnv->library=openCL_library;
2405  LoadOpenCLDevices(clEnv);
2406  if (clEnv->number_devices > 0)
2407  AutoSelectOpenCLDevices(clEnv);
2408  }
2409  clEnv->initialized=MagickTrue;
2410  UnlockSemaphoreInfo(clEnv->lock);
2411  return(HasOpenCLDevices(clEnv,exception));
2412 }
2413 
2414 /*
2415 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2416 % %
2417 % %
2418 % %
2419 % L o a d O p e n C L L i b r a r y %
2420 % %
2421 % %
2422 % %
2423 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2424 %
2425 % LoadOpenCLLibrary() load and binds the OpenCL library.
2426 %
2427 % The format of the LoadOpenCLLibrary method is:
2428 %
2429 % MagickBooleanType LoadOpenCLLibrary(void)
2430 %
2431 */
2432 
2433 void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2434 {
2435  if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2436  return (void *) NULL;
2437 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2438  return (void *) GetProcAddress((HMODULE)library,functionName);
2439 #else
2440  return (void *) dlsym(library,functionName);
2441 #endif
2442 }
2443 
2444 static MagickBooleanType BindOpenCLFunctions()
2445 {
2446 #ifdef MAGICKCORE_OPENCL_MACOSX
2447 #define BIND(X) openCL_library->X= &X;
2448 #else
2449  (void) ResetMagickMemory(openCL_library,0,sizeof(MagickLibrary));
2450 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2451  openCL_library->library=(void *)LoadLibraryA("OpenCL.dll");
2452 #else
2453  openCL_library->library=(void *)dlopen("libOpenCL.so", RTLD_NOW);
2454 #endif
2455 #define BIND(X) \
2456  if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2457  return(MagickFalse);
2458 #endif
2459 
2460  if (openCL_library->library == (void*) NULL)
2461  return(MagickFalse);
2462 
2463  BIND(clGetPlatformIDs);
2464  BIND(clGetPlatformInfo);
2465 
2466  BIND(clGetDeviceIDs);
2467  BIND(clGetDeviceInfo);
2468 
2469  BIND(clCreateBuffer);
2470  BIND(clReleaseMemObject);
2471  BIND(clRetainMemObject);
2472 
2473  BIND(clCreateContext);
2474  BIND(clReleaseContext);
2475 
2476  BIND(clCreateCommandQueue);
2477  BIND(clReleaseCommandQueue);
2478  BIND(clFlush);
2479  BIND(clFinish);
2480 
2481  BIND(clCreateProgramWithSource);
2482  BIND(clCreateProgramWithBinary);
2483  BIND(clReleaseProgram);
2484  BIND(clBuildProgram);
2485  BIND(clGetProgramBuildInfo);
2486  BIND(clGetProgramInfo);
2487 
2488  BIND(clCreateKernel);
2489  BIND(clReleaseKernel);
2490  BIND(clSetKernelArg);
2491  BIND(clGetKernelInfo);
2492 
2493  BIND(clEnqueueReadBuffer);
2494  BIND(clEnqueueMapBuffer);
2495  BIND(clEnqueueUnmapMemObject);
2496  BIND(clEnqueueNDRangeKernel);
2497 
2498  BIND(clGetEventInfo);
2499  BIND(clWaitForEvents);
2500  BIND(clReleaseEvent);
2501  BIND(clRetainEvent);
2502  BIND(clSetEventCallback);
2503 
2504  BIND(clGetEventProfilingInfo);
2505 
2506  return(MagickTrue);
2507 }
2508 
2509 static MagickBooleanType LoadOpenCLLibrary(void)
2510 {
2511  openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2512  if (openCL_library == (MagickLibrary *) NULL)
2513  return(MagickFalse);
2514 
2515  if (BindOpenCLFunctions() == MagickFalse)
2516  {
2517  openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2518  return(MagickFalse);
2519  }
2520 
2521  return(MagickTrue);
2522 }
2523 
2524 /*
2525 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2526 % %
2527 % %
2528 % %
2529 + O p e n C L T e r m i n u s %
2530 % %
2531 % %
2532 % %
2533 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2534 %
2535 % OpenCLTerminus() destroys the OpenCL component.
2536 %
2537 % The format of the OpenCLTerminus method is:
2538 %
2539 % OpenCLTerminus(void)
2540 %
2541 */
2542 
2543 MagickPrivate void OpenCLTerminus()
2544 {
2545  DumpOpenCLProfileData();
2546  if (cache_directory != (char *) NULL)
2547  cache_directory=DestroyString(cache_directory);
2548  if (cache_directory_lock != (SemaphoreInfo *) NULL)
2549  RelinquishSemaphoreInfo(&cache_directory_lock);
2550  if (default_CLEnv != (MagickCLEnv) NULL)
2551  default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2552  if (openCL_lock != (SemaphoreInfo *) NULL)
2553  RelinquishSemaphoreInfo(&openCL_lock);
2554  if (openCL_library != (MagickLibrary *) NULL)
2555  {
2556  if (openCL_library->library != (void *) NULL)
2557  (void) lt_dlclose(openCL_library->library);
2558  openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2559  }
2560 }
2561 
2562 /*
2563 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2564 % %
2565 % %
2566 % %
2567 + O p e n C L T h r o w M a g i c k E x c e p t i o n %
2568 % %
2569 % %
2570 % %
2571 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2572 %
2573 % OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2574 % configuration file. If an error occurs, MagickFalse is returned
2575 % otherwise MagickTrue.
2576 %
2577 % The format of the OpenCLThrowMagickException method is:
2578 %
2579 % MagickBooleanType ThrowFileException(ExceptionInfo *exception,
2580 % const char *module,const char *function,const size_t line,
2581 % const ExceptionType severity,const char *tag,const char *format,...)
2582 %
2583 % A description of each parameter follows:
2584 %
2585 % o exception: the exception info.
2586 %
2587 % o filename: the source module filename.
2588 %
2589 % o function: the function name.
2590 %
2591 % o line: the line number of the source module.
2592 %
2593 % o severity: Specifies the numeric error category.
2594 %
2595 % o tag: the locale tag.
2596 %
2597 % o format: the output format.
2598 %
2599 */
2600 
2601 MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2602  MagickCLDevice device,ExceptionInfo *exception,const char *module,
2603  const char *function,const size_t line,const ExceptionType severity,
2604  const char *tag,const char *format,...)
2605 {
2607  status;
2608 
2609  assert(device != (MagickCLDevice) NULL);
2610  assert(exception != (ExceptionInfo *) NULL);
2611  assert(exception->signature == MagickCoreSignature);
2612 
2613  status=MagickTrue;
2614  if (severity != 0)
2615  {
2616  if (device->type == CL_DEVICE_TYPE_CPU)
2617  {
2618  /* Workaround for Intel OpenCL CPU runtime bug */
2619  /* Turn off OpenCL when a problem is detected! */
2620  if (strncmp(device->platform_name, "Intel",5) == 0)
2621  default_CLEnv->enabled=MagickFalse;
2622  }
2623  }
2624 
2625 #ifdef OPENCLLOG_ENABLED
2626  {
2627  va_list
2628  operands;
2629  va_start(operands,format);
2630  status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2631  format,operands);
2632  va_end(operands);
2633  }
2634 #else
2635  magick_unreferenced(module);
2636  magick_unreferenced(function);
2637  magick_unreferenced(line);
2638  magick_unreferenced(tag);
2639  magick_unreferenced(format);
2640 #endif
2641 
2642  return(status);
2643 }
2644 
2645 /*
2646 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2647 % %
2648 % %
2649 % %
2650 + R e c o r d P r o f i l e D a t a %
2651 % %
2652 % %
2653 % %
2654 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2655 %
2656 % RecordProfileData() records profile data.
2657 %
2658 % The format of the RecordProfileData method is:
2659 %
2660 % void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2661 % cl_event event)
2662 %
2663 % A description of each parameter follows:
2664 %
2665 % o device: the OpenCL device that did the operation.
2666 %
2667 % o event: the event that contains the profiling data.
2668 %
2669 */
2670 
2671 MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2672  cl_kernel kernel,cl_event event)
2673 {
2674  char
2675  *name;
2676 
2677  cl_int
2678  status;
2679 
2680  cl_ulong
2681  elapsed,
2682  end,
2683  start;
2684 
2686  profile_record;
2687 
2688  size_t
2689  i,
2690  length;
2691 
2692  if (device->profile_kernels == MagickFalse)
2693  return(MagickFalse);
2694  status=openCL_library->clWaitForEvents(1,&event);
2695  if (status != CL_SUCCESS)
2696  return(MagickFalse);
2697  status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2698  &length);
2699  if (status != CL_SUCCESS)
2700  return(MagickTrue);
2701  name=AcquireQuantumMemory(length,sizeof(*name));
2702  if (name == (char *) NULL)
2703  return(MagickTrue);
2704  start=end=elapsed=0;
2705  status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2706  name,(size_t *) NULL);
2707  status|=openCL_library->clGetEventProfilingInfo(event,
2708  CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2709  status|=openCL_library->clGetEventProfilingInfo(event,
2710  CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2711  if (status != CL_SUCCESS)
2712  {
2713  name=DestroyString(name);
2714  return(MagickTrue);
2715  }
2716  start/=1000; // usecs
2717  end/=1000; // usecs
2718  elapsed=end-start;
2719  LockSemaphoreInfo(device->lock);
2720  i=0;
2721  profile_record=(KernelProfileRecord) NULL;
2722  if (device->profile_records != (KernelProfileRecord *) NULL)
2723  {
2724  while (device->profile_records[i] != (KernelProfileRecord) NULL)
2725  {
2726  if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2727  {
2728  profile_record=device->profile_records[i];
2729  break;
2730  }
2731  i++;
2732  }
2733  }
2734  if (profile_record != (KernelProfileRecord) NULL)
2735  name=DestroyString(name);
2736  else
2737  {
2738  profile_record=AcquireMagickMemory(sizeof(*profile_record));
2739  (void) ResetMagickMemory(profile_record,0,sizeof(*profile_record));
2740  profile_record->kernel_name=name;
2741  device->profile_records=ResizeMagickMemory(device->profile_records,(i+2)*
2742  sizeof(*device->profile_records));
2743  device->profile_records[i]=profile_record;
2744  device->profile_records[i+1]=(KernelProfileRecord) NULL;
2745  }
2746  if ((elapsed < profile_record->min) || (profile_record->count == 0))
2747  profile_record->min=elapsed;
2748  if (elapsed > profile_record->max)
2749  profile_record->max=elapsed;
2750  profile_record->total+=elapsed;
2751  profile_record->count+=1;
2752  UnlockSemaphoreInfo(device->lock);
2753  return(MagickTrue);
2754 }
2755 
2756 /*
2757 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2758 % %
2759 % %
2760 % %
2761 + R e l e a s e O p e n C L C o m m a n d Q u e u e %
2762 % %
2763 % %
2764 % %
2765 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2766 %
2767 % ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2768 %
2769 % The format of the ReleaseOpenCLCommandQueue method is:
2770 %
2771 % void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2772 % cl_command_queue queue)
2773 %
2774 % A description of each parameter follows:
2775 %
2776 % o device: the OpenCL device.
2777 %
2778 % o queue: the OpenCL queue to be released.
2779 */
2780 
2781 MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2782  cl_command_queue queue)
2783 {
2784  if (queue == (cl_command_queue) NULL)
2785  return;
2786 
2787  assert(device != (MagickCLDevice) NULL);
2788  LockSemaphoreInfo(device->lock);
2789  if ((device->profile_kernels != MagickFalse) ||
2790  (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2791  {
2792  UnlockSemaphoreInfo(device->lock);
2793  openCL_library->clFinish(queue);
2794  (void) openCL_library->clReleaseCommandQueue(queue);
2795  }
2796  else
2797  {
2798  openCL_library->clFlush(queue);
2799  device->command_queues[++device->command_queues_index]=queue;
2800  UnlockSemaphoreInfo(device->lock);
2801  }
2802 }
2803 
2804 /*
2805 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2806 % %
2807 % %
2808 % %
2809 + R e l e a s e M a g i c k C L D e v i c e %
2810 % %
2811 % %
2812 % %
2813 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2814 %
2815 % ReleaseOpenCLDevice() returns the OpenCL device to the environment
2816 %
2817 % The format of the ReleaseOpenCLDevice method is:
2818 %
2819 % void ReleaseOpenCLDevice(MagickCLDevice device)
2820 %
2821 % A description of each parameter follows:
2822 %
2823 % o device: the OpenCL device to be released.
2824 %
2825 */
2826 
2827 MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2828 {
2829  assert(device != (MagickCLDevice) NULL);
2830  LockSemaphoreInfo(openCL_lock);
2831  device->requested--;
2832  UnlockSemaphoreInfo(openCL_lock);
2833 }
2834 
2835 /*
2836 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2837 % %
2838 % %
2839 % %
2840 + R e l i n q u i s h M a g i c k C L C a c h e I n f o %
2841 % %
2842 % %
2843 % %
2844 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2845 %
2846 % RelinquishMagickCLCacheInfo() frees memory acquired with
2847 % AcquireMagickCLCacheInfo()
2848 %
2849 % The format of the RelinquishMagickCLCacheInfo method is:
2850 %
2851 % MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2852 % const MagickBooleanType relinquish_pixels)
2853 %
2854 % A description of each parameter follows:
2855 %
2856 % o info: the OpenCL cache info.
2857 %
2858 % o relinquish_pixels: the pixels will be relinquish when set to true.
2859 %
2860 */
2861 
2862 static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2863  cl_event magick_unused(event),
2864  cl_int magick_unused(event_command_exec_status),void *user_data)
2865 {
2867  info;
2868 
2869  Quantum
2870  *pixels;
2871 
2872  ssize_t
2873  i;
2874 
2875  magick_unreferenced(event);
2876  magick_unreferenced(event_command_exec_status);
2877  info=(MagickCLCacheInfo) user_data;
2878  for (i=(ssize_t)info->event_count-1; i >= 0; i--)
2879  {
2880  cl_int
2881  event_status;
2882 
2883  cl_uint
2884  status;
2885 
2886  status=openCL_library->clGetEventInfo(info->events[i],
2887  CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
2888  NULL);
2889  if ((status == CL_SUCCESS) && (event_status != CL_COMPLETE))
2890  {
2891  openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
2892  &DestroyMagickCLCacheInfoAndPixels,info);
2893  return;
2894  }
2895  }
2896  pixels=info->pixels;
2897  DestroyMagickCLCacheInfo(info);
2898  (void) RelinquishAlignedMemory(pixels);
2899 }
2900 
2901 MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2902  MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2903 {
2904  if (info == (MagickCLCacheInfo) NULL)
2905  return((MagickCLCacheInfo) NULL);
2906  if (relinquish_pixels != MagickFalse)
2907  DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2908  else
2909  DestroyMagickCLCacheInfo(info);
2910  return((MagickCLCacheInfo) NULL);
2911 }
2912 
2913 /*
2914 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2915 % %
2916 % %
2917 % %
2918 % R e l i n q u i s h M a g i c k C L D e v i c e %
2919 % %
2920 % %
2921 % %
2922 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2923 %
2924 % RelinquishMagickCLDevice() releases the OpenCL device
2925 %
2926 % The format of the RelinquishMagickCLDevice method is:
2927 %
2928 % MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2929 %
2930 % A description of each parameter follows:
2931 %
2932 % o device: the OpenCL device to be released.
2933 %
2934 */
2935 
2936 static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2937 {
2938  if (device == (MagickCLDevice) NULL)
2939  return((MagickCLDevice) NULL);
2940 
2941  device->platform_name=RelinquishMagickMemory(device->platform_name);
2942  device->vendor_name=RelinquishMagickMemory(device->vendor_name);
2943  device->name=RelinquishMagickMemory(device->name);
2944  device->version=RelinquishMagickMemory(device->version);
2945  if (device->program != (cl_program) NULL)
2946  (void) openCL_library->clReleaseProgram(device->program);
2947  while (device->command_queues_index >= 0)
2948  (void) openCL_library->clReleaseCommandQueue(
2949  device->command_queues[device->command_queues_index--]);
2950  RelinquishSemaphoreInfo(&device->lock);
2951  return((MagickCLDevice) RelinquishMagickMemory(device));
2952 }
2953 
2954 /*
2955 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2956 % %
2957 % %
2958 % %
2959 % R e l i n q u i s h M a g i c k C L E n v %
2960 % %
2961 % %
2962 % %
2963 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2964 %
2965 % RelinquishMagickCLEnv() releases the OpenCL environment
2966 %
2967 % The format of the RelinquishMagickCLEnv method is:
2968 %
2969 % MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
2970 %
2971 % A description of each parameter follows:
2972 %
2973 % o clEnv: the OpenCL environment to be released.
2974 %
2975 */
2976 
2977 static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
2978 {
2979  if (clEnv == (MagickCLEnv) NULL)
2980  return((MagickCLEnv) NULL);
2981 
2982  RelinquishSemaphoreInfo(&clEnv->lock);
2983  RelinquishMagickCLDevices(clEnv);
2984  if (clEnv->contexts != (cl_context *) NULL)
2985  {
2986  ssize_t
2987  i;
2988 
2989  for (i=0; i < clEnv->number_contexts; i++)
2990  (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
2991  clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
2992  }
2993  return((MagickCLEnv) RelinquishMagickMemory(clEnv));
2994 }
2995 
2996 /*
2997 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2998 % %
2999 % %
3000 % %
3001 + R e q u e s t O p e n C L D e v i c e %
3002 % %
3003 % %
3004 % %
3005 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3006 %
3007 % RequestOpenCLDevice() returns one of the enabled OpenCL devices.
3008 %
3009 % The format of the RequestOpenCLDevice method is:
3010 %
3011 % MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3012 %
3013 % A description of each parameter follows:
3014 %
3015 % o clEnv: the OpenCL environment.
3016 */
3017 
3018 MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3019 {
3021  device;
3022 
3023  double
3024  score,
3025  best_score;
3026 
3027  size_t
3028  i;
3029 
3030  if (clEnv == (MagickCLEnv) NULL)
3031  return((MagickCLDevice) NULL);
3032 
3033  if (clEnv->number_devices == 1)
3034  {
3035  if (clEnv->devices[0]->enabled)
3036  return(clEnv->devices[0]);
3037  else
3038  return((MagickCLDevice) NULL);
3039  }
3040 
3041  device=(MagickCLDevice) NULL;
3042  best_score=0.0;
3043  LockSemaphoreInfo(openCL_lock);
3044  for (i = 0; i < clEnv->number_devices; i++)
3045  {
3046  if (clEnv->devices[i]->enabled == MagickFalse)
3047  continue;
3048 
3049  score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3050  clEnv->devices[i]->requested);
3051  if ((device == (MagickCLDevice) NULL) || (score < best_score))
3052  {
3053  device=clEnv->devices[i];
3054  best_score=score;
3055  }
3056  }
3057  if (device != (MagickCLDevice)NULL)
3058  device->requested++;
3059  UnlockSemaphoreInfo(openCL_lock);
3060 
3061  return(device);
3062 }
3063 
3064 /*
3065 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3066 % %
3067 % %
3068 % %
3069 % S e t O p e n C L D e v i c e E n a b l e d %
3070 % %
3071 % %
3072 % %
3073 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3074 %
3075 % SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
3076 %
3077 % The format of the SetOpenCLDeviceEnabled method is:
3078 %
3079 % void SetOpenCLDeviceEnabled(MagickCLDevice device,
3080 % MagickBooleanType value)
3081 %
3082 % A description of each parameter follows:
3083 %
3084 % o device: the OpenCL device.
3085 %
3086 % o value: determines if the device should be enabled or disabled.
3087 */
3088 
3090  const MagickBooleanType value)
3091 {
3092  if (device == (MagickCLDevice) NULL)
3093  return;
3094  device->enabled=value;
3095 }
3096 
3097 /*
3098 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3099 % %
3100 % %
3101 % %
3102 % S e t O p e n C L K e r n e l P r o f i l e E n a b l e d %
3103 % %
3104 % %
3105 % %
3106 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3107 %
3108 % SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
3109 % kernel profiling of a device.
3110 %
3111 % The format of the SetOpenCLKernelProfileEnabled method is:
3112 %
3113 % void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3114 % MagickBooleanType value)
3115 %
3116 % A description of each parameter follows:
3117 %
3118 % o device: the OpenCL device.
3119 %
3120 % o value: determines if kernel profiling for the device should be enabled
3121 % or disabled.
3122 */
3123 
3125  const MagickBooleanType value)
3126 {
3127  if (device == (MagickCLDevice) NULL)
3128  return;
3129  device->profile_kernels=value;
3130 }
3131 
3132 /*
3133 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3134 % %
3135 % %
3136 % %
3137 % S e t O p e n C L E n a b l e d %
3138 % %
3139 % %
3140 % %
3141 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3142 %
3143 % SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3144 %
3145 % The format of the SetOpenCLEnabled method is:
3146 %
3147 % void SetOpenCLEnabled(MagickBooleanType)
3148 %
3149 % A description of each parameter follows:
3150 %
3151 % o value: specify true to enable OpenCL acceleration
3152 */
3153 
3155 {
3156  MagickCLEnv
3157  clEnv;
3158 
3159  clEnv=GetCurrentOpenCLEnv();
3160  if (clEnv == (MagickCLEnv) NULL)
3161  return(MagickFalse);
3162  clEnv->enabled=value;
3163  return(clEnv->enabled);
3164 }
3165 
3166 #else
3167 
3169  const MagickCLDevice magick_unused(device))
3170 {
3171  magick_unreferenced(device);
3172  return(0.0);
3173 }
3174 
3176  const MagickCLDevice magick_unused(device))
3177 {
3178  magick_unreferenced(device);
3179  return(MagickFalse);
3180 }
3181 
3183  const MagickCLDevice magick_unused(device))
3184 {
3185  magick_unreferenced(device);
3186  return((const char *) NULL);
3187 }
3188 
3190  ExceptionInfo *magick_unused(exception))
3191 {
3192  magick_unreferenced(exception);
3193  if (length != (size_t *) NULL)
3194  *length=0;
3195  return((MagickCLDevice *) NULL);
3196 }
3197 
3199  const MagickCLDevice magick_unused(device))
3200 {
3201  magick_unreferenced(device);
3202  return(UndefinedCLDeviceType);
3203 }
3204 
3206  const MagickCLDevice magick_unused(device),size_t *length)
3207 {
3208  magick_unreferenced(device);
3209  if (length != (size_t *) NULL)
3210  *length=0;
3211  return((const KernelProfileRecord *) NULL);
3212 }
3213 
3215  const MagickCLDevice magick_unused(device))
3216 {
3217  magick_unreferenced(device);
3218  return((const char *) NULL);
3219 }
3220 
3222 {
3223  return(MagickFalse);
3224 }
3225 
3227  MagickCLDevice magick_unused(device),
3228  const MagickBooleanType magick_unused(value))
3229 {
3230  magick_unreferenced(device);
3231  magick_unreferenced(value);
3232 }
3233 
3235  const MagickBooleanType magick_unused(value))
3236 {
3237  magick_unreferenced(value);
3238  return(MagickFalse);
3239 }
3240 
3242  MagickCLDevice magick_unused(device),
3243  const MagickBooleanType magick_unused(value))
3244 {
3245  magick_unreferenced(device);
3246  magick_unreferenced(value);
3247 }
3248 #endif
MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType magick_unused(value))
Definition: opencl.c:3234
MagickExport Image * ResizeImage(const Image *image, const size_t columns, const size_t rows, const FilterType filter, ExceptionInfo *exception)
Definition: resize.c:2859
MagickExport Image * BlurImage(const Image *image, const double radius, const double sigma, ExceptionInfo *exception)
Definition: effect.c:770
static MagickThreadType GetMagickThreadId(void)
MagickExport ImageInfo * AcquireImageInfo(void)
Definition: image.c:341
MagickExport MagickBooleanType GetPathAttributes(const char *path, void *attributes)
Definition: utility.c:1165
struct _MagickCLDevice * MagickCLDevice
Definition: opencl.h:44
MagickCLDeviceType
Definition: opencl.h:25
unsigned long min
Definition: opencl.h:38
MagickExport void UnlockSemaphoreInfo(SemaphoreInfo *semaphore_info)
Definition: semaphore.c:450
MagickExport Image * UnsharpMaskImage(const Image *image, const double radius, const double sigma, const double gain, const double threshold, ExceptionInfo *exception)
Definition: effect.c:3906
char * kernel_name
Definition: opencl.h:35
#define ThrowFatalException(severity, tag)
MagickExport const char * GetOpenCLDeviceVendorName(const MagickCLDevice)
static int StringToInteger(const char *magick_restrict value)
size_t signature
Definition: exception.h:123
MagickExport SemaphoreInfo * AcquireSemaphoreInfo(void)
Definition: semaphore.c:192
static double StringToDouble(const char *magick_restrict string, char **magick_restrict sentinal)
#define MagickPI
Definition: image-private.h:30
MagickExport ExceptionInfo * AcquireExceptionInfo(void)
Definition: exception.c:108
MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice magick_unused(device), const MagickBooleanType magick_unused(value))
Definition: opencl.c:3226
MagickExport ssize_t FormatLocaleString(char *magick_restrict string, const size_t length, const char *magick_restrict format,...)
Definition: locale.c:473
#define MAGICKCORE_QUANTUM_DEPTH
Definition: magick-type.h:28
MagickCLCacheInfo opencl
MagickExport void * ResizeMagickMemory(void *memory, const size_t size)
Definition: memory.c:1217
void * MagickCLCacheInfo
#define MagickEpsilon
Definition: magick-type.h:110
MagickExport void * ResizeQuantumMemory(void *memory, const size_t count, const size_t quantum)
Definition: memory.c:1277
MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice magick_unused(device), const MagickBooleanType magick_unused(value))
Definition: opencl.c:3241
unsigned long max
Definition: opencl.h:38
char * size
Definition: image.h:381
Definition: image.h:151
ExceptionType
Definition: exception.h:27
#define MagickCoreSignature
MagickExport void LockSemaphoreInfo(SemaphoreInfo *semaphore_info)
Definition: semaphore.c:293
MagickExport MagickCLDevice * GetOpenCLDevices(size_t *length, ExceptionInfo *magick_unused(exception))
Definition: opencl.c:3189
MagickExport unsigned char * GetStringInfoDatum(const StringInfo *string_info)
Definition: string.c:1283
MagickExport MagickBooleanType ThrowMagickExceptionList(ExceptionInfo *exception, const char *module, const char *function, const size_t line, const ExceptionType severity, const char *tag, const char *format, va_list operands)
Definition: exception.c:1007
MagickBooleanType
Definition: magick-type.h:156
#define DirectorySeparator
Definition: studio.h:254
unsigned int MagickStatusType
Definition: magick-type.h:119
MagickExport char * AcquireString(const char *source)
Definition: string.c:124
MagickExport void * FileToBlob(const char *filename, const size_t extent, size_t *length, ExceptionInfo *exception)
Definition: blob.c:1270
MagickExport void * ResetMagickMemory(void *memory, int byte, const size_t size)
Definition: memory.c:1164
pid_t MagickThreadType
Definition: thread_.h:34
static int remove_utf8(const char *path)
MagickExport StringInfo * DestroyStringInfo(StringInfo *string_info)
Definition: string.c:837
MagickExport void * AcquireQuantumMemory(const size_t count, const size_t quantum)
Definition: memory.c:529
char filename[MagickPathExtent]
Definition: image.h:471
static FILE * fopen_utf8(const char *path, const char *mode)
MagickExport int LocaleNCompare(const char *p, const char *q, const size_t length)
Definition: locale.c:1509
#define MaxTextExtent
#define magick_unused(x)
size_t MagickSizeType
Definition: magick-type.h:128
#define MagickPathExtent
void * cache
Definition: image.h:294
MagickExport void * RelinquishAlignedMemory(void *memory)
Definition: memory.c:1001
MagickExport MagickBooleanType static void * AcquireCriticalMemory(const size_t size)
MagickExport Image * ReadImage(const ImageInfo *image_info, ExceptionInfo *exception)
Definition: constitute.c:358
unsigned long count
Definition: opencl.h:38
MagickExport MagickBooleanType ThrowMagickException(ExceptionInfo *exception, const char *module, const char *function, const size_t line, const ExceptionType severity, const char *tag, const char *format,...)
Definition: exception.c:1058
#define QuantumScale
Definition: magick-type.h:113
MagickExport char * GetEnvironmentValue(const char *name)
Definition: string.c:1250
MagickExport StringInfo * ConfigureFileToStringInfo(const char *filename)
Definition: string.c:576
MagickExport MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice magick_unused(device))
Definition: opencl.c:3198
unsigned long total
Definition: opencl.h:38
MagickExport size_t CopyMagickString(char *destination, const char *source, const size_t length)
Definition: string.c:742
#define MaxMap
Definition: magick-type.h:75
MagickExport const KernelProfileRecord * GetOpenCLKernelProfileRecords(const MagickCLDevice magick_unused(device), size_t *length)
Definition: opencl.c:3205
MagickExport int LocaleCompare(const char *p, const char *q)
Definition: locale.c:1409
#define GetMagickModule()
Definition: log.h:28
MagickExport const char * GetOpenCLDeviceName(const MagickCLDevice magick_unused(device))
Definition: opencl.c:3182
MagickExport double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice magick_unused(device))
Definition: opencl.c:3168
MagickExport void GetNextToken(const char *start, const char **end, const size_t extent, char *token)
Definition: token.c:171
unsigned short Quantum
Definition: magick-type.h:82
MagickExport char * DestroyString(char *string)
Definition: string.c:810
MagickExport void * AcquireMagickMemory(const size_t size)
Definition: memory.c:458
struct _KernelProfileRecord * KernelProfileRecord
MagickExport void ActivateSemaphoreInfo(SemaphoreInfo **semaphore_info)
Definition: semaphore.c:97
#define MagickMin(x, y)
Definition: image-private.h:27
MagickExport void * RelinquishMagickMemory(void *memory)
Definition: memory.c:1038
#define magick_unreferenced(x)
MagickExport MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice magick_unused(device))
Definition: opencl.c:3175
MagickExport char * CloneString(char **destination, const char *source)
Definition: string.c:270
#define MagickPrivate
MagickExport const char * GetOpenCLDeviceVersion(const MagickCLDevice magick_unused(device))
Definition: opencl.c:3214
#define MagickExport
MagickExport void RelinquishSemaphoreInfo(SemaphoreInfo **semaphore_info)
Definition: semaphore.c:351
MagickExport Image * DestroyImage(Image *image)
Definition: image.c:1182
MagickExport char * ConstantString(const char *source)
Definition: string.c:687
MagickExport MagickBooleanType BlobToFile(char *filename, const void *blob, const size_t length, ExceptionInfo *exception)
Definition: blob.c:339
MagickExport MagickBooleanType GetOpenCLEnabled(void)
Definition: opencl.c:3221
#define QuantumRange
Definition: magick-type.h:83
MagickExport ExceptionInfo * DestroyExceptionInfo(ExceptionInfo *exception)
Definition: exception.c:417