MagickCore 7.1.2-13
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 device_benchmark->platform_name=(char *) RelinquishMagickMemory(
758 device_benchmark->platform_name);
759 device_benchmark->vendor_name=(char *) RelinquishMagickMemory(
760 device_benchmark->vendor_name);
761 device_benchmark->name=(char *) RelinquishMagickMemory(
762 device_benchmark->name);
763 device_benchmark->version=(char *) RelinquishMagickMemory(
764 device_benchmark->version);
765 return((MagickCLDeviceBenchmark *) RelinquishMagickMemory(
766 device_benchmark));
767}
768
769static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
770{
771 char
772 keyword[MagickPathExtent],
773 *token;
774
775 const char
776 *q;
777
778 MagickCLDeviceBenchmark
779 *device_benchmark;
780
781 size_t
782 i,
783 extent;
784
785 if (xml == (char *) NULL)
786 return;
787 device_benchmark=(MagickCLDeviceBenchmark *) NULL;
788 token=AcquireString(xml);
789 extent=strlen(token)+MagickPathExtent;
790 for (q=(char *) xml; *q != '\0'; )
791 {
792 /*
793 Interpret XML.
794 */
795 (void) GetNextToken(q,&q,extent,token);
796 if (*token == '\0')
797 break;
798 (void) CopyMagickString(keyword,token,MagickPathExtent);
799 if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
800 {
801 /*
802 Doctype element.
803 */
804 while ((LocaleNCompare(q,"]>",2) != 0) && (*q != '\0'))
805 (void) GetNextToken(q,&q,extent,token);
806 continue;
807 }
808 if (LocaleNCompare(keyword,"<!--",4) == 0)
809 {
810 /*
811 Comment element.
812 */
813 while ((LocaleNCompare(q,"->",2) != 0) && (*q != '\0'))
814 (void) GetNextToken(q,&q,extent,token);
815 continue;
816 }
817 if (LocaleCompare(keyword,"<device") == 0)
818 {
819 /*
820 Device element.
821 */
822 device_benchmark=(MagickCLDeviceBenchmark *) AcquireQuantumMemory(1,
823 sizeof(*device_benchmark));
824 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
825 break;
826 (void) memset(device_benchmark,0,sizeof(*device_benchmark));
827 device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
828 continue;
829 }
830 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
831 continue;
832 if (LocaleCompare(keyword,"/>") == 0)
833 {
834 if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
835 {
836 if (LocaleCompare(device_benchmark->name,"CPU") == 0)
837 clEnv->cpu_score=device_benchmark->score;
838 else
839 {
840 MagickCLDevice
841 device;
842
843 /*
844 Set the score for all devices that match this device.
845 */
846 for (i = 0; i < clEnv->number_devices; i++)
847 {
848 device=clEnv->devices[i];
849 if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
850 device->score=device_benchmark->score;
851 }
852 }
853 }
854 device_benchmark=RelinquishDeviceBenchmark(device_benchmark);
855 continue;
856 }
857 (void) GetNextToken(q,(const char **) NULL,extent,token);
858 if (*token != '=')
859 continue;
860 (void) GetNextToken(q,&q,extent,token);
861 (void) GetNextToken(q,&q,extent,token);
862 switch (*keyword)
863 {
864 case 'M':
865 case 'm':
866 {
867 if (LocaleCompare((char *) keyword,"maxClockFrequency") == 0)
868 {
869 device_benchmark->max_clock_frequency=StringToInteger(token);
870 break;
871 }
872 if (LocaleCompare((char *) keyword,"maxComputeUnits") == 0)
873 {
874 device_benchmark->max_compute_units=StringToInteger(token);
875 break;
876 }
877 break;
878 }
879 case 'N':
880 case 'n':
881 {
882 if (LocaleCompare((char *) keyword,"name") == 0)
883 device_benchmark->name=ConstantString(token);
884 break;
885 }
886 case 'P':
887 case 'p':
888 {
889 if (LocaleCompare((char *) keyword,"platform") == 0)
890 device_benchmark->platform_name=ConstantString(token);
891 break;
892 }
893 case 'S':
894 case 's':
895 {
896 if (LocaleCompare((char *) keyword,"score") == 0)
897 device_benchmark->score=StringToDouble(token,(char **) NULL);
898 break;
899 }
900 case 'V':
901 case 'v':
902 {
903 if (LocaleCompare((char *) keyword,"vendor") == 0)
904 device_benchmark->vendor_name=ConstantString(token);
905 if (LocaleCompare((char *) keyword,"version") == 0)
906 device_benchmark->version=ConstantString(token);
907 break;
908 }
909 default:
910 break;
911 }
912 }
913 token=(char *) RelinquishMagickMemory(token);
914 device_benchmark=RelinquishDeviceBenchmark(device_benchmark);
915}
916
917static MagickBooleanType CanWriteProfileToFile(const char *filename)
918{
919 FILE
920 *profileFile;
921
922 profileFile=fopen_utf8(filename,"ab");
923
924 if (profileFile == (FILE *) NULL)
925 {
926 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
927 "Unable to save profile to: \"%s\"",filename);
928 return(MagickFalse);
929 }
930
931 fclose(profileFile);
932 return(MagickTrue);
933}
934#endif
935
936static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
937{
938#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
939 char
940 filename[MagickPathExtent];
941
942 StringInfo
943 *option;
944
945 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
946 GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
947
948 /*
949 We don't run the benchmark when we can not write out a device profile. The
950 first GPU device will be used.
951 */
952 if (CanWriteProfileToFile(filename) == MagickFalse)
953#endif
954 {
955 size_t
956 i;
957
958 for (i = 0; i < clEnv->number_devices; i++)
959 clEnv->devices[i]->score=1.0;
960
961 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
962 return(MagickFalse);
963 }
964#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
965 option=ConfigureFileToStringInfo(filename);
966 LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(option));
967 option=DestroyStringInfo(option);
968 return(MagickTrue);
969#endif
970}
971
972static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
973{
974 char
975 *option;
976
977 double
978 best_score;
979
980 MagickBooleanType
981 benchmark;
982
983 size_t
984 i;
985
986 option=GetEnvironmentValue("MAGICK_OCL_DEVICE");
987 if (option != (const char *) NULL)
988 {
989 if (strcmp(option,"GPU") == 0)
990 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
991 else if (strcmp(option,"CPU") == 0)
992 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
993 option=DestroyString(option);
994 }
995
996 if (LoadOpenCLBenchmarks(clEnv) == MagickFalse)
997 return;
998
999 benchmark=MagickFalse;
1000 if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1001 benchmark=MagickTrue;
1002 else
1003 {
1004 for (i = 0; i < clEnv->number_devices; i++)
1005 {
1006 if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1007 {
1008 benchmark=MagickTrue;
1009 break;
1010 }
1011 }
1012 }
1013
1014 if (benchmark != MagickFalse)
1015 BenchmarkOpenCLDevices(clEnv);
1016
1017 best_score=clEnv->cpu_score;
1018 for (i = 0; i < clEnv->number_devices; i++)
1019 best_score=MagickMin(clEnv->devices[i]->score,best_score);
1020
1021 for (i = 0; i < clEnv->number_devices; i++)
1022 {
1023 if (clEnv->devices[i]->score != best_score)
1024 clEnv->devices[i]->enabled=MagickFalse;
1025 }
1026}
1027
1028/*
1029%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1030% %
1031% %
1032% %
1033% B e n c h m a r k O p e n C L D e v i c e s %
1034% %
1035% %
1036% %
1037%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1038%
1039% BenchmarkOpenCLDevices() benchmarks the OpenCL devices and the CPU to help
1040% the automatic selection of the best device.
1041%
1042% The format of the BenchmarkOpenCLDevices method is:
1043%
1044% void BenchmarkOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
1045%
1046% A description of each parameter follows:
1047%
1048% o clEnv: the OpenCL environment.
1049%
1050% o exception: return any errors or warnings
1051*/
1052
1053static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
1054{
1055 AccelerateTimer
1056 timer;
1057
1058 ExceptionInfo
1059 *exception;
1060
1061 Image
1062 *inputImage;
1063
1064 ImageInfo
1065 *imageInfo;
1066
1067 size_t
1068 i;
1069
1070 exception=AcquireExceptionInfo();
1071 imageInfo=AcquireImageInfo();
1072 CloneString(&imageInfo->size,"2048x1536");
1073 (void) CopyMagickString(imageInfo->filename,"xc:none",MagickPathExtent);
1074 inputImage=ReadImage(imageInfo,exception);
1075 if (inputImage == (Image *) NULL)
1076 return(0.0);
1077
1078 InitAccelerateTimer(&timer);
1079
1080 for (i=0; i<=2; i++)
1081 {
1082 Image
1083 *blurredImage,
1084 *resizedImage,
1085 *unsharpedImage;
1086
1087 if (i > 0)
1088 StartAccelerateTimer(&timer);
1089
1090 blurredImage=BlurImage(inputImage,10.0f,3.5f,exception);
1091 unsharpedImage=UnsharpMaskImage(blurredImage,2.0f,2.0f,50.0f,10.0f,
1092 exception);
1093 resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
1094 exception);
1095
1096 /*
1097 We need this to get a proper performance benchmark, the operations
1098 are executed asynchronous.
1099 */
1100 if (is_cpu == MagickFalse)
1101 {
1102 CacheInfo
1103 *cache_info;
1104
1105 cache_info=(CacheInfo *) resizedImage->cache;
1106 if (cache_info->opencl != (MagickCLCacheInfo) NULL)
1107 openCL_library->clWaitForEvents(cache_info->opencl->event_count,
1108 cache_info->opencl->events);
1109 }
1110
1111 if (i > 0)
1112 StopAccelerateTimer(&timer);
1113
1114 if (blurredImage != (Image *) NULL)
1115 DestroyImage(blurredImage);
1116 if (unsharpedImage != (Image *) NULL)
1117 DestroyImage(unsharpedImage);
1118 if (resizedImage != (Image *) NULL)
1119 DestroyImage(resizedImage);
1120 }
1121 DestroyImage(inputImage);
1122 return(ReadAccelerateTimer(&timer));
1123}
1124
1125static void RunDeviceBenchmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
1126 MagickCLDevice device)
1127{
1128 testEnv->devices[0]=device;
1129 default_CLEnv=testEnv;
1130 device->score=RunOpenCLBenchmark(MagickFalse);
1131 default_CLEnv=clEnv;
1132 testEnv->devices[0]=(MagickCLDevice) NULL;
1133}
1134
1135static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
1136{
1137 char
1138 filename[MagickPathExtent];
1139
1140 FILE
1141 *cache_file;
1142
1143 MagickCLDevice
1144 device;
1145
1146 size_t
1147 i,
1148 j;
1149
1150 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1151 GetOpenCLCacheDirectory(),DirectorySeparator,
1152 IMAGEMAGICK_PROFILE_FILE);
1153
1154 cache_file=fopen_utf8(filename,"wb");
1155 if (cache_file == (FILE *) NULL)
1156 return;
1157 fwrite("<devices>\n",sizeof(char),10,cache_file);
1158 fprintf(cache_file," <device name=\"CPU\" score=\"%.4g\"/>\n",
1159 clEnv->cpu_score);
1160 for (i = 0; i < clEnv->number_devices; i++)
1161 {
1162 MagickBooleanType
1163 duplicate;
1164
1165 device=clEnv->devices[i];
1166 duplicate=MagickFalse;
1167 for (j = 0; j < i; j++)
1168 {
1169 if (IsSameOpenCLDevice(clEnv->devices[j],device))
1170 {
1171 duplicate=MagickTrue;
1172 break;
1173 }
1174 }
1175
1176 if (duplicate)
1177 continue;
1178
1179 if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1180 fprintf(cache_file," <device platform=\"%s\" vendor=\"%s\" name=\"%s\"\
1181 version=\"%s\" maxClockFrequency=\"%d\" maxComputeUnits=\"%d\"\
1182 score=\"%.4g\"/>\n",
1183 device->platform_name,device->vendor_name,device->name,device->version,
1184 (int)device->max_clock_frequency,(int)device->max_compute_units,
1185 device->score);
1186 }
1187 fwrite("</devices>",sizeof(char),10,cache_file);
1188
1189 fclose(cache_file);
1190}
1191
1192static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
1193{
1194 MagickCLDevice
1195 device;
1196
1197 MagickCLEnv
1198 testEnv;
1199
1200 size_t
1201 i,
1202 j;
1203
1204 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
1205 "Starting benchmark");
1206 testEnv=AcquireMagickCLEnv();
1207 testEnv->library=openCL_library;
1208 testEnv->devices=(MagickCLDevice *) AcquireCriticalMemory(
1209 sizeof(MagickCLDevice));
1210 testEnv->number_devices=1;
1211 testEnv->benchmark_thread_id=GetMagickThreadId();
1212 testEnv->initialized=MagickTrue;
1213
1214 for (i = 0; i < clEnv->number_devices; i++)
1215 clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
1216
1217 for (i = 0; i < clEnv->number_devices; i++)
1218 {
1219 device=clEnv->devices[i];
1220 if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1221 RunDeviceBenchmark(clEnv,testEnv,device);
1222
1223 /* Set the score on all the other devices that are the same */
1224 for (j = i+1; j < clEnv->number_devices; j++)
1225 {
1226 MagickCLDevice
1227 other_device;
1228
1229 other_device=clEnv->devices[j];
1230 if (IsSameOpenCLDevice(device,other_device))
1231 other_device->score=device->score;
1232 }
1233 }
1234
1235 testEnv->enabled=MagickFalse;
1236 default_CLEnv=testEnv;
1237 clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
1238 default_CLEnv=clEnv;
1239
1240 testEnv=RelinquishMagickCLEnv(testEnv);
1241 CacheOpenCLBenchmarks(clEnv);
1242}
1243
1244/*
1245%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1246% %
1247% %
1248% %
1249% C o m p i l e O p e n C L K e r n e l %
1250% %
1251% %
1252% %
1253%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1254%
1255% CompileOpenCLKernel() compiles the kernel for the specified device. The
1256% kernel will be cached on disk to reduce the compilation time.
1257%
1258% The format of the CompileOpenCLKernel method is:
1259%
1260% MagickBooleanType AcquireOpenCLKernel(MagickCLDevice clEnv,
1261% unsigned int signature,const char *kernel,const char *options,
1262% ExceptionInfo *exception)
1263%
1264% A description of each parameter follows:
1265%
1266% o device: the OpenCL device.
1267%
1268% o kernel: the source code of the kernel.
1269%
1270% o options: options for the compiler.
1271%
1272% o signature: a number to uniquely identify the kernel
1273%
1274% o exception: return any errors or warnings in this structure.
1275%
1276*/
1277
1278static void CacheOpenCLKernel(MagickCLDevice device,char *filename,
1279 ExceptionInfo *exception)
1280{
1281 cl_uint
1282 status;
1283
1284 size_t
1285 binaryProgramSize;
1286
1287 unsigned char
1288 *binaryProgram;
1289
1290 status=openCL_library->clGetProgramInfo(device->program,
1291 CL_PROGRAM_BINARY_SIZES,sizeof(size_t),&binaryProgramSize,NULL);
1292 if (status != CL_SUCCESS)
1293 return;
1294 binaryProgram=(unsigned char*) AcquireQuantumMemory(1,binaryProgramSize);
1295 if (binaryProgram == (unsigned char *) NULL)
1296 {
1297 (void) ThrowMagickException(exception,GetMagickModule(),
1298 ResourceLimitError,"MemoryAllocationFailed","`%s'",filename);
1299 return;
1300 }
1301 status=openCL_library->clGetProgramInfo(device->program,
1302 CL_PROGRAM_BINARIES,sizeof(unsigned char*),&binaryProgram,NULL);
1303 if (status == CL_SUCCESS)
1304 {
1305 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
1306 "Creating cache file: \"%s\"",filename);
1307 (void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
1308 }
1309 binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1310}
1311
1312static MagickBooleanType LoadCachedOpenCLKernels(MagickCLDevice device,
1313 const char *filename)
1314{
1315 cl_int
1316 binaryStatus,
1317 status;
1318
1319 ExceptionInfo
1320 *sans_exception;
1321
1322 size_t
1323 length;
1324
1325 unsigned char
1326 *binaryProgram;
1327
1328 sans_exception=AcquireExceptionInfo();
1329 binaryProgram=(unsigned char *) FileToBlob(filename,SIZE_MAX,&length,
1330 sans_exception);
1331 sans_exception=DestroyExceptionInfo(sans_exception);
1332 if (binaryProgram == (unsigned char *) NULL)
1333 return(MagickFalse);
1334 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
1335 "Loaded cached kernels: \"%s\"",filename);
1336 device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
1337 &device->deviceID,&length,(const unsigned char**)&binaryProgram,
1338 &binaryStatus,&status);
1339 binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1340 return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
1341 MagickTrue);
1342}
1343
1344static void LogOpenCLBuildFailure(MagickCLDevice device,const char *kernel,
1345 ExceptionInfo *exception)
1346{
1347 char
1348 filename[MagickPathExtent],
1349 *log;
1350
1351 size_t
1352 log_size;
1353
1354 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1355 GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.cl");
1356
1357 (void) remove_utf8(filename);
1358 (void) BlobToFile(filename,kernel,strlen(kernel),exception);
1359
1360 openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1361 CL_PROGRAM_BUILD_LOG,0,NULL,&log_size);
1362 log=(char*)AcquireCriticalMemory(log_size);
1363 openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1364 CL_PROGRAM_BUILD_LOG,log_size,log,&log_size);
1365
1366 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1367 GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.log");
1368
1369 (void) remove_utf8(filename);
1370 (void) BlobToFile(filename,log,log_size,exception);
1371 log=(char*)RelinquishMagickMemory(log);
1372}
1373
1374static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
1375 const char *kernel,const char *options,size_t signature,
1376 ExceptionInfo *exception)
1377{
1378 char
1379 deviceName[MagickPathExtent],
1380 filename[MagickPathExtent],
1381 *ptr;
1382
1383 cl_int
1384 status;
1385
1386 MagickBooleanType
1387 loaded;
1388
1389 size_t
1390 length;
1391
1392 (void) CopyMagickString(deviceName,device->name,MagickPathExtent);
1393 ptr=deviceName;
1394 /* Strip out illegal characters for file names */
1395 while (*ptr != '\0')
1396 {
1397 if ((*ptr == ' ') || (*ptr == '\\') || (*ptr == '/') || (*ptr == ':') ||
1398 (*ptr == '*') || (*ptr == '?') || (*ptr == '"') || (*ptr == '<') ||
1399 (*ptr == '>' || *ptr == '|'))
1400 *ptr = '_';
1401 ptr++;
1402 }
1403 (void) FormatLocaleString(filename,MagickPathExtent,
1404 "%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
1405 DirectorySeparator,"magick_opencl",deviceName,(unsigned int) signature,
1406 (double) sizeof(char*)*8);
1407 loaded=LoadCachedOpenCLKernels(device,filename);
1408 if (loaded == MagickFalse)
1409 {
1410 /* Binary CL program unavailable, compile the program from source */
1411 length=strlen(kernel);
1412 device->program=openCL_library->clCreateProgramWithSource(
1413 device->context,1,&kernel,&length,&status);
1414 if (status != CL_SUCCESS)
1415 return(MagickFalse);
1416 }
1417
1418 status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
1419 options,NULL,NULL);
1420 if (status != CL_SUCCESS)
1421 {
1422 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
1423 "clBuildProgram failed.","(%d)",(int)status);
1424 LogOpenCLBuildFailure(device,kernel,exception);
1425 return(MagickFalse);
1426 }
1427
1428 /* Save the binary to a file to avoid re-compilation of the kernels */
1429 if (loaded == MagickFalse)
1430 CacheOpenCLKernel(device,filename,exception);
1431
1432 return(MagickTrue);
1433}
1434
1435static cl_event* CopyOpenCLEvents(MagickCLCacheInfo first,
1436 MagickCLCacheInfo second,cl_uint *event_count)
1437{
1438 cl_event
1439 *events;
1440
1441 size_t
1442 i;
1443
1444 size_t
1445 j;
1446
1447 assert(first != (MagickCLCacheInfo) NULL);
1448 assert(event_count != (cl_uint *) NULL);
1449 events=(cl_event *) NULL;
1450 LockSemaphoreInfo(first->events_semaphore);
1451 if (second != (MagickCLCacheInfo) NULL)
1452 LockSemaphoreInfo(second->events_semaphore);
1453 *event_count=first->event_count;
1454 if (second != (MagickCLCacheInfo) NULL)
1455 *event_count+=second->event_count;
1456 if (*event_count > 0)
1457 {
1458 events=(cl_event *) AcquireQuantumMemory(*event_count,sizeof(*events));
1459 if (events == (cl_event *) NULL)
1460 *event_count=0;
1461 else
1462 {
1463 j=0;
1464 for (i=0; i < first->event_count; i++, j++)
1465 events[j]=first->events[i];
1466 if (second != (MagickCLCacheInfo) NULL)
1467 {
1468 for (i=0; i < second->event_count; i++, j++)
1469 events[j]=second->events[i];
1470 }
1471 }
1472 }
1473 UnlockSemaphoreInfo(first->events_semaphore);
1474 if (second != (MagickCLCacheInfo) NULL)
1475 UnlockSemaphoreInfo(second->events_semaphore);
1476 return(events);
1477}
1478
1479/*
1480%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1481% %
1482% %
1483% %
1484+ C o p y M a g i c k C L C a c h e I n f o %
1485% %
1486% %
1487% %
1488%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1489%
1490% CopyMagickCLCacheInfo() copies the memory from the device into host memory.
1491%
1492% The format of the CopyMagickCLCacheInfo method is:
1493%
1494% void CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1495%
1496% A description of each parameter follows:
1497%
1498% o info: the OpenCL cache info.
1499%
1500*/
1501MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1502{
1503 cl_command_queue
1504 queue;
1505
1506 cl_event
1507 *events;
1508
1509 cl_uint
1510 event_count;
1511
1512 Quantum
1513 *pixels;
1514
1515 if (info == (MagickCLCacheInfo) NULL)
1516 return((MagickCLCacheInfo) NULL);
1517 events=CopyOpenCLEvents(info,(MagickCLCacheInfo) NULL,&event_count);
1518 if (events != (cl_event *) NULL)
1519 {
1520 queue=AcquireOpenCLCommandQueue(info->device);
1521 pixels=(Quantum *) openCL_library->clEnqueueMapBuffer(queue,info->buffer,
1522 CL_TRUE,CL_MAP_READ | CL_MAP_WRITE,0,(size_t) info->length,event_count,
1523 events,
1524 (cl_event *) NULL,(cl_int *) NULL);
1525 assert(pixels == info->pixels);
1526 ReleaseOpenCLCommandQueue(info->device,queue);
1527 events=(cl_event *) RelinquishMagickMemory(events);
1528 }
1529 return(RelinquishMagickCLCacheInfo(info,MagickFalse));
1530}
1531
1532/*
1533%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1534% %
1535% %
1536% %
1537+ D u m p O p e n C L P r o f i l e D a t a %
1538% %
1539% %
1540% %
1541%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1542%
1543% DumpOpenCLProfileData() dumps the kernel profile data.
1544%
1545% The format of the DumpProfileData method is:
1546%
1547% void DumpProfileData()
1548%
1549*/
1550
1551MagickPrivate void DumpOpenCLProfileData()
1552{
1553#define OpenCLLog(message) \
1554 fwrite(message,sizeof(char),strlen(message),log); \
1555 fwrite("\n",sizeof(char),1,log);
1556
1557 char
1558 buf[4096],
1559 filename[MagickPathExtent],
1560 indent[160];
1561
1562 FILE
1563 *log;
1564
1565 size_t
1566 i,
1567 j;
1568
1569 if (default_CLEnv == (MagickCLEnv) NULL)
1570 return;
1571
1572 for (i = 0; i < default_CLEnv->number_devices; i++)
1573 if (default_CLEnv->devices[i]->profile_kernels != MagickFalse)
1574 break;
1575 if (i == default_CLEnv->number_devices)
1576 return;
1577
1578 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1579 GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
1580
1581 log=fopen_utf8(filename,"wb");
1582 if (log == (FILE *) NULL)
1583 return;
1584 for (i = 0; i < default_CLEnv->number_devices; i++)
1585 {
1586 MagickCLDevice
1587 device;
1588
1589 device=default_CLEnv->devices[i];
1590 if ((device->profile_kernels == MagickFalse) ||
1591 (device->profile_records == (KernelProfileRecord *) NULL))
1592 continue;
1593
1594 OpenCLLog("====================================================");
1595 fprintf(log,"Device: %s\n",device->name);
1596 fprintf(log,"Version: %s\n",device->version);
1597 OpenCLLog("====================================================");
1598 OpenCLLog(" average calls min max");
1599 OpenCLLog(" ------- ----- --- ---");
1600 j=0;
1601 while (device->profile_records[j] != (KernelProfileRecord) NULL)
1602 {
1603 KernelProfileRecord
1604 profile;
1605
1606 profile=device->profile_records[j];
1607 (void) CopyMagickString(indent," ",
1608 sizeof(indent));
1609 (void) CopyMagickString(indent,profile->kernel_name,MagickMin(strlen(
1610 profile->kernel_name),strlen(indent)));
1611 (void) FormatLocaleString(buf,sizeof(buf),"%s %7d %7d %7d %7d",indent,
1612 (int) (profile->total/profile->count),(int) profile->count,
1613 (int) profile->min,(int) profile->max);
1614 OpenCLLog(buf);
1615 j++;
1616 }
1617 OpenCLLog("====================================================");
1618 fwrite("\n\n",sizeof(char),2,log);
1619 }
1620 fclose(log);
1621}
1622/*
1623%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1624% %
1625% %
1626% %
1627+ E n q u e u e O p e n C L K e r n e l %
1628% %
1629% %
1630% %
1631%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1632%
1633% EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
1634% events with the images.
1635%
1636% The format of the EnqueueOpenCLKernel method is:
1637%
1638% MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
1639% const size_t *global_work_offset,const size_t *global_work_size,
1640% const size_t *local_work_size,const Image *input_image,
1641% const Image *output_image,ExceptionInfo *exception)
1642%
1643% A description of each parameter follows:
1644%
1645% o kernel: the OpenCL kernel.
1646%
1647% o work_dim: the number of dimensions used to specify the global work-items
1648% and work-items in the work-group.
1649%
1650% o offset: can be used to specify an array of work_dim unsigned values
1651% that describe the offset used to calculate the global ID of a
1652% work-item.
1653%
1654% o gsize: points to an array of work_dim unsigned values that describe the
1655% number of global work-items in work_dim dimensions that will
1656% execute the kernel function.
1657%
1658% o lsize: points to an array of work_dim unsigned values that describe the
1659% number of work-items that make up a work-group that will execute
1660% the kernel specified by kernel.
1661%
1662% o input_image: the input image of the operation.
1663%
1664% o output_image: the output or secondary image of the operation.
1665%
1666% o exception: return any errors or warnings in this structure.
1667%
1668*/
1669
1670static MagickBooleanType RegisterCacheEvent(MagickCLCacheInfo info,
1671 cl_event event)
1672{
1673 assert(info != (MagickCLCacheInfo) NULL);
1674 assert(event != (cl_event) NULL);
1675 if (openCL_library->clRetainEvent(event) != CL_SUCCESS)
1676 {
1677 openCL_library->clWaitForEvents(1,&event);
1678 return(MagickFalse);
1679 }
1680 LockSemaphoreInfo(info->events_semaphore);
1681 if (info->events == (cl_event *) NULL)
1682 {
1683 info->events=(cl_event *) AcquireMagickMemory(sizeof(*info->events));
1684 info->event_count=1;
1685 }
1686 else
1687 info->events=(cl_event *) ResizeQuantumMemory(info->events,
1688 ++info->event_count,sizeof(*info->events));
1689 if (info->events == (cl_event *) NULL)
1690 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
1691 info->events[info->event_count-1]=event;
1692 UnlockSemaphoreInfo(info->events_semaphore);
1693 return(MagickTrue);
1694}
1695
1696MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1697 cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
1698 const size_t *lsize,const Image *input_image,const Image *output_image,
1699 MagickBooleanType flush,ExceptionInfo *exception)
1700{
1701 CacheInfo
1702 *output_info,
1703 *input_info;
1704
1705 cl_event
1706 event,
1707 *events;
1708
1709 cl_int
1710 status;
1711
1712 cl_uint
1713 event_count;
1714
1715 assert(input_image != (const Image *) NULL);
1716 input_info=(CacheInfo *) input_image->cache;
1717 assert(input_info != (CacheInfo *) NULL);
1718 assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1719 output_info=(CacheInfo *) NULL;
1720 if (output_image == (const Image *) NULL)
1721 events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
1722 &event_count);
1723 else
1724 {
1725 output_info=(CacheInfo *) output_image->cache;
1726 assert(output_info != (CacheInfo *) NULL);
1727 assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1728 events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
1729 &event_count);
1730 }
1731 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1732 gsize,lsize,event_count,events,&event);
1733 /* This can fail due to memory issues and calling clFinish might help. */
1734 if ((status != CL_SUCCESS) && (event_count > 0))
1735 {
1736 openCL_library->clFinish(queue);
1737 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
1738 offset,gsize,lsize,event_count,events,&event);
1739 }
1740 events=(cl_event *) RelinquishMagickMemory(events);
1741 if (status != CL_SUCCESS)
1742 {
1743 (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1744 GetMagickModule(),ResourceLimitWarning,
1745 "clEnqueueNDRangeKernel failed.","'%s'",".");
1746 return(MagickFalse);
1747 }
1748 if (flush != MagickFalse)
1749 openCL_library->clFlush(queue);
1750 if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1751 {
1752 if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
1753 {
1754 if (output_info != (CacheInfo *) NULL)
1755 (void) RegisterCacheEvent(output_info->opencl,event);
1756 }
1757 }
1758 openCL_library->clReleaseEvent(event);
1759 return(MagickTrue);
1760}
1761
1762/*
1763%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1764% %
1765% %
1766% %
1767+ G e t C u r r e n t O p e n C L E n v %
1768% %
1769% %
1770% %
1771%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1772%
1773% GetCurrentOpenCLEnv() returns the current OpenCL env
1774%
1775% The format of the GetCurrentOpenCLEnv method is:
1776%
1777% MagickCLEnv GetCurrentOpenCLEnv()
1778%
1779*/
1780
1781MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
1782{
1783 if (default_CLEnv != (MagickCLEnv) NULL)
1784 {
1785 if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1786 (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1787 return((MagickCLEnv) NULL);
1788 else
1789 return(default_CLEnv);
1790 }
1791
1792 if (GetOpenCLCacheDirectory() == (char *) NULL)
1793 return((MagickCLEnv) NULL);
1794
1795 if (openCL_lock == (SemaphoreInfo *) NULL)
1796 ActivateSemaphoreInfo(&openCL_lock);
1797
1798 LockSemaphoreInfo(openCL_lock);
1799 if (default_CLEnv == (MagickCLEnv) NULL)
1800 default_CLEnv=AcquireMagickCLEnv();
1801 UnlockSemaphoreInfo(openCL_lock);
1802
1803 return(default_CLEnv);
1804}
1805
1806/*
1807%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1808% %
1809% %
1810% %
1811% 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 %
1812% %
1813% %
1814% %
1815%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1816%
1817% GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
1818% device. The score is determined by the duration of the micro benchmark so
1819% that means a lower score is better than a higher score.
1820%
1821% The format of the GetOpenCLDeviceBenchmarkScore method is:
1822%
1823% double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
1824%
1825% A description of each parameter follows:
1826%
1827% o device: the OpenCL device.
1828*/
1829
1830MagickExport double GetOpenCLDeviceBenchmarkScore(
1831 const MagickCLDevice device)
1832{
1833 if (device == (MagickCLDevice) NULL)
1834 return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1835 return(device->score);
1836}
1837
1838/*
1839%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1840% %
1841% %
1842% %
1843% G e t O p e n C L D e v i c e E n a b l e d %
1844% %
1845% %
1846% %
1847%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1848%
1849% GetOpenCLDeviceEnabled() returns true if the device is enabled.
1850%
1851% The format of the GetOpenCLDeviceEnabled method is:
1852%
1853% MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
1854%
1855% A description of each parameter follows:
1856%
1857% o device: the OpenCL device.
1858*/
1859
1860MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
1861 const MagickCLDevice device)
1862{
1863 if (device == (MagickCLDevice) NULL)
1864 return(MagickFalse);
1865 return(device->enabled);
1866}
1867
1868/*
1869%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1870% %
1871% %
1872% %
1873% G e t O p e n C L D e v i c e N a m e %
1874% %
1875% %
1876% %
1877%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1878%
1879% GetOpenCLDeviceName() returns the name of the device.
1880%
1881% The format of the GetOpenCLDeviceName method is:
1882%
1883% const char *GetOpenCLDeviceName(const MagickCLDevice device)
1884%
1885% A description of each parameter follows:
1886%
1887% o device: the OpenCL device.
1888*/
1889
1890MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
1891{
1892 if (device == (MagickCLDevice) NULL)
1893 return((const char *) NULL);
1894 return(device->name);
1895}
1896
1897/*
1898%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1899% %
1900% %
1901% %
1902% 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 %
1903% %
1904% %
1905% %
1906%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1907%
1908% GetOpenCLDeviceVendorName() returns the vendor name of the device.
1909%
1910% The format of the GetOpenCLDeviceVendorName method is:
1911%
1912% const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1913%
1914% A description of each parameter follows:
1915%
1916% o device: the OpenCL device.
1917*/
1918
1919MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1920{
1921 if (device == (MagickCLDevice) NULL)
1922 return((const char *) NULL);
1923 return(device->vendor_name);
1924}
1925
1926/*
1927%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1928% %
1929% %
1930% %
1931% G e t O p e n C L D e v i c e s %
1932% %
1933% %
1934% %
1935%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1936%
1937% GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
1938% value of length to the number of devices that are available.
1939%
1940% The format of the GetOpenCLDevices method is:
1941%
1942% const MagickCLDevice *GetOpenCLDevices(size_t *length,
1943% ExceptionInfo *exception)
1944%
1945% A description of each parameter follows:
1946%
1947% o length: the number of device.
1948%
1949% o exception: return any errors or warnings in this structure.
1950%
1951*/
1952
1953MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
1954 ExceptionInfo *exception)
1955{
1956 MagickCLEnv
1957 clEnv;
1958
1959 clEnv=GetCurrentOpenCLEnv();
1960 if (clEnv == (MagickCLEnv) NULL)
1961 {
1962 if (length != (size_t *) NULL)
1963 *length=0;
1964 return((MagickCLDevice *) NULL);
1965 }
1966 InitializeOpenCL(clEnv,exception);
1967 if (length != (size_t *) NULL)
1968 *length=clEnv->number_devices;
1969 return(clEnv->devices);
1970}
1971
1972/*
1973%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1974% %
1975% %
1976% %
1977% G e t O p e n C L D e v i c e T y p e %
1978% %
1979% %
1980% %
1981%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1982%
1983% GetOpenCLDeviceType() returns the type of the device.
1984%
1985% The format of the GetOpenCLDeviceType method is:
1986%
1987% MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
1988%
1989% A description of each parameter follows:
1990%
1991% o device: the OpenCL device.
1992*/
1993
1994MagickExport MagickCLDeviceType GetOpenCLDeviceType(
1995 const MagickCLDevice device)
1996{
1997 if (device == (MagickCLDevice) NULL)
1998 return(UndefinedCLDeviceType);
1999 if (device->type == CL_DEVICE_TYPE_GPU)
2000 return(GpuCLDeviceType);
2001 if (device->type == CL_DEVICE_TYPE_CPU)
2002 return(CpuCLDeviceType);
2003 return(UndefinedCLDeviceType);
2004}
2005
2006/*
2007%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2008% %
2009% %
2010% %
2011% G e t O p e n C L D e v i c e V e r s i o n %
2012% %
2013% %
2014% %
2015%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2016%
2017% GetOpenCLDeviceVersion() returns the version of the device.
2018%
2019% The format of the GetOpenCLDeviceName method is:
2020%
2021% const char *GetOpenCLDeviceVersion(MagickCLDevice device)
2022%
2023% A description of each parameter follows:
2024%
2025% o device: the OpenCL device.
2026*/
2027
2028MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
2029{
2030 if (device == (MagickCLDevice) NULL)
2031 return((const char *) NULL);
2032 return(device->version);
2033}
2034
2035/*
2036%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2037% %
2038% %
2039% %
2040% G e t O p e n C L E n a b l e d %
2041% %
2042% %
2043% %
2044%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2045%
2046% GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
2047%
2048% The format of the GetOpenCLEnabled method is:
2049%
2050% MagickBooleanType GetOpenCLEnabled()
2051%
2052*/
2053
2054MagickExport MagickBooleanType GetOpenCLEnabled(void)
2055{
2056 MagickCLEnv
2057 clEnv;
2058
2059 clEnv=GetCurrentOpenCLEnv();
2060 if (clEnv == (MagickCLEnv) NULL)
2061 return(MagickFalse);
2062 return(clEnv->enabled);
2063}
2064
2065/*
2066%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2067% %
2068% %
2069% %
2070% 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 %
2071% %
2072% %
2073% %
2074%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2075%
2076% GetOpenCLKernelProfileRecords() returns the profile records for the
2077% specified device and sets length to the number of profile records.
2078%
2079% The format of the GetOpenCLKernelProfileRecords method is:
2080%
2081% const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
2082%
2083% A description of each parameter follows:
2084%
2085% o length: the number of profiles records.
2086*/
2087
2088MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
2089 const MagickCLDevice device,size_t *length)
2090{
2091 if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
2092 (KernelProfileRecord *) NULL))
2093 {
2094 if (length != (size_t *) NULL)
2095 *length=0;
2096 return((const KernelProfileRecord *) NULL);
2097 }
2098 if (length != (size_t *) NULL)
2099 {
2100 *length=0;
2101 LockSemaphoreInfo(device->lock);
2102 while (device->profile_records[*length] != (KernelProfileRecord) NULL)
2103 *length=*length+1;
2104 UnlockSemaphoreInfo(device->lock);
2105 }
2106 return(device->profile_records);
2107}
2108
2109/*
2110%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2111% %
2112% %
2113% %
2114% H a s O p e n C L D e v i c e s %
2115% %
2116% %
2117% %
2118%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2119%
2120% HasOpenCLDevices() checks if the OpenCL environment has devices that are
2121% enabled and compiles the kernel for the device when necessary. False will be
2122% returned if no enabled devices could be found
2123%
2124% The format of the HasOpenCLDevices method is:
2125%
2126% MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2127% ExceptionInfo exception)
2128%
2129% A description of each parameter follows:
2130%
2131% o clEnv: the OpenCL environment.
2132%
2133% o exception: return any errors or warnings in this structure.
2134%
2135*/
2136
2137static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2138 ExceptionInfo *exception)
2139{
2140 char
2141 *accelerateKernelsBuffer,
2142 options[MagickPathExtent];
2143
2144 MagickBooleanType
2145 status;
2146
2147 size_t
2148 i;
2149
2150 size_t
2151 signature;
2152
2153 /* Check if there are enabled devices */
2154 for (i = 0; i < clEnv->number_devices; i++)
2155 {
2156 if ((clEnv->devices[i]->enabled != MagickFalse))
2157 break;
2158 }
2159 if (i == clEnv->number_devices)
2160 return(MagickFalse);
2161
2162 /* Check if we need to compile a kernel for one of the devices */
2163 status=MagickTrue;
2164 for (i = 0; i < clEnv->number_devices; i++)
2165 {
2166 if ((clEnv->devices[i]->enabled != MagickFalse) &&
2167 (clEnv->devices[i]->program == (cl_program) NULL))
2168 {
2169 status=MagickFalse;
2170 break;
2171 }
2172 }
2173 if (status != MagickFalse)
2174 return(MagickTrue);
2175
2176 /* Get additional options */
2177 (void) FormatLocaleString(options,MagickPathExtent,CLOptions,
2178 (float)QuantumRange,(float)CLCharQuantumScale,(float)MagickEpsilon,
2179 (float)MagickPI,(unsigned int)MaxMap,(unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2180
2181 signature=StringSignature(options);
2182 accelerateKernelsBuffer=(char*) AcquireQuantumMemory(1,
2183 strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2184 if (accelerateKernelsBuffer == (char*) NULL)
2185 return(MagickFalse);
2186 (void) FormatLocaleString(accelerateKernelsBuffer,strlen(accelerateKernels)+
2187 strlen(accelerateKernels2)+1,"%s%s",accelerateKernels,accelerateKernels2);
2188 signature^=StringSignature(accelerateKernelsBuffer);
2189
2190 status=MagickTrue;
2191 for (i = 0; i < clEnv->number_devices; i++)
2192 {
2193 MagickCLDevice
2194 device;
2195
2196 size_t
2197 device_signature;
2198
2199 device=clEnv->devices[i];
2200 if ((device->enabled == MagickFalse) ||
2201 (device->program != (cl_program) NULL))
2202 continue;
2203
2204 LockSemaphoreInfo(device->lock);
2205 if (device->program != (cl_program) NULL)
2206 {
2207 UnlockSemaphoreInfo(device->lock);
2208 continue;
2209 }
2210 device_signature=signature;
2211 device_signature^=StringSignature(device->platform_name);
2212 status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2213 device_signature,exception);
2214 UnlockSemaphoreInfo(device->lock);
2215 if (status == MagickFalse)
2216 break;
2217 }
2218 accelerateKernelsBuffer=(char *) RelinquishMagickMemory(
2219 accelerateKernelsBuffer);
2220 return(status);
2221}
2222
2223/*
2224%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2225% %
2226% %
2227% %
2228+ I n i t i a l i z e O p e n C L %
2229% %
2230% %
2231% %
2232%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2233%
2234% InitializeOpenCL() is used to initialize the OpenCL environment. This method
2235% makes sure the devices are properly initialized and benchmarked.
2236%
2237% The format of the InitializeOpenCL method is:
2238%
2239% MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2240%
2241% A description of each parameter follows:
2242%
2243% o exception: return any errors or warnings in this structure.
2244%
2245*/
2246
2247static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2248{
2249 char
2250 version[MagickPathExtent];
2251
2252 cl_uint
2253 num;
2254
2255 if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2256 MagickPathExtent,version,NULL) != CL_SUCCESS)
2257 return(0);
2258 if (strncmp(version,"OpenCL 1.0 ",11) == 0)
2259 return(0);
2260 if (clEnv->library->clGetDeviceIDs(platform,
2261 CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2262 return(0);
2263 return(num);
2264}
2265
2266static inline char *GetOpenCLPlatformString(cl_platform_id platform,
2267 cl_platform_info param_name)
2268{
2269 char
2270 *value;
2271
2272 size_t
2273 length;
2274
2275 openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
2276 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2277 openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
2278 return(value);
2279}
2280
2281static inline char *GetOpenCLDeviceString(cl_device_id device,
2282 cl_device_info param_name)
2283{
2284 char
2285 *value;
2286
2287 size_t
2288 length;
2289
2290 openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
2291 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2292 openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
2293 return(value);
2294}
2295
2296static void LoadOpenCLDevices(MagickCLEnv clEnv)
2297{
2298 cl_context_properties
2299 properties[3];
2300
2301 cl_device_id
2302 *devices;
2303
2304 cl_int
2305 status;
2306
2307 cl_platform_id
2308 *platforms;
2309
2310 cl_uint
2311 i,
2312 j,
2313 next,
2314 number_devices,
2315 number_platforms;
2316
2317 number_platforms=0;
2318 if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2319 return;
2320 if (number_platforms == 0)
2321 return;
2322 platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
2323 sizeof(cl_platform_id));
2324 if (platforms == (cl_platform_id *) NULL)
2325 return;
2326 if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2327 {
2328 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2329 return;
2330 }
2331 for (i = 0; i < number_platforms; i++)
2332 {
2333 number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2334 if (number_devices == 0)
2335 platforms[i]=(cl_platform_id) NULL;
2336 else
2337 clEnv->number_devices+=number_devices;
2338 }
2339 if (clEnv->number_devices == 0)
2340 {
2341 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2342 return;
2343 }
2344 clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2345 sizeof(MagickCLDevice));
2346 if (clEnv->devices == (MagickCLDevice *) NULL)
2347 {
2348 RelinquishMagickCLDevices(clEnv);
2349 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2350 return;
2351 }
2352 (void) memset(clEnv->devices,0,clEnv->number_devices*sizeof(MagickCLDevice));
2353 devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2354 sizeof(cl_device_id));
2355 if (devices == (cl_device_id *) NULL)
2356 {
2357 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2358 RelinquishMagickCLDevices(clEnv);
2359 return;
2360 }
2361 (void) memset(devices,0,clEnv->number_devices*sizeof(cl_device_id));
2362 clEnv->number_contexts=(size_t) number_platforms;
2363 clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2364 sizeof(cl_context));
2365 if (clEnv->contexts == (cl_context *) NULL)
2366 {
2367 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2368 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2369 RelinquishMagickCLDevices(clEnv);
2370 return;
2371 }
2372 (void) memset(clEnv->contexts,0,clEnv->number_contexts*sizeof(cl_context));
2373 next=0;
2374 for (i = 0; i < number_platforms; i++)
2375 {
2376 if (platforms[i] == (cl_platform_id) NULL)
2377 continue;
2378
2379 status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2380 CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2381 if (status != CL_SUCCESS)
2382 continue;
2383
2384 properties[0]=CL_CONTEXT_PLATFORM;
2385 properties[1]=(cl_context_properties) platforms[i];
2386 properties[2]=0;
2387 clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2388 devices,NULL,NULL,&status);
2389 if (status != CL_SUCCESS)
2390 continue;
2391
2392 for (j = 0; j < number_devices; j++,next++)
2393 {
2394 MagickCLDevice
2395 device;
2396
2397 device=AcquireMagickCLDevice();
2398 if (device == (MagickCLDevice) NULL)
2399 break;
2400
2401 device->context=clEnv->contexts[i];
2402 device->deviceID=devices[j];
2403
2404 device->platform_name=GetOpenCLPlatformString(platforms[i],
2405 CL_PLATFORM_NAME);
2406
2407 device->vendor_name=GetOpenCLPlatformString(platforms[i],
2408 CL_PLATFORM_VENDOR);
2409
2410 device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);
2411
2412 device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);
2413
2414 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2415 sizeof(cl_uint),&device->max_clock_frequency,NULL);
2416
2417 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2418 sizeof(cl_uint),&device->max_compute_units,NULL);
2419
2420 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2421 sizeof(cl_device_type),&device->type,NULL);
2422
2423 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2424 sizeof(cl_ulong),&device->local_memory_size,NULL);
2425
2426 clEnv->devices[next]=device;
2427 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
2428 "Found device: %s (%s)",device->name,device->platform_name);
2429 }
2430 }
2431 if (next != clEnv->number_devices)
2432 RelinquishMagickCLDevices(clEnv);
2433 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2434 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2435}
2436
2437MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2438 ExceptionInfo *exception)
2439{
2440 LockSemaphoreInfo(clEnv->lock);
2441 if (clEnv->initialized != MagickFalse)
2442 {
2443 UnlockSemaphoreInfo(clEnv->lock);
2444 return(HasOpenCLDevices(clEnv,exception));
2445 }
2446 if (LoadOpenCLLibrary() != MagickFalse)
2447 {
2448 clEnv->library=openCL_library;
2449 LoadOpenCLDevices(clEnv);
2450 if (clEnv->number_devices > 0)
2451 AutoSelectOpenCLDevices(clEnv);
2452 }
2453 clEnv->initialized=MagickTrue;
2454 UnlockSemaphoreInfo(clEnv->lock);
2455 return(HasOpenCLDevices(clEnv,exception));
2456}
2457
2458/*
2459%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2460% %
2461% %
2462% %
2463% L o a d O p e n C L L i b r a r y %
2464% %
2465% %
2466% %
2467%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2468%
2469% LoadOpenCLLibrary() load and binds the OpenCL library.
2470%
2471% The format of the LoadOpenCLLibrary method is:
2472%
2473% MagickBooleanType LoadOpenCLLibrary(void)
2474%
2475*/
2476
2477void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2478{
2479 if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2480 return (void *) NULL;
2481 return lt_dlsym(library,functionName);
2482}
2483
2484static MagickBooleanType BindOpenCLFunctions()
2485{
2486#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
2487#define BIND(X) openCL_library->X= &X;
2488#else
2489 (void) memset(openCL_library,0,sizeof(MagickLibrary));
2490#ifdef MAGICKCORE_WINDOWS_SUPPORT
2491 openCL_library->library=(void *)lt_dlopen("OpenCL.dll");
2492#else
2493 openCL_library->library=(void *)lt_dlopen("libOpenCL.so");
2494#endif
2495#define BIND(X) \
2496 if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2497 return(MagickFalse);
2498#endif
2499
2500 if (openCL_library->library == (void*) NULL)
2501 return(MagickFalse);
2502
2503 BIND(clGetPlatformIDs);
2504 BIND(clGetPlatformInfo);
2505
2506 BIND(clGetDeviceIDs);
2507 BIND(clGetDeviceInfo);
2508
2509 BIND(clCreateBuffer);
2510 BIND(clReleaseMemObject);
2511 BIND(clRetainMemObject);
2512
2513 BIND(clCreateContext);
2514 BIND(clReleaseContext);
2515
2516 BIND(clCreateCommandQueue);
2517 BIND(clReleaseCommandQueue);
2518 BIND(clFlush);
2519 BIND(clFinish);
2520
2521 BIND(clCreateProgramWithSource);
2522 BIND(clCreateProgramWithBinary);
2523 BIND(clReleaseProgram);
2524 BIND(clBuildProgram);
2525 BIND(clGetProgramBuildInfo);
2526 BIND(clGetProgramInfo);
2527
2528 BIND(clCreateKernel);
2529 BIND(clReleaseKernel);
2530 BIND(clSetKernelArg);
2531 BIND(clGetKernelInfo);
2532
2533 BIND(clEnqueueReadBuffer);
2534 BIND(clEnqueueMapBuffer);
2535 BIND(clEnqueueUnmapMemObject);
2536 BIND(clEnqueueNDRangeKernel);
2537
2538 BIND(clGetEventInfo);
2539 BIND(clWaitForEvents);
2540 BIND(clReleaseEvent);
2541 BIND(clRetainEvent);
2542 BIND(clSetEventCallback);
2543
2544 BIND(clGetEventProfilingInfo);
2545
2546 return(MagickTrue);
2547}
2548
2549static MagickBooleanType LoadOpenCLLibrary(void)
2550{
2551 openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2552 if (openCL_library == (MagickLibrary *) NULL)
2553 return(MagickFalse);
2554
2555 if (BindOpenCLFunctions() == MagickFalse)
2556 {
2557 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2558 return(MagickFalse);
2559 }
2560
2561 return(MagickTrue);
2562}
2563
2564/*
2565%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2566% %
2567% %
2568% %
2569+ O p e n C L T e r m i n u s %
2570% %
2571% %
2572% %
2573%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2574%
2575% OpenCLTerminus() destroys the OpenCL component.
2576%
2577% The format of the OpenCLTerminus method is:
2578%
2579% OpenCLTerminus(void)
2580%
2581*/
2582
2583MagickPrivate void OpenCLTerminus()
2584{
2585 DumpOpenCLProfileData();
2586 if (cache_directory != (char *) NULL)
2587 cache_directory=DestroyString(cache_directory);
2588 if (cache_directory_lock != (SemaphoreInfo *) NULL)
2589 RelinquishSemaphoreInfo(&cache_directory_lock);
2590 if (default_CLEnv != (MagickCLEnv) NULL)
2591 default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2592 if (openCL_lock != (SemaphoreInfo *) NULL)
2593 RelinquishSemaphoreInfo(&openCL_lock);
2594 if (openCL_library != (MagickLibrary *) NULL)
2595 {
2596 if (openCL_library->library != (void *) NULL)
2597 (void) lt_dlclose(openCL_library->library);
2598 openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2599 }
2600}
2601
2602/*
2603%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2604% %
2605% %
2606% %
2607+ 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 %
2608% %
2609% %
2610% %
2611%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2612%
2613% OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2614% configuration file. If an error occurs, MagickFalse is returned
2615% otherwise MagickTrue.
2616%
2617% The format of the OpenCLThrowMagickException method is:
2618%
2619% MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2620% const char *module,const char *function,const size_t line,
2621% const ExceptionType severity,const char *tag,const char *format,...)
2622%
2623% A description of each parameter follows:
2624%
2625% o exception: the exception info.
2626%
2627% o filename: the source module filename.
2628%
2629% o function: the function name.
2630%
2631% o line: the line number of the source module.
2632%
2633% o severity: Specifies the numeric error category.
2634%
2635% o tag: the locale tag.
2636%
2637% o format: the output format.
2638%
2639*/
2640
2641MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2642 MagickCLDevice device,ExceptionInfo *exception,const char *module,
2643 const char *function,const size_t line,const ExceptionType severity,
2644 const char *tag,const char *format,...)
2645{
2646 MagickBooleanType
2647 status;
2648
2649 assert(device != (MagickCLDevice) NULL);
2650 assert(exception != (ExceptionInfo *) NULL);
2651 assert(exception->signature == MagickCoreSignature);
2652 (void) exception;
2653 status=MagickTrue;
2654 if (severity != 0)
2655 {
2656 if (device->type == CL_DEVICE_TYPE_CPU)
2657 {
2658 /* Workaround for Intel OpenCL CPU runtime bug */
2659 /* Turn off OpenCL when a problem is detected! */
2660 if (strncmp(device->platform_name,"Intel",5) == 0)
2661 default_CLEnv->enabled=MagickFalse;
2662 }
2663 }
2664
2665#ifdef OPENCLLOG_ENABLED
2666 {
2667 va_list
2668 operands;
2669 va_start(operands,format);
2670 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2671 format,operands);
2672 va_end(operands);
2673 }
2674#else
2675 magick_unreferenced(module);
2676 magick_unreferenced(function);
2677 magick_unreferenced(line);
2678 magick_unreferenced(tag);
2679 magick_unreferenced(format);
2680#endif
2681
2682 return(status);
2683}
2684
2685/*
2686%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2687% %
2688% %
2689% %
2690+ R e c o r d P r o f i l e D a t a %
2691% %
2692% %
2693% %
2694%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2695%
2696% RecordProfileData() records profile data.
2697%
2698% The format of the RecordProfileData method is:
2699%
2700% void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2701% cl_event event)
2702%
2703% A description of each parameter follows:
2704%
2705% o device: the OpenCL device that did the operation.
2706%
2707% o event: the event that contains the profiling data.
2708%
2709*/
2710
2711MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2712 cl_kernel kernel,cl_event event)
2713{
2714 char
2715 *name;
2716
2717 cl_int
2718 status;
2719
2720 cl_ulong
2721 elapsed,
2722 end,
2723 start;
2724
2725 KernelProfileRecord
2726 profile_record;
2727
2728 size_t
2729 i,
2730 length;
2731
2732 if (device->profile_kernels == MagickFalse)
2733 return(MagickFalse);
2734 status=openCL_library->clWaitForEvents(1,&event);
2735 if (status != CL_SUCCESS)
2736 return(MagickFalse);
2737 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2738 &length);
2739 if (status != CL_SUCCESS)
2740 return(MagickTrue);
2741 name=(char *) AcquireQuantumMemory(length,sizeof(*name));
2742 if (name == (char *) NULL)
2743 return(MagickTrue);
2744 start=end=elapsed=0;
2745 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2746 name,(size_t *) NULL);
2747 status|=openCL_library->clGetEventProfilingInfo(event,
2748 CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2749 status|=openCL_library->clGetEventProfilingInfo(event,
2750 CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2751 if (status != CL_SUCCESS)
2752 {
2753 name=DestroyString(name);
2754 return(MagickTrue);
2755 }
2756 start/=1000; /* usecs */
2757 end/=1000;
2758 elapsed=end-start;
2759 LockSemaphoreInfo(device->lock);
2760 i=0;
2761 profile_record=(KernelProfileRecord) NULL;
2762 if (device->profile_records != (KernelProfileRecord *) NULL)
2763 {
2764 while (device->profile_records[i] != (KernelProfileRecord) NULL)
2765 {
2766 if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2767 {
2768 profile_record=device->profile_records[i];
2769 break;
2770 }
2771 i++;
2772 }
2773 }
2774 if (profile_record != (KernelProfileRecord) NULL)
2775 name=DestroyString(name);
2776 else
2777 {
2778 profile_record=(KernelProfileRecord) AcquireCriticalMemory(
2779 sizeof(*profile_record));
2780 (void) memset(profile_record,0,sizeof(*profile_record));
2781 profile_record->kernel_name=name;
2782 device->profile_records=(KernelProfileRecord *) ResizeQuantumMemory(
2783 device->profile_records,(i+2),sizeof(*device->profile_records));
2784 if (device->profile_records == (KernelProfileRecord *) NULL)
2785 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
2786 device->profile_records[i]=profile_record;
2787 device->profile_records[i+1]=(KernelProfileRecord) NULL;
2788 }
2789 if ((elapsed < profile_record->min) || (profile_record->count == 0))
2790 profile_record->min=(unsigned long) elapsed;
2791 if (elapsed > profile_record->max)
2792 profile_record->max=(unsigned long) elapsed;
2793 profile_record->total+=(unsigned long) elapsed;
2794 profile_record->count+=1;
2795 UnlockSemaphoreInfo(device->lock);
2796 return(MagickTrue);
2797}
2798
2799/*
2800%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2801% %
2802% %
2803% %
2804+ 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 %
2805% %
2806% %
2807% %
2808%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2809%
2810% ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2811%
2812% The format of the ReleaseOpenCLCommandQueue method is:
2813%
2814% void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2815% cl_command_queue queue)
2816%
2817% A description of each parameter follows:
2818%
2819% o device: the OpenCL device.
2820%
2821% o queue: the OpenCL queue to be released.
2822*/
2823
2824MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2825 cl_command_queue queue)
2826{
2827 if (queue == (cl_command_queue) NULL)
2828 return;
2829
2830 assert(device != (MagickCLDevice) NULL);
2831 LockSemaphoreInfo(device->lock);
2832 if ((device->profile_kernels != MagickFalse) ||
2833 (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2834 {
2835 UnlockSemaphoreInfo(device->lock);
2836 openCL_library->clFinish(queue);
2837 (void) openCL_library->clReleaseCommandQueue(queue);
2838 }
2839 else
2840 {
2841 openCL_library->clFlush(queue);
2842 device->command_queues[++device->command_queues_index]=queue;
2843 UnlockSemaphoreInfo(device->lock);
2844 }
2845}
2846
2847/*
2848%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2849% %
2850% %
2851% %
2852+ R e l e a s e M a g i c k C L D e v i c e %
2853% %
2854% %
2855% %
2856%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2857%
2858% ReleaseOpenCLDevice() returns the OpenCL device to the environment
2859%
2860% The format of the ReleaseOpenCLDevice method is:
2861%
2862% void ReleaseOpenCLDevice(MagickCLDevice device)
2863%
2864% A description of each parameter follows:
2865%
2866% o device: the OpenCL device to be released.
2867%
2868*/
2869
2870MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2871{
2872 assert(device != (MagickCLDevice) NULL);
2873 LockSemaphoreInfo(openCL_lock);
2874 device->requested--;
2875 UnlockSemaphoreInfo(openCL_lock);
2876}
2877
2878/*
2879%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2880% %
2881% %
2882% %
2883+ 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 %
2884% %
2885% %
2886% %
2887%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2888%
2889% RelinquishMagickCLCacheInfo() frees memory acquired with
2890% AcquireMagickCLCacheInfo()
2891%
2892% The format of the RelinquishMagickCLCacheInfo method is:
2893%
2894% MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2895% const MagickBooleanType relinquish_pixels)
2896%
2897% A description of each parameter follows:
2898%
2899% o info: the OpenCL cache info.
2900%
2901% o relinquish_pixels: the pixels will be relinquish when set to true.
2902%
2903*/
2904
2905static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2906 cl_event magick_unused(event),
2907 cl_int magick_unused(event_command_exec_status),void *user_data)
2908{
2909 MagickCLCacheInfo
2910 info;
2911
2912 Quantum
2913 *pixels;
2914
2915 ssize_t
2916 i;
2917
2918 magick_unreferenced(event);
2919 magick_unreferenced(event_command_exec_status);
2920 info=(MagickCLCacheInfo) user_data;
2921 for (i=(ssize_t)info->event_count-1; i >= 0; i--)
2922 {
2923 cl_int
2924 event_status;
2925
2926 cl_uint
2927 status;
2928
2929 status=openCL_library->clGetEventInfo(info->events[i],
2930 CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
2931 NULL);
2932 if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
2933 {
2934 openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
2935 &DestroyMagickCLCacheInfoAndPixels,info);
2936 return;
2937 }
2938 }
2939 pixels=info->pixels;
2940 RelinquishMagickResource(MemoryResource,info->length);
2941 DestroyMagickCLCacheInfo(info);
2942 (void) RelinquishAlignedMemory(pixels);
2943}
2944
2945MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2946 MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2947{
2948 if (info == (MagickCLCacheInfo) NULL)
2949 return((MagickCLCacheInfo) NULL);
2950 if (relinquish_pixels != MagickFalse)
2951 DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2952 else
2953 DestroyMagickCLCacheInfo(info);
2954 return((MagickCLCacheInfo) NULL);
2955}
2956
2957/*
2958%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2959% %
2960% %
2961% %
2962% R e l i n q u i s h M a g i c k C L D e v i c e %
2963% %
2964% %
2965% %
2966%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2967%
2968% RelinquishMagickCLDevice() releases the OpenCL device
2969%
2970% The format of the RelinquishMagickCLDevice method is:
2971%
2972% MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2973%
2974% A description of each parameter follows:
2975%
2976% o device: the OpenCL device to be released.
2977%
2978*/
2979
2980static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2981{
2982 if (device == (MagickCLDevice) NULL)
2983 return((MagickCLDevice) NULL);
2984
2985 device->platform_name=(char *) RelinquishMagickMemory(device->platform_name);
2986 device->vendor_name=(char *) RelinquishMagickMemory(device->vendor_name);
2987 device->name=(char *) RelinquishMagickMemory(device->name);
2988 device->version=(char *) RelinquishMagickMemory(device->version);
2989 if (device->program != (cl_program) NULL)
2990 (void) openCL_library->clReleaseProgram(device->program);
2991 while (device->command_queues_index >= 0)
2992 (void) openCL_library->clReleaseCommandQueue(
2993 device->command_queues[device->command_queues_index--]);
2994 RelinquishSemaphoreInfo(&device->lock);
2995 return((MagickCLDevice) RelinquishMagickMemory(device));
2996}
2997
2998/*
2999%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3000% %
3001% %
3002% %
3003% R e l i n q u i s h M a g i c k C L E n v %
3004% %
3005% %
3006% %
3007%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3008%
3009% RelinquishMagickCLEnv() releases the OpenCL environment
3010%
3011% The format of the RelinquishMagickCLEnv method is:
3012%
3013% MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
3014%
3015% A description of each parameter follows:
3016%
3017% o clEnv: the OpenCL environment to be released.
3018%
3019*/
3020
3021static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
3022{
3023 if (clEnv == (MagickCLEnv) NULL)
3024 return((MagickCLEnv) NULL);
3025
3026 RelinquishSemaphoreInfo(&clEnv->lock);
3027 RelinquishMagickCLDevices(clEnv);
3028 if (clEnv->contexts != (cl_context *) NULL)
3029 {
3030 ssize_t
3031 i;
3032
3033 for (i=0; i < (ssize_t) clEnv->number_contexts; i++)
3034 if (clEnv->contexts[i] != (cl_context) NULL)
3035 (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
3036 clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
3037 }
3038 return((MagickCLEnv) RelinquishMagickMemory(clEnv));
3039}
3040
3041/*
3042%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3043% %
3044% %
3045% %
3046+ R e q u e s t O p e n C L D e v i c e %
3047% %
3048% %
3049% %
3050%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3051%
3052% RequestOpenCLDevice() returns one of the enabled OpenCL devices.
3053%
3054% The format of the RequestOpenCLDevice method is:
3055%
3056% MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3057%
3058% A description of each parameter follows:
3059%
3060% o clEnv: the OpenCL environment.
3061*/
3062
3063MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3064{
3065 MagickCLDevice
3066 device;
3067
3068 double
3069 score,
3070 best_score;
3071
3072 size_t
3073 i;
3074
3075 if (clEnv == (MagickCLEnv) NULL)
3076 return((MagickCLDevice) NULL);
3077
3078 if (clEnv->number_devices == 1)
3079 {
3080 if (clEnv->devices[0]->enabled)
3081 return(clEnv->devices[0]);
3082 else
3083 return((MagickCLDevice) NULL);
3084 }
3085
3086 device=(MagickCLDevice) NULL;
3087 best_score=0.0;
3088 LockSemaphoreInfo(openCL_lock);
3089 for (i = 0; i < clEnv->number_devices; i++)
3090 {
3091 if (clEnv->devices[i]->enabled == MagickFalse)
3092 continue;
3093
3094 score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3095 clEnv->devices[i]->requested);
3096 if ((device == (MagickCLDevice) NULL) || (score < best_score))
3097 {
3098 device=clEnv->devices[i];
3099 best_score=score;
3100 }
3101 }
3102 if (device != (MagickCLDevice)NULL)
3103 device->requested++;
3104 UnlockSemaphoreInfo(openCL_lock);
3105
3106 return(device);
3107}
3108
3109/*
3110%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3111% %
3112% %
3113% %
3114% S e t O p e n C L D e v i c e E n a b l e d %
3115% %
3116% %
3117% %
3118%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3119%
3120% SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
3121%
3122% The format of the SetOpenCLDeviceEnabled method is:
3123%
3124% void SetOpenCLDeviceEnabled(MagickCLDevice device,
3125% MagickBooleanType value)
3126%
3127% A description of each parameter follows:
3128%
3129% o device: the OpenCL device.
3130%
3131% o value: determines if the device should be enabled or disabled.
3132*/
3133
3134MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
3135 const MagickBooleanType value)
3136{
3137 if (device == (MagickCLDevice) NULL)
3138 return;
3139 device->enabled=value;
3140}
3141
3142/*
3143%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3144% %
3145% %
3146% %
3147% 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 %
3148% %
3149% %
3150% %
3151%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3152%
3153% SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
3154% kernel profiling of a device.
3155%
3156% The format of the SetOpenCLKernelProfileEnabled method is:
3157%
3158% void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3159% MagickBooleanType value)
3160%
3161% A description of each parameter follows:
3162%
3163% o device: the OpenCL device.
3164%
3165% o value: determines if kernel profiling for the device should be enabled
3166% or disabled.
3167*/
3168
3169MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3170 const MagickBooleanType value)
3171{
3172 if (device == (MagickCLDevice) NULL)
3173 return;
3174 device->profile_kernels=value;
3175}
3176
3177/*
3178%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3179% %
3180% %
3181% %
3182% S e t O p e n C L E n a b l e d %
3183% %
3184% %
3185% %
3186%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3187%
3188% SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3189%
3190% The format of the SetOpenCLEnabled method is:
3191%
3192% void SetOpenCLEnabled(MagickBooleanType)
3193%
3194% A description of each parameter follows:
3195%
3196% o value: specify true to enable OpenCL acceleration
3197*/
3198
3199MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
3200{
3201 MagickCLEnv
3202 clEnv;
3203
3204 clEnv=GetCurrentOpenCLEnv();
3205 if (clEnv == (MagickCLEnv) NULL)
3206 return(MagickFalse);
3207 clEnv->enabled=value;
3208 return(clEnv->enabled);
3209}
3210
3211#else
3212
3213MagickExport double GetOpenCLDeviceBenchmarkScore(
3214 const MagickCLDevice magick_unused(device))
3215{
3216 magick_unreferenced(device);
3217 return(0.0);
3218}
3219
3220MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3221 const MagickCLDevice magick_unused(device))
3222{
3223 magick_unreferenced(device);
3224 return(MagickFalse);
3225}
3226
3227MagickExport const char *GetOpenCLDeviceName(
3228 const MagickCLDevice magick_unused(device))
3229{
3230 magick_unreferenced(device);
3231 return((const char *) NULL);
3232}
3233
3234MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
3235 ExceptionInfo *magick_unused(exception))
3236{
3237 magick_unreferenced(exception);
3238 if (length != (size_t *) NULL)
3239 *length=0;
3240 return((MagickCLDevice *) NULL);
3241}
3242
3243MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3244 const MagickCLDevice magick_unused(device))
3245{
3246 magick_unreferenced(device);
3247 return(UndefinedCLDeviceType);
3248}
3249
3250MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
3251 const MagickCLDevice magick_unused(device),size_t *length)
3252{
3253 magick_unreferenced(device);
3254 if (length != (size_t *) NULL)
3255 *length=0;
3256 return((const KernelProfileRecord *) NULL);
3257}
3258
3259MagickExport const char *GetOpenCLDeviceVersion(
3260 const MagickCLDevice magick_unused(device))
3261{
3262 magick_unreferenced(device);
3263 return((const char *) NULL);
3264}
3265
3266MagickExport MagickBooleanType GetOpenCLEnabled(void)
3267{
3268 return(MagickFalse);
3269}
3270
3271MagickExport void SetOpenCLDeviceEnabled(
3272 MagickCLDevice magick_unused(device),
3273 const MagickBooleanType magick_unused(value))
3274{
3275 magick_unreferenced(device);
3276 magick_unreferenced(value);
3277}
3278
3279MagickExport MagickBooleanType SetOpenCLEnabled(
3280 const MagickBooleanType magick_unused(value))
3281{
3282 magick_unreferenced(value);
3283 return(MagickFalse);
3284}
3285
3286MagickExport void SetOpenCLKernelProfileEnabled(
3287 MagickCLDevice magick_unused(device),
3288 const MagickBooleanType magick_unused(value))
3289{
3290 magick_unreferenced(device);
3291 magick_unreferenced(value);
3292}
3293#endif