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