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