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