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