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