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