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