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