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