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