MagickCore 7.1.2
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_%.20g.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 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
1694 info->events[info->event_count-1]=event;
1695 UnlockSemaphoreInfo(info->events_semaphore);
1696 return(MagickTrue);
1697}
1698
1699MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1700 cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
1701 const size_t *lsize,const Image *input_image,const Image *output_image,
1702 MagickBooleanType flush,ExceptionInfo *exception)
1703{
1704 CacheInfo
1705 *output_info,
1706 *input_info;
1707
1708 cl_event
1709 event,
1710 *events;
1711
1712 cl_int
1713 status;
1714
1715 cl_uint
1716 event_count;
1717
1718 assert(input_image != (const Image *) NULL);
1719 input_info=(CacheInfo *) input_image->cache;
1720 assert(input_info != (CacheInfo *) NULL);
1721 assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1722 output_info=(CacheInfo *) NULL;
1723 if (output_image == (const Image *) NULL)
1724 events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
1725 &event_count);
1726 else
1727 {
1728 output_info=(CacheInfo *) output_image->cache;
1729 assert(output_info != (CacheInfo *) NULL);
1730 assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1731 events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
1732 &event_count);
1733 }
1734 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1735 gsize,lsize,event_count,events,&event);
1736 /* This can fail due to memory issues and calling clFinish might help. */
1737 if ((status != CL_SUCCESS) && (event_count > 0))
1738 {
1739 openCL_library->clFinish(queue);
1740 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
1741 offset,gsize,lsize,event_count,events,&event);
1742 }
1743 events=(cl_event *) RelinquishMagickMemory(events);
1744 if (status != CL_SUCCESS)
1745 {
1746 (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1747 GetMagickModule(),ResourceLimitWarning,
1748 "clEnqueueNDRangeKernel failed.","'%s'",".");
1749 return(MagickFalse);
1750 }
1751 if (flush != MagickFalse)
1752 openCL_library->clFlush(queue);
1753 if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1754 {
1755 if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
1756 {
1757 if (output_info != (CacheInfo *) NULL)
1758 (void) RegisterCacheEvent(output_info->opencl,event);
1759 }
1760 }
1761 openCL_library->clReleaseEvent(event);
1762 return(MagickTrue);
1763}
1764
1765/*
1766%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1767% %
1768% %
1769% %
1770+ G e t C u r r e n t O p e n C L E n v %
1771% %
1772% %
1773% %
1774%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1775%
1776% GetCurrentOpenCLEnv() returns the current OpenCL env
1777%
1778% The format of the GetCurrentOpenCLEnv method is:
1779%
1780% MagickCLEnv GetCurrentOpenCLEnv()
1781%
1782*/
1783
1784MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
1785{
1786 if (default_CLEnv != (MagickCLEnv) NULL)
1787 {
1788 if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1789 (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1790 return((MagickCLEnv) NULL);
1791 else
1792 return(default_CLEnv);
1793 }
1794
1795 if (GetOpenCLCacheDirectory() == (char *) NULL)
1796 return((MagickCLEnv) NULL);
1797
1798 if (openCL_lock == (SemaphoreInfo *) NULL)
1799 ActivateSemaphoreInfo(&openCL_lock);
1800
1801 LockSemaphoreInfo(openCL_lock);
1802 if (default_CLEnv == (MagickCLEnv) NULL)
1803 default_CLEnv=AcquireMagickCLEnv();
1804 UnlockSemaphoreInfo(openCL_lock);
1805
1806 return(default_CLEnv);
1807}
1808
1809/*
1810%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1811% %
1812% %
1813% %
1814% 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 %
1815% %
1816% %
1817% %
1818%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1819%
1820% GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
1821% device. The score is determined by the duration of the micro benchmark so
1822% that means a lower score is better than a higher score.
1823%
1824% The format of the GetOpenCLDeviceBenchmarkScore method is:
1825%
1826% double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
1827%
1828% A description of each parameter follows:
1829%
1830% o device: the OpenCL device.
1831*/
1832
1833MagickExport double GetOpenCLDeviceBenchmarkScore(
1834 const MagickCLDevice device)
1835{
1836 if (device == (MagickCLDevice) NULL)
1837 return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1838 return(device->score);
1839}
1840
1841/*
1842%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1843% %
1844% %
1845% %
1846% G e t O p e n C L D e v i c e E n a b l e d %
1847% %
1848% %
1849% %
1850%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1851%
1852% GetOpenCLDeviceEnabled() returns true if the device is enabled.
1853%
1854% The format of the GetOpenCLDeviceEnabled method is:
1855%
1856% MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
1857%
1858% A description of each parameter follows:
1859%
1860% o device: the OpenCL device.
1861*/
1862
1863MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
1864 const MagickCLDevice device)
1865{
1866 if (device == (MagickCLDevice) NULL)
1867 return(MagickFalse);
1868 return(device->enabled);
1869}
1870
1871/*
1872%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1873% %
1874% %
1875% %
1876% G e t O p e n C L D e v i c e N a m e %
1877% %
1878% %
1879% %
1880%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1881%
1882% GetOpenCLDeviceName() returns the name of the device.
1883%
1884% The format of the GetOpenCLDeviceName method is:
1885%
1886% const char *GetOpenCLDeviceName(const MagickCLDevice device)
1887%
1888% A description of each parameter follows:
1889%
1890% o device: the OpenCL device.
1891*/
1892
1893MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
1894{
1895 if (device == (MagickCLDevice) NULL)
1896 return((const char *) NULL);
1897 return(device->name);
1898}
1899
1900/*
1901%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1902% %
1903% %
1904% %
1905% 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 %
1906% %
1907% %
1908% %
1909%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1910%
1911% GetOpenCLDeviceVendorName() returns the vendor name of the device.
1912%
1913% The format of the GetOpenCLDeviceVendorName method is:
1914%
1915% const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1916%
1917% A description of each parameter follows:
1918%
1919% o device: the OpenCL device.
1920*/
1921
1922MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1923{
1924 if (device == (MagickCLDevice) NULL)
1925 return((const char *) NULL);
1926 return(device->vendor_name);
1927}
1928
1929/*
1930%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1931% %
1932% %
1933% %
1934% G e t O p e n C L D e v i c e s %
1935% %
1936% %
1937% %
1938%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1939%
1940% GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
1941% value of length to the number of devices that are available.
1942%
1943% The format of the GetOpenCLDevices method is:
1944%
1945% const MagickCLDevice *GetOpenCLDevices(size_t *length,
1946% ExceptionInfo *exception)
1947%
1948% A description of each parameter follows:
1949%
1950% o length: the number of device.
1951%
1952% o exception: return any errors or warnings in this structure.
1953%
1954*/
1955
1956MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
1957 ExceptionInfo *exception)
1958{
1959 MagickCLEnv
1960 clEnv;
1961
1962 clEnv=GetCurrentOpenCLEnv();
1963 if (clEnv == (MagickCLEnv) NULL)
1964 {
1965 if (length != (size_t *) NULL)
1966 *length=0;
1967 return((MagickCLDevice *) NULL);
1968 }
1969 InitializeOpenCL(clEnv,exception);
1970 if (length != (size_t *) NULL)
1971 *length=clEnv->number_devices;
1972 return(clEnv->devices);
1973}
1974
1975/*
1976%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1977% %
1978% %
1979% %
1980% G e t O p e n C L D e v i c e T y p e %
1981% %
1982% %
1983% %
1984%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1985%
1986% GetOpenCLDeviceType() returns the type of the device.
1987%
1988% The format of the GetOpenCLDeviceType method is:
1989%
1990% MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
1991%
1992% A description of each parameter follows:
1993%
1994% o device: the OpenCL device.
1995*/
1996
1997MagickExport MagickCLDeviceType GetOpenCLDeviceType(
1998 const MagickCLDevice device)
1999{
2000 if (device == (MagickCLDevice) NULL)
2001 return(UndefinedCLDeviceType);
2002 if (device->type == CL_DEVICE_TYPE_GPU)
2003 return(GpuCLDeviceType);
2004 if (device->type == CL_DEVICE_TYPE_CPU)
2005 return(CpuCLDeviceType);
2006 return(UndefinedCLDeviceType);
2007}
2008
2009/*
2010%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2011% %
2012% %
2013% %
2014% G e t O p e n C L D e v i c e V e r s i o n %
2015% %
2016% %
2017% %
2018%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2019%
2020% GetOpenCLDeviceVersion() returns the version of the device.
2021%
2022% The format of the GetOpenCLDeviceName method is:
2023%
2024% const char *GetOpenCLDeviceVersion(MagickCLDevice device)
2025%
2026% A description of each parameter follows:
2027%
2028% o device: the OpenCL device.
2029*/
2030
2031MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
2032{
2033 if (device == (MagickCLDevice) NULL)
2034 return((const char *) NULL);
2035 return(device->version);
2036}
2037
2038/*
2039%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2040% %
2041% %
2042% %
2043% G e t O p e n C L E n a b l e d %
2044% %
2045% %
2046% %
2047%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2048%
2049% GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
2050%
2051% The format of the GetOpenCLEnabled method is:
2052%
2053% MagickBooleanType GetOpenCLEnabled()
2054%
2055*/
2056
2057MagickExport MagickBooleanType GetOpenCLEnabled(void)
2058{
2059 MagickCLEnv
2060 clEnv;
2061
2062 clEnv=GetCurrentOpenCLEnv();
2063 if (clEnv == (MagickCLEnv) NULL)
2064 return(MagickFalse);
2065 return(clEnv->enabled);
2066}
2067
2068/*
2069%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2070% %
2071% %
2072% %
2073% 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 %
2074% %
2075% %
2076% %
2077%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2078%
2079% GetOpenCLKernelProfileRecords() returns the profile records for the
2080% specified device and sets length to the number of profile records.
2081%
2082% The format of the GetOpenCLKernelProfileRecords method is:
2083%
2084% const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
2085%
2086% A description of each parameter follows:
2087%
2088% o length: the number of profiles records.
2089*/
2090
2091MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
2092 const MagickCLDevice device,size_t *length)
2093{
2094 if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
2095 (KernelProfileRecord *) NULL))
2096 {
2097 if (length != (size_t *) NULL)
2098 *length=0;
2099 return((const KernelProfileRecord *) NULL);
2100 }
2101 if (length != (size_t *) NULL)
2102 {
2103 *length=0;
2104 LockSemaphoreInfo(device->lock);
2105 while (device->profile_records[*length] != (KernelProfileRecord) NULL)
2106 *length=*length+1;
2107 UnlockSemaphoreInfo(device->lock);
2108 }
2109 return(device->profile_records);
2110}
2111
2112/*
2113%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2114% %
2115% %
2116% %
2117% H a s O p e n C L D e v i c e s %
2118% %
2119% %
2120% %
2121%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2122%
2123% HasOpenCLDevices() checks if the OpenCL environment has devices that are
2124% enabled and compiles the kernel for the device when necessary. False will be
2125% returned if no enabled devices could be found
2126%
2127% The format of the HasOpenCLDevices method is:
2128%
2129% MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2130% ExceptionInfo exception)
2131%
2132% A description of each parameter follows:
2133%
2134% o clEnv: the OpenCL environment.
2135%
2136% o exception: return any errors or warnings in this structure.
2137%
2138*/
2139
2140static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2141 ExceptionInfo *exception)
2142{
2143 char
2144 *accelerateKernelsBuffer,
2145 options[MagickPathExtent];
2146
2147 MagickBooleanType
2148 status;
2149
2150 size_t
2151 i;
2152
2153 size_t
2154 signature;
2155
2156 /* Check if there are enabled devices */
2157 for (i = 0; i < clEnv->number_devices; i++)
2158 {
2159 if ((clEnv->devices[i]->enabled != MagickFalse))
2160 break;
2161 }
2162 if (i == clEnv->number_devices)
2163 return(MagickFalse);
2164
2165 /* Check if we need to compile a kernel for one of the devices */
2166 status=MagickTrue;
2167 for (i = 0; i < clEnv->number_devices; i++)
2168 {
2169 if ((clEnv->devices[i]->enabled != MagickFalse) &&
2170 (clEnv->devices[i]->program == (cl_program) NULL))
2171 {
2172 status=MagickFalse;
2173 break;
2174 }
2175 }
2176 if (status != MagickFalse)
2177 return(MagickTrue);
2178
2179 /* Get additional options */
2180 (void) FormatLocaleString(options,MagickPathExtent,CLOptions,
2181 (float)QuantumRange,(float)CLCharQuantumScale,(float)MagickEpsilon,
2182 (float)MagickPI,(unsigned int)MaxMap,(unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2183
2184 signature=StringSignature(options);
2185 accelerateKernelsBuffer=(char*) AcquireQuantumMemory(1,
2186 strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2187 if (accelerateKernelsBuffer == (char*) NULL)
2188 return(MagickFalse);
2189 (void) FormatLocaleString(accelerateKernelsBuffer,strlen(accelerateKernels)+
2190 strlen(accelerateKernels2)+1,"%s%s",accelerateKernels,accelerateKernels2);
2191 signature^=StringSignature(accelerateKernelsBuffer);
2192
2193 status=MagickTrue;
2194 for (i = 0; i < clEnv->number_devices; i++)
2195 {
2196 MagickCLDevice
2197 device;
2198
2199 size_t
2200 device_signature;
2201
2202 device=clEnv->devices[i];
2203 if ((device->enabled == MagickFalse) ||
2204 (device->program != (cl_program) NULL))
2205 continue;
2206
2207 LockSemaphoreInfo(device->lock);
2208 if (device->program != (cl_program) NULL)
2209 {
2210 UnlockSemaphoreInfo(device->lock);
2211 continue;
2212 }
2213 device_signature=signature;
2214 device_signature^=StringSignature(device->platform_name);
2215 status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2216 device_signature,exception);
2217 UnlockSemaphoreInfo(device->lock);
2218 if (status == MagickFalse)
2219 break;
2220 }
2221 accelerateKernelsBuffer=(char *) RelinquishMagickMemory(
2222 accelerateKernelsBuffer);
2223 return(status);
2224}
2225
2226/*
2227%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2228% %
2229% %
2230% %
2231+ I n i t i a l i z e O p e n C L %
2232% %
2233% %
2234% %
2235%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2236%
2237% InitializeOpenCL() is used to initialize the OpenCL environment. This method
2238% makes sure the devices are properly initialized and benchmarked.
2239%
2240% The format of the InitializeOpenCL method is:
2241%
2242% MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2243%
2244% A description of each parameter follows:
2245%
2246% o exception: return any errors or warnings in this structure.
2247%
2248*/
2249
2250static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2251{
2252 char
2253 version[MagickPathExtent];
2254
2255 cl_uint
2256 num;
2257
2258 if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2259 MagickPathExtent,version,NULL) != CL_SUCCESS)
2260 return(0);
2261 if (strncmp(version,"OpenCL 1.0 ",11) == 0)
2262 return(0);
2263 if (clEnv->library->clGetDeviceIDs(platform,
2264 CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2265 return(0);
2266 return(num);
2267}
2268
2269static inline char *GetOpenCLPlatformString(cl_platform_id platform,
2270 cl_platform_info param_name)
2271{
2272 char
2273 *value;
2274
2275 size_t
2276 length;
2277
2278 openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
2279 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2280 openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
2281 return(value);
2282}
2283
2284static inline char *GetOpenCLDeviceString(cl_device_id device,
2285 cl_device_info param_name)
2286{
2287 char
2288 *value;
2289
2290 size_t
2291 length;
2292
2293 openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
2294 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2295 openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
2296 return(value);
2297}
2298
2299static void LoadOpenCLDevices(MagickCLEnv clEnv)
2300{
2301 cl_context_properties
2302 properties[3];
2303
2304 cl_device_id
2305 *devices;
2306
2307 cl_int
2308 status;
2309
2310 cl_platform_id
2311 *platforms;
2312
2313 cl_uint
2314 i,
2315 j,
2316 next,
2317 number_devices,
2318 number_platforms;
2319
2320 number_platforms=0;
2321 if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2322 return;
2323 if (number_platforms == 0)
2324 return;
2325 platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
2326 sizeof(cl_platform_id));
2327 if (platforms == (cl_platform_id *) NULL)
2328 return;
2329 if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2330 {
2331 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2332 return;
2333 }
2334 for (i = 0; i < number_platforms; i++)
2335 {
2336 number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2337 if (number_devices == 0)
2338 platforms[i]=(cl_platform_id) NULL;
2339 else
2340 clEnv->number_devices+=number_devices;
2341 }
2342 if (clEnv->number_devices == 0)
2343 {
2344 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2345 return;
2346 }
2347 clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2348 sizeof(MagickCLDevice));
2349 if (clEnv->devices == (MagickCLDevice *) NULL)
2350 {
2351 RelinquishMagickCLDevices(clEnv);
2352 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2353 return;
2354 }
2355 (void) memset(clEnv->devices,0,clEnv->number_devices*sizeof(MagickCLDevice));
2356 devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2357 sizeof(cl_device_id));
2358 if (devices == (cl_device_id *) NULL)
2359 {
2360 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2361 RelinquishMagickCLDevices(clEnv);
2362 return;
2363 }
2364 (void) memset(devices,0,clEnv->number_devices*sizeof(cl_device_id));
2365 clEnv->number_contexts=(size_t) number_platforms;
2366 clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2367 sizeof(cl_context));
2368 if (clEnv->contexts == (cl_context *) NULL)
2369 {
2370 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2371 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2372 RelinquishMagickCLDevices(clEnv);
2373 return;
2374 }
2375 (void) memset(clEnv->contexts,0,clEnv->number_contexts*sizeof(cl_context));
2376 next=0;
2377 for (i = 0; i < number_platforms; i++)
2378 {
2379 if (platforms[i] == (cl_platform_id) NULL)
2380 continue;
2381
2382 status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2383 CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2384 if (status != CL_SUCCESS)
2385 continue;
2386
2387 properties[0]=CL_CONTEXT_PLATFORM;
2388 properties[1]=(cl_context_properties) platforms[i];
2389 properties[2]=0;
2390 clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2391 devices,NULL,NULL,&status);
2392 if (status != CL_SUCCESS)
2393 continue;
2394
2395 for (j = 0; j < number_devices; j++,next++)
2396 {
2397 MagickCLDevice
2398 device;
2399
2400 device=AcquireMagickCLDevice();
2401 if (device == (MagickCLDevice) NULL)
2402 break;
2403
2404 device->context=clEnv->contexts[i];
2405 device->deviceID=devices[j];
2406
2407 device->platform_name=GetOpenCLPlatformString(platforms[i],
2408 CL_PLATFORM_NAME);
2409
2410 device->vendor_name=GetOpenCLPlatformString(platforms[i],
2411 CL_PLATFORM_VENDOR);
2412
2413 device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);
2414
2415 device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);
2416
2417 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2418 sizeof(cl_uint),&device->max_clock_frequency,NULL);
2419
2420 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2421 sizeof(cl_uint),&device->max_compute_units,NULL);
2422
2423 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2424 sizeof(cl_device_type),&device->type,NULL);
2425
2426 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2427 sizeof(cl_ulong),&device->local_memory_size,NULL);
2428
2429 clEnv->devices[next]=device;
2430 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
2431 "Found device: %s (%s)",device->name,device->platform_name);
2432 }
2433 }
2434 if (next != clEnv->number_devices)
2435 RelinquishMagickCLDevices(clEnv);
2436 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2437 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2438}
2439
2440MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2441 ExceptionInfo *exception)
2442{
2443 LockSemaphoreInfo(clEnv->lock);
2444 if (clEnv->initialized != MagickFalse)
2445 {
2446 UnlockSemaphoreInfo(clEnv->lock);
2447 return(HasOpenCLDevices(clEnv,exception));
2448 }
2449 if (LoadOpenCLLibrary() != MagickFalse)
2450 {
2451 clEnv->library=openCL_library;
2452 LoadOpenCLDevices(clEnv);
2453 if (clEnv->number_devices > 0)
2454 AutoSelectOpenCLDevices(clEnv);
2455 }
2456 clEnv->initialized=MagickTrue;
2457 UnlockSemaphoreInfo(clEnv->lock);
2458 return(HasOpenCLDevices(clEnv,exception));
2459}
2460
2461/*
2462%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2463% %
2464% %
2465% %
2466% L o a d O p e n C L L i b r a r y %
2467% %
2468% %
2469% %
2470%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2471%
2472% LoadOpenCLLibrary() load and binds the OpenCL library.
2473%
2474% The format of the LoadOpenCLLibrary method is:
2475%
2476% MagickBooleanType LoadOpenCLLibrary(void)
2477%
2478*/
2479
2480void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2481{
2482 if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2483 return (void *) NULL;
2484 return lt_dlsym(library,functionName);
2485}
2486
2487static MagickBooleanType BindOpenCLFunctions()
2488{
2489#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
2490#define BIND(X) openCL_library->X= &X;
2491#else
2492 (void) memset(openCL_library,0,sizeof(MagickLibrary));
2493#ifdef MAGICKCORE_WINDOWS_SUPPORT
2494 openCL_library->library=(void *)lt_dlopen("OpenCL.dll");
2495#else
2496 openCL_library->library=(void *)lt_dlopen("libOpenCL.so");
2497#endif
2498#define BIND(X) \
2499 if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2500 return(MagickFalse);
2501#endif
2502
2503 if (openCL_library->library == (void*) NULL)
2504 return(MagickFalse);
2505
2506 BIND(clGetPlatformIDs);
2507 BIND(clGetPlatformInfo);
2508
2509 BIND(clGetDeviceIDs);
2510 BIND(clGetDeviceInfo);
2511
2512 BIND(clCreateBuffer);
2513 BIND(clReleaseMemObject);
2514 BIND(clRetainMemObject);
2515
2516 BIND(clCreateContext);
2517 BIND(clReleaseContext);
2518
2519 BIND(clCreateCommandQueue);
2520 BIND(clReleaseCommandQueue);
2521 BIND(clFlush);
2522 BIND(clFinish);
2523
2524 BIND(clCreateProgramWithSource);
2525 BIND(clCreateProgramWithBinary);
2526 BIND(clReleaseProgram);
2527 BIND(clBuildProgram);
2528 BIND(clGetProgramBuildInfo);
2529 BIND(clGetProgramInfo);
2530
2531 BIND(clCreateKernel);
2532 BIND(clReleaseKernel);
2533 BIND(clSetKernelArg);
2534 BIND(clGetKernelInfo);
2535
2536 BIND(clEnqueueReadBuffer);
2537 BIND(clEnqueueMapBuffer);
2538 BIND(clEnqueueUnmapMemObject);
2539 BIND(clEnqueueNDRangeKernel);
2540
2541 BIND(clGetEventInfo);
2542 BIND(clWaitForEvents);
2543 BIND(clReleaseEvent);
2544 BIND(clRetainEvent);
2545 BIND(clSetEventCallback);
2546
2547 BIND(clGetEventProfilingInfo);
2548
2549 return(MagickTrue);
2550}
2551
2552static MagickBooleanType LoadOpenCLLibrary(void)
2553{
2554 openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2555 if (openCL_library == (MagickLibrary *) NULL)
2556 return(MagickFalse);
2557
2558 if (BindOpenCLFunctions() == MagickFalse)
2559 {
2560 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2561 return(MagickFalse);
2562 }
2563
2564 return(MagickTrue);
2565}
2566
2567/*
2568%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2569% %
2570% %
2571% %
2572+ O p e n C L T e r m i n u s %
2573% %
2574% %
2575% %
2576%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2577%
2578% OpenCLTerminus() destroys the OpenCL component.
2579%
2580% The format of the OpenCLTerminus method is:
2581%
2582% OpenCLTerminus(void)
2583%
2584*/
2585
2586MagickPrivate void OpenCLTerminus()
2587{
2588 DumpOpenCLProfileData();
2589 if (cache_directory != (char *) NULL)
2590 cache_directory=DestroyString(cache_directory);
2591 if (cache_directory_lock != (SemaphoreInfo *) NULL)
2592 RelinquishSemaphoreInfo(&cache_directory_lock);
2593 if (default_CLEnv != (MagickCLEnv) NULL)
2594 default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2595 if (openCL_lock != (SemaphoreInfo *) NULL)
2596 RelinquishSemaphoreInfo(&openCL_lock);
2597 if (openCL_library != (MagickLibrary *) NULL)
2598 {
2599 if (openCL_library->library != (void *) NULL)
2600 (void) lt_dlclose(openCL_library->library);
2601 openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2602 }
2603}
2604
2605/*
2606%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2607% %
2608% %
2609% %
2610+ 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 %
2611% %
2612% %
2613% %
2614%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2615%
2616% OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2617% configuration file. If an error occurs, MagickFalse is returned
2618% otherwise MagickTrue.
2619%
2620% The format of the OpenCLThrowMagickException method is:
2621%
2622% MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2623% const char *module,const char *function,const size_t line,
2624% const ExceptionType severity,const char *tag,const char *format,...)
2625%
2626% A description of each parameter follows:
2627%
2628% o exception: the exception info.
2629%
2630% o filename: the source module filename.
2631%
2632% o function: the function name.
2633%
2634% o line: the line number of the source module.
2635%
2636% o severity: Specifies the numeric error category.
2637%
2638% o tag: the locale tag.
2639%
2640% o format: the output format.
2641%
2642*/
2643
2644MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2645 MagickCLDevice device,ExceptionInfo *exception,const char *module,
2646 const char *function,const size_t line,const ExceptionType severity,
2647 const char *tag,const char *format,...)
2648{
2649 MagickBooleanType
2650 status;
2651
2652 assert(device != (MagickCLDevice) NULL);
2653 assert(exception != (ExceptionInfo *) NULL);
2654 assert(exception->signature == MagickCoreSignature);
2655 (void) exception;
2656 status=MagickTrue;
2657 if (severity != 0)
2658 {
2659 if (device->type == CL_DEVICE_TYPE_CPU)
2660 {
2661 /* Workaround for Intel OpenCL CPU runtime bug */
2662 /* Turn off OpenCL when a problem is detected! */
2663 if (strncmp(device->platform_name,"Intel",5) == 0)
2664 default_CLEnv->enabled=MagickFalse;
2665 }
2666 }
2667
2668#ifdef OPENCLLOG_ENABLED
2669 {
2670 va_list
2671 operands;
2672 va_start(operands,format);
2673 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2674 format,operands);
2675 va_end(operands);
2676 }
2677#else
2678 magick_unreferenced(module);
2679 magick_unreferenced(function);
2680 magick_unreferenced(line);
2681 magick_unreferenced(tag);
2682 magick_unreferenced(format);
2683#endif
2684
2685 return(status);
2686}
2687
2688/*
2689%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2690% %
2691% %
2692% %
2693+ R e c o r d P r o f i l e D a t a %
2694% %
2695% %
2696% %
2697%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2698%
2699% RecordProfileData() records profile data.
2700%
2701% The format of the RecordProfileData method is:
2702%
2703% void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2704% cl_event event)
2705%
2706% A description of each parameter follows:
2707%
2708% o device: the OpenCL device that did the operation.
2709%
2710% o event: the event that contains the profiling data.
2711%
2712*/
2713
2714MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2715 cl_kernel kernel,cl_event event)
2716{
2717 char
2718 *name;
2719
2720 cl_int
2721 status;
2722
2723 cl_ulong
2724 elapsed,
2725 end,
2726 start;
2727
2728 KernelProfileRecord
2729 profile_record;
2730
2731 size_t
2732 i,
2733 length;
2734
2735 if (device->profile_kernels == MagickFalse)
2736 return(MagickFalse);
2737 status=openCL_library->clWaitForEvents(1,&event);
2738 if (status != CL_SUCCESS)
2739 return(MagickFalse);
2740 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2741 &length);
2742 if (status != CL_SUCCESS)
2743 return(MagickTrue);
2744 name=(char *) AcquireQuantumMemory(length,sizeof(*name));
2745 if (name == (char *) NULL)
2746 return(MagickTrue);
2747 start=end=elapsed=0;
2748 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2749 name,(size_t *) NULL);
2750 status|=openCL_library->clGetEventProfilingInfo(event,
2751 CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2752 status|=openCL_library->clGetEventProfilingInfo(event,
2753 CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2754 if (status != CL_SUCCESS)
2755 {
2756 name=DestroyString(name);
2757 return(MagickTrue);
2758 }
2759 start/=1000; /* usecs */
2760 end/=1000;
2761 elapsed=end-start;
2762 LockSemaphoreInfo(device->lock);
2763 i=0;
2764 profile_record=(KernelProfileRecord) NULL;
2765 if (device->profile_records != (KernelProfileRecord *) NULL)
2766 {
2767 while (device->profile_records[i] != (KernelProfileRecord) NULL)
2768 {
2769 if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2770 {
2771 profile_record=device->profile_records[i];
2772 break;
2773 }
2774 i++;
2775 }
2776 }
2777 if (profile_record != (KernelProfileRecord) NULL)
2778 name=DestroyString(name);
2779 else
2780 {
2781 profile_record=(KernelProfileRecord) AcquireCriticalMemory(
2782 sizeof(*profile_record));
2783 (void) memset(profile_record,0,sizeof(*profile_record));
2784 profile_record->kernel_name=name;
2785 device->profile_records=(KernelProfileRecord *) ResizeQuantumMemory(
2786 device->profile_records,(i+2),sizeof(*device->profile_records));
2787 if (device->profile_records == (KernelProfileRecord *) NULL)
2788 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
2789 device->profile_records[i]=profile_record;
2790 device->profile_records[i+1]=(KernelProfileRecord) NULL;
2791 }
2792 if ((elapsed < profile_record->min) || (profile_record->count == 0))
2793 profile_record->min=(unsigned long) elapsed;
2794 if (elapsed > profile_record->max)
2795 profile_record->max=(unsigned long) elapsed;
2796 profile_record->total+=(unsigned long) elapsed;
2797 profile_record->count+=1;
2798 UnlockSemaphoreInfo(device->lock);
2799 return(MagickTrue);
2800}
2801
2802/*
2803%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2804% %
2805% %
2806% %
2807+ 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 %
2808% %
2809% %
2810% %
2811%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2812%
2813% ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2814%
2815% The format of the ReleaseOpenCLCommandQueue method is:
2816%
2817% void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2818% cl_command_queue queue)
2819%
2820% A description of each parameter follows:
2821%
2822% o device: the OpenCL device.
2823%
2824% o queue: the OpenCL queue to be released.
2825*/
2826
2827MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2828 cl_command_queue queue)
2829{
2830 if (queue == (cl_command_queue) NULL)
2831 return;
2832
2833 assert(device != (MagickCLDevice) NULL);
2834 LockSemaphoreInfo(device->lock);
2835 if ((device->profile_kernels != MagickFalse) ||
2836 (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2837 {
2838 UnlockSemaphoreInfo(device->lock);
2839 openCL_library->clFinish(queue);
2840 (void) openCL_library->clReleaseCommandQueue(queue);
2841 }
2842 else
2843 {
2844 openCL_library->clFlush(queue);
2845 device->command_queues[++device->command_queues_index]=queue;
2846 UnlockSemaphoreInfo(device->lock);
2847 }
2848}
2849
2850/*
2851%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2852% %
2853% %
2854% %
2855+ R e l e a s e M a g i c k C L D e v i c e %
2856% %
2857% %
2858% %
2859%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2860%
2861% ReleaseOpenCLDevice() returns the OpenCL device to the environment
2862%
2863% The format of the ReleaseOpenCLDevice method is:
2864%
2865% void ReleaseOpenCLDevice(MagickCLDevice device)
2866%
2867% A description of each parameter follows:
2868%
2869% o device: the OpenCL device to be released.
2870%
2871*/
2872
2873MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2874{
2875 assert(device != (MagickCLDevice) NULL);
2876 LockSemaphoreInfo(openCL_lock);
2877 device->requested--;
2878 UnlockSemaphoreInfo(openCL_lock);
2879}
2880
2881/*
2882%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2883% %
2884% %
2885% %
2886+ 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 %
2887% %
2888% %
2889% %
2890%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2891%
2892% RelinquishMagickCLCacheInfo() frees memory acquired with
2893% AcquireMagickCLCacheInfo()
2894%
2895% The format of the RelinquishMagickCLCacheInfo method is:
2896%
2897% MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2898% const MagickBooleanType relinquish_pixels)
2899%
2900% A description of each parameter follows:
2901%
2902% o info: the OpenCL cache info.
2903%
2904% o relinquish_pixels: the pixels will be relinquish when set to true.
2905%
2906*/
2907
2908static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2909 cl_event magick_unused(event),
2910 cl_int magick_unused(event_command_exec_status),void *user_data)
2911{
2912 MagickCLCacheInfo
2913 info;
2914
2915 Quantum
2916 *pixels;
2917
2918 ssize_t
2919 i;
2920
2921 magick_unreferenced(event);
2922 magick_unreferenced(event_command_exec_status);
2923 info=(MagickCLCacheInfo) user_data;
2924 for (i=(ssize_t)info->event_count-1; i >= 0; i--)
2925 {
2926 cl_int
2927 event_status;
2928
2929 cl_uint
2930 status;
2931
2932 status=openCL_library->clGetEventInfo(info->events[i],
2933 CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
2934 NULL);
2935 if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
2936 {
2937 openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
2938 &DestroyMagickCLCacheInfoAndPixels,info);
2939 return;
2940 }
2941 }
2942 pixels=info->pixels;
2943 RelinquishMagickResource(MemoryResource,info->length);
2944 DestroyMagickCLCacheInfo(info);
2945 (void) RelinquishAlignedMemory(pixels);
2946}
2947
2948MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2949 MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2950{
2951 if (info == (MagickCLCacheInfo) NULL)
2952 return((MagickCLCacheInfo) NULL);
2953 if (relinquish_pixels != MagickFalse)
2954 DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2955 else
2956 DestroyMagickCLCacheInfo(info);
2957 return((MagickCLCacheInfo) NULL);
2958}
2959
2960/*
2961%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2962% %
2963% %
2964% %
2965% R e l i n q u i s h M a g i c k C L D e v i c e %
2966% %
2967% %
2968% %
2969%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2970%
2971% RelinquishMagickCLDevice() releases the OpenCL device
2972%
2973% The format of the RelinquishMagickCLDevice method is:
2974%
2975% MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2976%
2977% A description of each parameter follows:
2978%
2979% o device: the OpenCL device to be released.
2980%
2981*/
2982
2983static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2984{
2985 if (device == (MagickCLDevice) NULL)
2986 return((MagickCLDevice) NULL);
2987
2988 device->platform_name=(char *) RelinquishMagickMemory(device->platform_name);
2989 device->vendor_name=(char *) RelinquishMagickMemory(device->vendor_name);
2990 device->name=(char *) RelinquishMagickMemory(device->name);
2991 device->version=(char *) RelinquishMagickMemory(device->version);
2992 if (device->program != (cl_program) NULL)
2993 (void) openCL_library->clReleaseProgram(device->program);
2994 while (device->command_queues_index >= 0)
2995 (void) openCL_library->clReleaseCommandQueue(
2996 device->command_queues[device->command_queues_index--]);
2997 RelinquishSemaphoreInfo(&device->lock);
2998 return((MagickCLDevice) RelinquishMagickMemory(device));
2999}
3000
3001/*
3002%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3003% %
3004% %
3005% %
3006% R e l i n q u i s h M a g i c k C L E n v %
3007% %
3008% %
3009% %
3010%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3011%
3012% RelinquishMagickCLEnv() releases the OpenCL environment
3013%
3014% The format of the RelinquishMagickCLEnv method is:
3015%
3016% MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
3017%
3018% A description of each parameter follows:
3019%
3020% o clEnv: the OpenCL environment to be released.
3021%
3022*/
3023
3024static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
3025{
3026 if (clEnv == (MagickCLEnv) NULL)
3027 return((MagickCLEnv) NULL);
3028
3029 RelinquishSemaphoreInfo(&clEnv->lock);
3030 RelinquishMagickCLDevices(clEnv);
3031 if (clEnv->contexts != (cl_context *) NULL)
3032 {
3033 ssize_t
3034 i;
3035
3036 for (i=0; i < (ssize_t) clEnv->number_contexts; i++)
3037 if (clEnv->contexts[i] != (cl_context) NULL)
3038 (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
3039 clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
3040 }
3041 return((MagickCLEnv) RelinquishMagickMemory(clEnv));
3042}
3043
3044/*
3045%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3046% %
3047% %
3048% %
3049+ R e q u e s t O p e n C L D e v i c e %
3050% %
3051% %
3052% %
3053%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3054%
3055% RequestOpenCLDevice() returns one of the enabled OpenCL devices.
3056%
3057% The format of the RequestOpenCLDevice method is:
3058%
3059% MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3060%
3061% A description of each parameter follows:
3062%
3063% o clEnv: the OpenCL environment.
3064*/
3065
3066MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3067{
3068 MagickCLDevice
3069 device;
3070
3071 double
3072 score,
3073 best_score;
3074
3075 size_t
3076 i;
3077
3078 if (clEnv == (MagickCLEnv) NULL)
3079 return((MagickCLDevice) NULL);
3080
3081 if (clEnv->number_devices == 1)
3082 {
3083 if (clEnv->devices[0]->enabled)
3084 return(clEnv->devices[0]);
3085 else
3086 return((MagickCLDevice) NULL);
3087 }
3088
3089 device=(MagickCLDevice) NULL;
3090 best_score=0.0;
3091 LockSemaphoreInfo(openCL_lock);
3092 for (i = 0; i < clEnv->number_devices; i++)
3093 {
3094 if (clEnv->devices[i]->enabled == MagickFalse)
3095 continue;
3096
3097 score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3098 clEnv->devices[i]->requested);
3099 if ((device == (MagickCLDevice) NULL) || (score < best_score))
3100 {
3101 device=clEnv->devices[i];
3102 best_score=score;
3103 }
3104 }
3105 if (device != (MagickCLDevice)NULL)
3106 device->requested++;
3107 UnlockSemaphoreInfo(openCL_lock);
3108
3109 return(device);
3110}
3111
3112/*
3113%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3114% %
3115% %
3116% %
3117% S e t O p e n C L D e v i c e E n a b l e d %
3118% %
3119% %
3120% %
3121%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3122%
3123% SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
3124%
3125% The format of the SetOpenCLDeviceEnabled method is:
3126%
3127% void SetOpenCLDeviceEnabled(MagickCLDevice device,
3128% MagickBooleanType value)
3129%
3130% A description of each parameter follows:
3131%
3132% o device: the OpenCL device.
3133%
3134% o value: determines if the device should be enabled or disabled.
3135*/
3136
3137MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
3138 const MagickBooleanType value)
3139{
3140 if (device == (MagickCLDevice) NULL)
3141 return;
3142 device->enabled=value;
3143}
3144
3145/*
3146%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3147% %
3148% %
3149% %
3150% 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 %
3151% %
3152% %
3153% %
3154%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3155%
3156% SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
3157% kernel profiling of a device.
3158%
3159% The format of the SetOpenCLKernelProfileEnabled method is:
3160%
3161% void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3162% MagickBooleanType value)
3163%
3164% A description of each parameter follows:
3165%
3166% o device: the OpenCL device.
3167%
3168% o value: determines if kernel profiling for the device should be enabled
3169% or disabled.
3170*/
3171
3172MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3173 const MagickBooleanType value)
3174{
3175 if (device == (MagickCLDevice) NULL)
3176 return;
3177 device->profile_kernels=value;
3178}
3179
3180/*
3181%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3182% %
3183% %
3184% %
3185% S e t O p e n C L E n a b l e d %
3186% %
3187% %
3188% %
3189%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3190%
3191% SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3192%
3193% The format of the SetOpenCLEnabled method is:
3194%
3195% void SetOpenCLEnabled(MagickBooleanType)
3196%
3197% A description of each parameter follows:
3198%
3199% o value: specify true to enable OpenCL acceleration
3200*/
3201
3202MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
3203{
3204 MagickCLEnv
3205 clEnv;
3206
3207 clEnv=GetCurrentOpenCLEnv();
3208 if (clEnv == (MagickCLEnv) NULL)
3209 return(MagickFalse);
3210 clEnv->enabled=value;
3211 return(clEnv->enabled);
3212}
3213
3214#else
3215
3216MagickExport double GetOpenCLDeviceBenchmarkScore(
3217 const MagickCLDevice magick_unused(device))
3218{
3219 magick_unreferenced(device);
3220 return(0.0);
3221}
3222
3223MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3224 const MagickCLDevice magick_unused(device))
3225{
3226 magick_unreferenced(device);
3227 return(MagickFalse);
3228}
3229
3230MagickExport const char *GetOpenCLDeviceName(
3231 const MagickCLDevice magick_unused(device))
3232{
3233 magick_unreferenced(device);
3234 return((const char *) NULL);
3235}
3236
3237MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
3238 ExceptionInfo *magick_unused(exception))
3239{
3240 magick_unreferenced(exception);
3241 if (length != (size_t *) NULL)
3242 *length=0;
3243 return((MagickCLDevice *) NULL);
3244}
3245
3246MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3247 const MagickCLDevice magick_unused(device))
3248{
3249 magick_unreferenced(device);
3250 return(UndefinedCLDeviceType);
3251}
3252
3253MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
3254 const MagickCLDevice magick_unused(device),size_t *length)
3255{
3256 magick_unreferenced(device);
3257 if (length != (size_t *) NULL)
3258 *length=0;
3259 return((const KernelProfileRecord *) NULL);
3260}
3261
3262MagickExport const char *GetOpenCLDeviceVersion(
3263 const MagickCLDevice magick_unused(device))
3264{
3265 magick_unreferenced(device);
3266 return((const char *) NULL);
3267}
3268
3269MagickExport MagickBooleanType GetOpenCLEnabled(void)
3270{
3271 return(MagickFalse);
3272}
3273
3274MagickExport void SetOpenCLDeviceEnabled(
3275 MagickCLDevice magick_unused(device),
3276 const MagickBooleanType magick_unused(value))
3277{
3278 magick_unreferenced(device);
3279 magick_unreferenced(value);
3280}
3281
3282MagickExport MagickBooleanType SetOpenCLEnabled(
3283 const MagickBooleanType magick_unused(value))
3284{
3285 magick_unreferenced(value);
3286 return(MagickFalse);
3287}
3288
3289MagickExport void SetOpenCLKernelProfileEnabled(
3290 MagickCLDevice magick_unused(device),
3291 const MagickBooleanType magick_unused(value))
3292{
3293 magick_unreferenced(device);
3294 magick_unreferenced(value);
3295}
3296#endif