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