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