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