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