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)CLCharQuantumScale,(float)MagickEpsilon,
2172 (float)MagickPI,(unsigned int)MaxMap,(unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2173
2174 signature=StringSignature(options);
2175 accelerateKernelsBuffer=(char*) AcquireQuantumMemory(1,
2176 strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2177 if (accelerateKernelsBuffer == (char*) NULL)
2178 return(MagickFalse);
2179 (void) FormatLocaleString(accelerateKernelsBuffer,strlen(accelerateKernels)+
2180 strlen(accelerateKernels2)+1,"%s%s",accelerateKernels,accelerateKernels2);
2181 signature^=StringSignature(accelerateKernelsBuffer);
2182
2183 status=MagickTrue;
2184 for (i = 0; i < clEnv->number_devices; i++)
2185 {
2186 MagickCLDevice
2187 device;
2188
2189 size_t
2190 device_signature;
2191
2192 device=clEnv->devices[i];
2193 if ((device->enabled == MagickFalse) ||
2194 (device->program != (cl_program) NULL))
2195 continue;
2196
2197 LockSemaphoreInfo(device->lock);
2198 if (device->program != (cl_program) NULL)
2199 {
2200 UnlockSemaphoreInfo(device->lock);
2201 continue;
2202 }
2203 device_signature=signature;
2204 device_signature^=StringSignature(device->platform_name);
2205 status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2206 device_signature,exception);
2207 UnlockSemaphoreInfo(device->lock);
2208 if (status == MagickFalse)
2209 break;
2210 }
2211 accelerateKernelsBuffer=(char *) RelinquishMagickMemory(
2212 accelerateKernelsBuffer);
2213 return(status);
2214}
2215
2216/*
2217%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2218% %
2219% %
2220% %
2221+ I n i t i a l i z e O p e n C L %
2222% %
2223% %
2224% %
2225%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2226%
2227% InitializeOpenCL() is used to initialize the OpenCL environment. This method
2228% makes sure the devices are properly initialized and benchmarked.
2229%
2230% The format of the InitializeOpenCL method is:
2231%
2232% MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2233%
2234% A description of each parameter follows:
2235%
2236% o exception: return any errors or warnings in this structure.
2237%
2238*/
2239
2240static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2241{
2242 char
2243 version[MagickPathExtent];
2244
2245 cl_uint
2246 num;
2247
2248 if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2249 MagickPathExtent,version,NULL) != CL_SUCCESS)
2250 return(0);
2251 if (strncmp(version,"OpenCL 1.0 ",11) == 0)
2252 return(0);
2253 if (clEnv->library->clGetDeviceIDs(platform,
2254 CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2255 return(0);
2256 return(num);
2257}
2258
2259static inline char *GetOpenCLPlatformString(cl_platform_id platform,
2260 cl_platform_info param_name)
2261{
2262 char
2263 *value;
2264
2265 size_t
2266 length;
2267
2268 openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
2269 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2270 openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
2271 return(value);
2272}
2273
2274static inline char *GetOpenCLDeviceString(cl_device_id device,
2275 cl_device_info param_name)
2276{
2277 char
2278 *value;
2279
2280 size_t
2281 length;
2282
2283 openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
2284 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2285 openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
2286 return(value);
2287}
2288
2289static void LoadOpenCLDevices(MagickCLEnv clEnv)
2290{
2291 cl_context_properties
2292 properties[3];
2293
2294 cl_device_id
2295 *devices;
2296
2297 cl_int
2298 status;
2299
2300 cl_platform_id
2301 *platforms;
2302
2303 cl_uint
2304 i,
2305 j,
2306 next,
2307 number_devices,
2308 number_platforms;
2309
2310 number_platforms=0;
2311 if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2312 return;
2313 if (number_platforms == 0)
2314 return;
2315 platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
2316 sizeof(cl_platform_id));
2317 if (platforms == (cl_platform_id *) NULL)
2318 return;
2319 if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2320 {
2321 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2322 return;
2323 }
2324 for (i = 0; i < number_platforms; i++)
2325 {
2326 number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2327 if (number_devices == 0)
2328 platforms[i]=(cl_platform_id) NULL;
2329 else
2330 clEnv->number_devices+=number_devices;
2331 }
2332 if (clEnv->number_devices == 0)
2333 {
2334 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2335 return;
2336 }
2337 clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2338 sizeof(MagickCLDevice));
2339 if (clEnv->devices == (MagickCLDevice *) NULL)
2340 {
2341 RelinquishMagickCLDevices(clEnv);
2342 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2343 return;
2344 }
2345 (void) memset(clEnv->devices,0,clEnv->number_devices*sizeof(MagickCLDevice));
2346 devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2347 sizeof(cl_device_id));
2348 if (devices == (cl_device_id *) NULL)
2349 {
2350 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2351 RelinquishMagickCLDevices(clEnv);
2352 return;
2353 }
2354 (void) memset(devices,0,clEnv->number_devices*sizeof(cl_device_id));
2355 clEnv->number_contexts=(size_t) number_platforms;
2356 clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2357 sizeof(cl_context));
2358 if (clEnv->contexts == (cl_context *) NULL)
2359 {
2360 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2361 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2362 RelinquishMagickCLDevices(clEnv);
2363 return;
2364 }
2365 (void) memset(clEnv->contexts,0,clEnv->number_contexts*sizeof(cl_context));
2366 next=0;
2367 for (i = 0; i < number_platforms; i++)
2368 {
2369 if (platforms[i] == (cl_platform_id) NULL)
2370 continue;
2371
2372 status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2373 CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2374 if (status != CL_SUCCESS)
2375 continue;
2376
2377 properties[0]=CL_CONTEXT_PLATFORM;
2378 properties[1]=(cl_context_properties) platforms[i];
2379 properties[2]=0;
2380 clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2381 devices,NULL,NULL,&status);
2382 if (status != CL_SUCCESS)
2383 continue;
2384
2385 for (j = 0; j < number_devices; j++,next++)
2386 {
2387 MagickCLDevice
2388 device;
2389
2390 device=AcquireMagickCLDevice();
2391 if (device == (MagickCLDevice) NULL)
2392 break;
2393
2394 device->context=clEnv->contexts[i];
2395 device->deviceID=devices[j];
2396
2397 device->platform_name=GetOpenCLPlatformString(platforms[i],
2398 CL_PLATFORM_NAME);
2399
2400 device->vendor_name=GetOpenCLPlatformString(platforms[i],
2401 CL_PLATFORM_VENDOR);
2402
2403 device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);
2404
2405 device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);
2406
2407 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2408 sizeof(cl_uint),&device->max_clock_frequency,NULL);
2409
2410 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2411 sizeof(cl_uint),&device->max_compute_units,NULL);
2412
2413 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2414 sizeof(cl_device_type),&device->type,NULL);
2415
2416 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2417 sizeof(cl_ulong),&device->local_memory_size,NULL);
2418
2419 clEnv->devices[next]=device;
2420 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
2421 "Found device: %s (%s)",device->name,device->platform_name);
2422 }
2423 }
2424 if (next != clEnv->number_devices)
2425 RelinquishMagickCLDevices(clEnv);
2426 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2427 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2428}
2429
2430MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2431 ExceptionInfo *exception)
2432{
2433 LockSemaphoreInfo(clEnv->lock);
2434 if (clEnv->initialized != MagickFalse)
2435 {
2436 UnlockSemaphoreInfo(clEnv->lock);
2437 return(HasOpenCLDevices(clEnv,exception));
2438 }
2439 if (LoadOpenCLLibrary() != MagickFalse)
2440 {
2441 clEnv->library=openCL_library;
2442 LoadOpenCLDevices(clEnv);
2443 if (clEnv->number_devices > 0)
2444 AutoSelectOpenCLDevices(clEnv);
2445 }
2446 clEnv->initialized=MagickTrue;
2447 UnlockSemaphoreInfo(clEnv->lock);
2448 return(HasOpenCLDevices(clEnv,exception));
2449}
2450
2451/*
2452%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2453% %
2454% %
2455% %
2456% L o a d O p e n C L L i b r a r y %
2457% %
2458% %
2459% %
2460%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2461%
2462% LoadOpenCLLibrary() load and binds the OpenCL library.
2463%
2464% The format of the LoadOpenCLLibrary method is:
2465%
2466% MagickBooleanType LoadOpenCLLibrary(void)
2467%
2468*/
2469
2470void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2471{
2472 if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2473 return (void *) NULL;
2474 return lt_dlsym(library,functionName);
2475}
2476
2477static MagickBooleanType BindOpenCLFunctions()
2478{
2479#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
2480#define BIND(X) openCL_library->X= &X;
2481#else
2482 (void) memset(openCL_library,0,sizeof(MagickLibrary));
2483#ifdef MAGICKCORE_WINDOWS_SUPPORT
2484 openCL_library->library=(void *)lt_dlopen("OpenCL.dll");
2485#else
2486 openCL_library->library=(void *)lt_dlopen("libOpenCL.so");
2487#endif
2488#define BIND(X) \
2489 if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2490 return(MagickFalse);
2491#endif
2492
2493 if (openCL_library->library == (void*) NULL)
2494 return(MagickFalse);
2495
2496 BIND(clGetPlatformIDs);
2497 BIND(clGetPlatformInfo);
2498
2499 BIND(clGetDeviceIDs);
2500 BIND(clGetDeviceInfo);
2501
2502 BIND(clCreateBuffer);
2503 BIND(clReleaseMemObject);
2504 BIND(clRetainMemObject);
2505
2506 BIND(clCreateContext);
2507 BIND(clReleaseContext);
2508
2509 BIND(clCreateCommandQueue);
2510 BIND(clReleaseCommandQueue);
2511 BIND(clFlush);
2512 BIND(clFinish);
2513
2514 BIND(clCreateProgramWithSource);
2515 BIND(clCreateProgramWithBinary);
2516 BIND(clReleaseProgram);
2517 BIND(clBuildProgram);
2518 BIND(clGetProgramBuildInfo);
2519 BIND(clGetProgramInfo);
2520
2521 BIND(clCreateKernel);
2522 BIND(clReleaseKernel);
2523 BIND(clSetKernelArg);
2524 BIND(clGetKernelInfo);
2525
2526 BIND(clEnqueueReadBuffer);
2527 BIND(clEnqueueMapBuffer);
2528 BIND(clEnqueueUnmapMemObject);
2529 BIND(clEnqueueNDRangeKernel);
2530
2531 BIND(clGetEventInfo);
2532 BIND(clWaitForEvents);
2533 BIND(clReleaseEvent);
2534 BIND(clRetainEvent);
2535 BIND(clSetEventCallback);
2536
2537 BIND(clGetEventProfilingInfo);
2538
2539 return(MagickTrue);
2540}
2541
2542static MagickBooleanType LoadOpenCLLibrary(void)
2543{
2544 openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2545 if (openCL_library == (MagickLibrary *) NULL)
2546 return(MagickFalse);
2547
2548 if (BindOpenCLFunctions() == MagickFalse)
2549 {
2550 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2551 return(MagickFalse);
2552 }
2553
2554 return(MagickTrue);
2555}
2556
2557/*
2558%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2559% %
2560% %
2561% %
2562+ O p e n C L T e r m i n u s %
2563% %
2564% %
2565% %
2566%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2567%
2568% OpenCLTerminus() destroys the OpenCL component.
2569%
2570% The format of the OpenCLTerminus method is:
2571%
2572% OpenCLTerminus(void)
2573%
2574*/
2575
2576MagickPrivate void OpenCLTerminus()
2577{
2578 DumpOpenCLProfileData();
2579 if (cache_directory != (char *) NULL)
2580 cache_directory=DestroyString(cache_directory);
2581 if (cache_directory_lock != (SemaphoreInfo *) NULL)
2582 RelinquishSemaphoreInfo(&cache_directory_lock);
2583 if (default_CLEnv != (MagickCLEnv) NULL)
2584 default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2585 if (openCL_lock != (SemaphoreInfo *) NULL)
2586 RelinquishSemaphoreInfo(&openCL_lock);
2587 if (openCL_library != (MagickLibrary *) NULL)
2588 {
2589 if (openCL_library->library != (void *) NULL)
2590 (void) lt_dlclose(openCL_library->library);
2591 openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2592 }
2593}
2594
2595/*
2596%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2597% %
2598% %
2599% %
2600+ 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 %
2601% %
2602% %
2603% %
2604%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2605%
2606% OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2607% configuration file. If an error occurs, MagickFalse is returned
2608% otherwise MagickTrue.
2609%
2610% The format of the OpenCLThrowMagickException method is:
2611%
2612% MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2613% const char *module,const char *function,const size_t line,
2614% const ExceptionType severity,const char *tag,const char *format,...)
2615%
2616% A description of each parameter follows:
2617%
2618% o exception: the exception info.
2619%
2620% o filename: the source module filename.
2621%
2622% o function: the function name.
2623%
2624% o line: the line number of the source module.
2625%
2626% o severity: Specifies the numeric error category.
2627%
2628% o tag: the locale tag.
2629%
2630% o format: the output format.
2631%
2632*/
2633
2634MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2635 MagickCLDevice device,ExceptionInfo *exception,const char *module,
2636 const char *function,const size_t line,const ExceptionType severity,
2637 const char *tag,const char *format,...)
2638{
2639 MagickBooleanType
2640 status;
2641
2642 assert(device != (MagickCLDevice) NULL);
2643 assert(exception != (ExceptionInfo *) NULL);
2644 assert(exception->signature == MagickCoreSignature);
2645 (void) exception;
2646 status=MagickTrue;
2647 if (severity != 0)
2648 {
2649 if (device->type == CL_DEVICE_TYPE_CPU)
2650 {
2651 /* Workaround for Intel OpenCL CPU runtime bug */
2652 /* Turn off OpenCL when a problem is detected! */
2653 if (strncmp(device->platform_name,"Intel",5) == 0)
2654 default_CLEnv->enabled=MagickFalse;
2655 }
2656 }
2657
2658#ifdef OPENCLLOG_ENABLED
2659 {
2660 va_list
2661 operands;
2662 va_start(operands,format);
2663 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2664 format,operands);
2665 va_end(operands);
2666 }
2667#else
2668 magick_unreferenced(module);
2669 magick_unreferenced(function);
2670 magick_unreferenced(line);
2671 magick_unreferenced(tag);
2672 magick_unreferenced(format);
2673#endif
2674
2675 return(status);
2676}
2677
2678/*
2679%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2680% %
2681% %
2682% %
2683+ R e c o r d P r o f i l e D a t a %
2684% %
2685% %
2686% %
2687%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2688%
2689% RecordProfileData() records profile data.
2690%
2691% The format of the RecordProfileData method is:
2692%
2693% void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2694% cl_event event)
2695%
2696% A description of each parameter follows:
2697%
2698% o device: the OpenCL device that did the operation.
2699%
2700% o event: the event that contains the profiling data.
2701%
2702*/
2703
2704MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2705 cl_kernel kernel,cl_event event)
2706{
2707 char
2708 *name;
2709
2710 cl_int
2711 status;
2712
2713 cl_ulong
2714 elapsed,
2715 end,
2716 start;
2717
2718 KernelProfileRecord
2719 profile_record;
2720
2721 size_t
2722 i,
2723 length;
2724
2725 if (device->profile_kernels == MagickFalse)
2726 return(MagickFalse);
2727 status=openCL_library->clWaitForEvents(1,&event);
2728 if (status != CL_SUCCESS)
2729 return(MagickFalse);
2730 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2731 &length);
2732 if (status != CL_SUCCESS)
2733 return(MagickTrue);
2734 name=(char *) AcquireQuantumMemory(length,sizeof(*name));
2735 if (name == (char *) NULL)
2736 return(MagickTrue);
2737 start=end=elapsed=0;
2738 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2739 name,(size_t *) NULL);
2740 status|=openCL_library->clGetEventProfilingInfo(event,
2741 CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2742 status|=openCL_library->clGetEventProfilingInfo(event,
2743 CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2744 if (status != CL_SUCCESS)
2745 {
2746 name=DestroyString(name);
2747 return(MagickTrue);
2748 }
2749 start/=1000; /* usecs */
2750 end/=1000;
2751 elapsed=end-start;
2752 LockSemaphoreInfo(device->lock);
2753 i=0;
2754 profile_record=(KernelProfileRecord) NULL;
2755 if (device->profile_records != (KernelProfileRecord *) NULL)
2756 {
2757 while (device->profile_records[i] != (KernelProfileRecord) NULL)
2758 {
2759 if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2760 {
2761 profile_record=device->profile_records[i];
2762 break;
2763 }
2764 i++;
2765 }
2766 }
2767 if (profile_record != (KernelProfileRecord) NULL)
2768 name=DestroyString(name);
2769 else
2770 {
2771 profile_record=(KernelProfileRecord) AcquireCriticalMemory(
2772 sizeof(*profile_record));
2773 (void) memset(profile_record,0,sizeof(*profile_record));
2774 profile_record->kernel_name=name;
2775 device->profile_records=(KernelProfileRecord *) ResizeQuantumMemory(
2776 device->profile_records,(i+2),sizeof(*device->profile_records));
2777 if (device->profile_records == (KernelProfileRecord *) NULL)
2778 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
2779 device->profile_records[i]=profile_record;
2780 device->profile_records[i+1]=(KernelProfileRecord) NULL;
2781 }
2782 if ((elapsed < profile_record->min) || (profile_record->count == 0))
2783 profile_record->min=elapsed;
2784 if (elapsed > profile_record->max)
2785 profile_record->max=elapsed;
2786 profile_record->total+=elapsed;
2787 profile_record->count+=1;
2788 UnlockSemaphoreInfo(device->lock);
2789 return(MagickTrue);
2790}
2791
2792/*
2793%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2794% %
2795% %
2796% %
2797+ 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 %
2798% %
2799% %
2800% %
2801%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2802%
2803% ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2804%
2805% The format of the ReleaseOpenCLCommandQueue method is:
2806%
2807% void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2808% cl_command_queue queue)
2809%
2810% A description of each parameter follows:
2811%
2812% o device: the OpenCL device.
2813%
2814% o queue: the OpenCL queue to be released.
2815*/
2816
2817MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2818 cl_command_queue queue)
2819{
2820 if (queue == (cl_command_queue) NULL)
2821 return;
2822
2823 assert(device != (MagickCLDevice) NULL);
2824 LockSemaphoreInfo(device->lock);
2825 if ((device->profile_kernels != MagickFalse) ||
2826 (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2827 {
2828 UnlockSemaphoreInfo(device->lock);
2829 openCL_library->clFinish(queue);
2830 (void) openCL_library->clReleaseCommandQueue(queue);
2831 }
2832 else
2833 {
2834 openCL_library->clFlush(queue);
2835 device->command_queues[++device->command_queues_index]=queue;
2836 UnlockSemaphoreInfo(device->lock);
2837 }
2838}
2839
2840/*
2841%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2842% %
2843% %
2844% %
2845+ R e l e a s e M a g i c k C L D e v i c e %
2846% %
2847% %
2848% %
2849%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2850%
2851% ReleaseOpenCLDevice() returns the OpenCL device to the environment
2852%
2853% The format of the ReleaseOpenCLDevice method is:
2854%
2855% void ReleaseOpenCLDevice(MagickCLDevice device)
2856%
2857% A description of each parameter follows:
2858%
2859% o device: the OpenCL device to be released.
2860%
2861*/
2862
2863MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2864{
2865 assert(device != (MagickCLDevice) NULL);
2866 LockSemaphoreInfo(openCL_lock);
2867 device->requested--;
2868 UnlockSemaphoreInfo(openCL_lock);
2869}
2870
2871/*
2872%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2873% %
2874% %
2875% %
2876+ 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 %
2877% %
2878% %
2879% %
2880%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2881%
2882% RelinquishMagickCLCacheInfo() frees memory acquired with
2883% AcquireMagickCLCacheInfo()
2884%
2885% The format of the RelinquishMagickCLCacheInfo method is:
2886%
2887% MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2888% const MagickBooleanType relinquish_pixels)
2889%
2890% A description of each parameter follows:
2891%
2892% o info: the OpenCL cache info.
2893%
2894% o relinquish_pixels: the pixels will be relinquish when set to true.
2895%
2896*/
2897
2898static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2899 cl_event magick_unused(event),
2900 cl_int magick_unused(event_command_exec_status),void *user_data)
2901{
2902 MagickCLCacheInfo
2903 info;
2904
2905 Quantum
2906 *pixels;
2907
2908 ssize_t
2909 i;
2910
2911 magick_unreferenced(event);
2912 magick_unreferenced(event_command_exec_status);
2913 info=(MagickCLCacheInfo) user_data;
2914 for (i=(ssize_t)info->event_count-1; i >= 0; i--)
2915 {
2916 cl_int
2917 event_status;
2918
2919 cl_uint
2920 status;
2921
2922 status=openCL_library->clGetEventInfo(info->events[i],
2923 CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
2924 NULL);
2925 if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
2926 {
2927 openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
2928 &DestroyMagickCLCacheInfoAndPixels,info);
2929 return;
2930 }
2931 }
2932 pixels=info->pixels;
2933 RelinquishMagickResource(MemoryResource,info->length);
2934 DestroyMagickCLCacheInfo(info);
2935 (void) RelinquishAlignedMemory(pixels);
2936}
2937
2938MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2939 MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2940{
2941 if (info == (MagickCLCacheInfo) NULL)
2942 return((MagickCLCacheInfo) NULL);
2943 if (relinquish_pixels != MagickFalse)
2944 DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2945 else
2946 DestroyMagickCLCacheInfo(info);
2947 return((MagickCLCacheInfo) NULL);
2948}
2949
2950/*
2951%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2952% %
2953% %
2954% %
2955% R e l i n q u i s h M a g i c k C L D e v i c e %
2956% %
2957% %
2958% %
2959%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2960%
2961% RelinquishMagickCLDevice() releases the OpenCL device
2962%
2963% The format of the RelinquishMagickCLDevice method is:
2964%
2965% MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2966%
2967% A description of each parameter follows:
2968%
2969% o device: the OpenCL device to be released.
2970%
2971*/
2972
2973static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2974{
2975 if (device == (MagickCLDevice) NULL)
2976 return((MagickCLDevice) NULL);
2977
2978 device->platform_name=(char *) RelinquishMagickMemory(device->platform_name);
2979 device->vendor_name=(char *) RelinquishMagickMemory(device->vendor_name);
2980 device->name=(char *) RelinquishMagickMemory(device->name);
2981 device->version=(char *) RelinquishMagickMemory(device->version);
2982 if (device->program != (cl_program) NULL)
2983 (void) openCL_library->clReleaseProgram(device->program);
2984 while (device->command_queues_index >= 0)
2985 (void) openCL_library->clReleaseCommandQueue(
2986 device->command_queues[device->command_queues_index--]);
2987 RelinquishSemaphoreInfo(&device->lock);
2988 return((MagickCLDevice) RelinquishMagickMemory(device));
2989}
2990
2991/*
2992%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2993% %
2994% %
2995% %
2996% R e l i n q u i s h M a g i c k C L E n v %
2997% %
2998% %
2999% %
3000%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3001%
3002% RelinquishMagickCLEnv() releases the OpenCL environment
3003%
3004% The format of the RelinquishMagickCLEnv method is:
3005%
3006% MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
3007%
3008% A description of each parameter follows:
3009%
3010% o clEnv: the OpenCL environment to be released.
3011%
3012*/
3013
3014static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
3015{
3016 if (clEnv == (MagickCLEnv) NULL)
3017 return((MagickCLEnv) NULL);
3018
3019 RelinquishSemaphoreInfo(&clEnv->lock);
3020 RelinquishMagickCLDevices(clEnv);
3021 if (clEnv->contexts != (cl_context *) NULL)
3022 {
3023 ssize_t
3024 i;
3025
3026 for (i=0; i < clEnv->number_contexts; i++)
3027 if (clEnv->contexts[i] != (cl_context) NULL)
3028 (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
3029 clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
3030 }
3031 return((MagickCLEnv) RelinquishMagickMemory(clEnv));
3032}
3033
3034/*
3035%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3036% %
3037% %
3038% %
3039+ R e q u e s t O p e n C L D e v i c e %
3040% %
3041% %
3042% %
3043%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3044%
3045% RequestOpenCLDevice() returns one of the enabled OpenCL devices.
3046%
3047% The format of the RequestOpenCLDevice method is:
3048%
3049% MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3050%
3051% A description of each parameter follows:
3052%
3053% o clEnv: the OpenCL environment.
3054*/
3055
3056MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3057{
3058 MagickCLDevice
3059 device;
3060
3061 double
3062 score,
3063 best_score;
3064
3065 size_t
3066 i;
3067
3068 if (clEnv == (MagickCLEnv) NULL)
3069 return((MagickCLDevice) NULL);
3070
3071 if (clEnv->number_devices == 1)
3072 {
3073 if (clEnv->devices[0]->enabled)
3074 return(clEnv->devices[0]);
3075 else
3076 return((MagickCLDevice) NULL);
3077 }
3078
3079 device=(MagickCLDevice) NULL;
3080 best_score=0.0;
3081 LockSemaphoreInfo(openCL_lock);
3082 for (i = 0; i < clEnv->number_devices; i++)
3083 {
3084 if (clEnv->devices[i]->enabled == MagickFalse)
3085 continue;
3086
3087 score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3088 clEnv->devices[i]->requested);
3089 if ((device == (MagickCLDevice) NULL) || (score < best_score))
3090 {
3091 device=clEnv->devices[i];
3092 best_score=score;
3093 }
3094 }
3095 if (device != (MagickCLDevice)NULL)
3096 device->requested++;
3097 UnlockSemaphoreInfo(openCL_lock);
3098
3099 return(device);
3100}
3101
3102/*
3103%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3104% %
3105% %
3106% %
3107% S e t O p e n C L D e v i c e E n a b l e d %
3108% %
3109% %
3110% %
3111%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3112%
3113% SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
3114%
3115% The format of the SetOpenCLDeviceEnabled method is:
3116%
3117% void SetOpenCLDeviceEnabled(MagickCLDevice device,
3118% MagickBooleanType value)
3119%
3120% A description of each parameter follows:
3121%
3122% o device: the OpenCL device.
3123%
3124% o value: determines if the device should be enabled or disabled.
3125*/
3126
3127MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
3128 const MagickBooleanType value)
3129{
3130 if (device == (MagickCLDevice) NULL)
3131 return;
3132 device->enabled=value;
3133}
3134
3135/*
3136%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3137% %
3138% %
3139% %
3140% 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 %
3141% %
3142% %
3143% %
3144%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3145%
3146% SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
3147% kernel profiling of a device.
3148%
3149% The format of the SetOpenCLKernelProfileEnabled method is:
3150%
3151% void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3152% MagickBooleanType value)
3153%
3154% A description of each parameter follows:
3155%
3156% o device: the OpenCL device.
3157%
3158% o value: determines if kernel profiling for the device should be enabled
3159% or disabled.
3160*/
3161
3162MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3163 const MagickBooleanType value)
3164{
3165 if (device == (MagickCLDevice) NULL)
3166 return;
3167 device->profile_kernels=value;
3168}
3169
3170/*
3171%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3172% %
3173% %
3174% %
3175% S e t O p e n C L E n a b l e d %
3176% %
3177% %
3178% %
3179%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3180%
3181% SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3182%
3183% The format of the SetOpenCLEnabled method is:
3184%
3185% void SetOpenCLEnabled(MagickBooleanType)
3186%
3187% A description of each parameter follows:
3188%
3189% o value: specify true to enable OpenCL acceleration
3190*/
3191
3192MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
3193{
3194 MagickCLEnv
3195 clEnv;
3196
3197 clEnv=GetCurrentOpenCLEnv();
3198 if (clEnv == (MagickCLEnv) NULL)
3199 return(MagickFalse);
3200 clEnv->enabled=value;
3201 return(clEnv->enabled);
3202}
3203
3204#else
3205
3206MagickExport double GetOpenCLDeviceBenchmarkScore(
3207 const MagickCLDevice magick_unused(device))
3208{
3209 magick_unreferenced(device);
3210 return(0.0);
3211}
3212
3213MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3214 const MagickCLDevice magick_unused(device))
3215{
3216 magick_unreferenced(device);
3217 return(MagickFalse);
3218}
3219
3220MagickExport const char *GetOpenCLDeviceName(
3221 const MagickCLDevice magick_unused(device))
3222{
3223 magick_unreferenced(device);
3224 return((const char *) NULL);
3225}
3226
3227MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
3228 ExceptionInfo *magick_unused(exception))
3229{
3230 magick_unreferenced(exception);
3231 if (length != (size_t *) NULL)
3232 *length=0;
3233 return((MagickCLDevice *) NULL);
3234}
3235
3236MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3237 const MagickCLDevice magick_unused(device))
3238{
3239 magick_unreferenced(device);
3240 return(UndefinedCLDeviceType);
3241}
3242
3243MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
3244 const MagickCLDevice magick_unused(device),size_t *length)
3245{
3246 magick_unreferenced(device);
3247 if (length != (size_t *) NULL)
3248 *length=0;
3249 return((const KernelProfileRecord *) NULL);
3250}
3251
3252MagickExport const char *GetOpenCLDeviceVersion(
3253 const MagickCLDevice magick_unused(device))
3254{
3255 magick_unreferenced(device);
3256 return((const char *) NULL);
3257}
3258
3259MagickExport MagickBooleanType GetOpenCLEnabled(void)
3260{
3261 return(MagickFalse);
3262}
3263
3264MagickExport void SetOpenCLDeviceEnabled(
3265 MagickCLDevice magick_unused(device),
3266 const MagickBooleanType magick_unused(value))
3267{
3268 magick_unreferenced(device);
3269 magick_unreferenced(value);
3270}
3271
3272MagickExport MagickBooleanType SetOpenCLEnabled(
3273 const MagickBooleanType magick_unused(value))
3274{
3275 magick_unreferenced(value);
3276 return(MagickFalse);
3277}
3278
3279MagickExport void SetOpenCLKernelProfileEnabled(
3280 MagickCLDevice magick_unused(device),
3281 const MagickBooleanType magick_unused(value))
3282{
3283 magick_unreferenced(device);
3284 magick_unreferenced(value);
3285}
3286#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