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