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