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