MagickCore 7.1.2-33
Convert, Edit, Or Compose Bitmap Images
Loading...
Searching...
No Matches
opencl.c
1/*
2%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3% %
4% %
5% %
6% OOO PPPP EEEEE N N CCCC L %
7% O O P P E NN N C L %
8% O O PPPP EEE N N N C L %
9% O O P E N NN C L %
10% OOO P EEEEE N N CCCC LLLLL %
11% %
12% %
13% MagickCore OpenCL Methods %
14% %
15% Software Design %
16% Cristy %
17% March 2000 %
18% %
19% %
20% Copyright @ 1999 ImageMagick Studio LLC, a non-profit organization %
21% dedicated to making software imaging solutions freely available. %
22% %
23% You may not use this file except in compliance with the License. You may %
24% obtain a copy of the License at %
25% %
26% https://imagemagick.org/license/ %
27% %
28% Unless required by applicable law or agreed to in writing, software %
29% distributed under the License is distributed on an "AS IS" BASIS, %
30% WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. %
31% See the License for the specific language governing permissions and %
32% limitations under the License. %
33% %
34%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
35%
36%
37%
38*/
39␌
40/*
41 Include declarations.
42*/
43#include "MagickCore/studio.h"
44#include "MagickCore/accelerate-kernels-private.h"
45#include "MagickCore/artifact.h"
46#include "MagickCore/cache.h"
47#include "MagickCore/cache-private.h"
48#include "MagickCore/color.h"
49#include "MagickCore/compare.h"
50#include "MagickCore/constitute.h"
51#include "MagickCore/configure.h"
52#include "MagickCore/distort.h"
53#include "MagickCore/draw.h"
54#include "MagickCore/effect.h"
55#include "MagickCore/exception.h"
56#include "MagickCore/exception-private.h"
57#include "MagickCore/fx.h"
58#include "MagickCore/gem.h"
59#include "MagickCore/geometry.h"
60#include "MagickCore/image.h"
61#include "MagickCore/image-private.h"
62#include "MagickCore/layer.h"
63#include "MagickCore/locale_.h"
64#include "MagickCore/mime-private.h"
65#include "MagickCore/memory_.h"
66#include "MagickCore/memory-private.h"
67#include "MagickCore/monitor.h"
68#include "MagickCore/montage.h"
69#include "MagickCore/morphology.h"
70#include "MagickCore/nt-base.h"
71#include "MagickCore/nt-base-private.h"
72#include "MagickCore/opencl.h"
73#include "MagickCore/opencl-private.h"
74#include "MagickCore/option.h"
75#include "MagickCore/policy.h"
76#include "MagickCore/property.h"
77#include "MagickCore/quantize.h"
78#include "MagickCore/quantum.h"
79#include "MagickCore/random_.h"
80#include "MagickCore/random-private.h"
81#include "MagickCore/resample.h"
82#include "MagickCore/resource_.h"
83#include "MagickCore/splay-tree.h"
84#include "MagickCore/semaphore.h"
85#include "MagickCore/statistic.h"
86#include "MagickCore/string_.h"
87#include "MagickCore/string-private.h"
88#include "MagickCore/token.h"
89#include "MagickCore/utility.h"
90#include "MagickCore/utility-private.h"
91#include "MagickCore/xml-tree-private.h"
92
93#if defined(MAGICKCORE_OPENCL_SUPPORT)
94#if defined(MAGICKCORE_LTDL_DELEGATE)
95#include "ltdl.h"
96#endif
97
98/*
99 Define declarations.
100*/
101#define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"
102
103/*
104 Typedef declarations.
105*/
106typedef struct
107{
108 long long freq;
109 long long clocks;
110 long long start;
111} AccelerateTimer;
112
113typedef struct
114{
115 char
116 *name,
117 *platform_name,
118 *vendor_name,
119 *version;
120
121 cl_uint
122 max_clock_frequency,
123 max_compute_units;
124
125 double
126 score;
127} MagickCLDeviceBenchmark;
128
129/*
130 Forward declarations.
131*/
132
133static MagickBooleanType
134 HasOpenCLDevices(MagickCLEnv,ExceptionInfo *),
135 LoadOpenCLLibrary(void);
136
137static MagickCLDevice
138 RelinquishMagickCLDevice(MagickCLDevice);
139
140static MagickCLEnv
141 RelinquishMagickCLEnv(MagickCLEnv);
142
143static void
144 BenchmarkOpenCLDevices(MagickCLEnv);
145
146/* OpenCL library */
147MagickLibrary
148 *openCL_library;
149
150/* Default OpenCL environment */
151MagickCLEnv
152 default_CLEnv;
153MagickThreadType
154 test_thread_id=0;
155SemaphoreInfo
156 *openCL_lock;
157
158/* Cached location of the OpenCL cache files */
159char
160 *cache_directory;
161SemaphoreInfo
162 *cache_directory_lock;
163
164static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
165 MagickCLDevice b)
166{
167 if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
168 (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
169 (LocaleCompare(a->name,b->name) == 0) &&
170 (LocaleCompare(a->version,b->version) == 0) &&
171 (a->max_clock_frequency == b->max_clock_frequency) &&
172 (a->max_compute_units == b->max_compute_units))
173 return(MagickTrue);
174
175 return(MagickFalse);
176}
177
178static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
179 MagickCLDeviceBenchmark *b)
180{
181 if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
182 (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
183 (LocaleCompare(a->name,b->name) == 0) &&
184 (LocaleCompare(a->version,b->version) == 0) &&
185 (a->max_clock_frequency == b->max_clock_frequency) &&
186 (a->max_compute_units == b->max_compute_units))
187 return(MagickTrue);
188
189 return(MagickFalse);
190}
191
192static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
193{
194 size_t
195 i;
196
197 if (clEnv->devices != (MagickCLDevice *) NULL)
198 {
199 for (i = 0; i < clEnv->number_devices; i++)
200 clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
201 clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
202 }
203 clEnv->number_devices=0;
204}
205
206static inline MagickBooleanType MagickCreateDirectory(const char *path)
207{
208 int
209 status;
210
211#ifdef MAGICKCORE_WINDOWS_SUPPORT
212 status=_mkdir(path);
213#else
214 status=mkdir(path,0777);
215#endif
216 return(status == 0 ? MagickTrue : MagickFalse);
217}
218
219static inline void InitAccelerateTimer(AccelerateTimer *timer)
220{
221#ifdef _WIN32
222 QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
223#else
224 timer->freq=(long long)1.0E3;
225#endif
226 timer->clocks=0;
227 timer->start=0;
228}
229
230static inline double ReadAccelerateTimer(AccelerateTimer *timer)
231{
232 return (double)timer->clocks/(double)timer->freq;
233}
234
235static inline void StartAccelerateTimer(AccelerateTimer* timer)
236{
237#ifdef _WIN32
238 QueryPerformanceCounter((LARGE_INTEGER*)&timer->start);
239#else
240 struct timeval
241 s;
242 gettimeofday(&s,0);
243 timer->start=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
244 (long long)1.0E3;
245#endif
246}
247
248static inline void StopAccelerateTimer(AccelerateTimer *timer)
249{
250 long long
251 n;
252
253 n=0;
254#ifdef _WIN32
255 QueryPerformanceCounter((LARGE_INTEGER*)&(n));
256#else
257 struct timeval
258 s;
259 gettimeofday(&s,0);
260 n=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
261 (long long)1.0E3;
262#endif
263 n-=timer->start;
264 timer->start=0;
265 timer->clocks+=n;
266}
267
268static const char *GetOpenCLCacheDirectory()
269{
270 if (cache_directory == (char *) NULL)
271 {
272 if (cache_directory_lock == (SemaphoreInfo *) NULL)
273 ActivateSemaphoreInfo(&cache_directory_lock);
274 LockSemaphoreInfo(cache_directory_lock);
275 if (cache_directory == (char *) NULL)
276 {
277 char
278 *home,
279 path[MagickPathExtent],
280 *temp;
281
282 MagickBooleanType
283 status;
284
285 struct stat
286 attributes;
287
288 temp=(char *) NULL;
289 home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
290 if (home == (char *) NULL)
291 {
292 home=GetEnvironmentValue("XDG_CACHE_HOME");
293#if defined(MAGICKCORE_WINDOWS_SUPPORT) || defined(__MINGW32__)
294 if (home == (char *) NULL)
295 home=GetEnvironmentValue("LOCALAPPDATA");
296 if (home == (char *) NULL)
297 home=GetEnvironmentValue("APPDATA");
298 if (home == (char *) NULL)
299 home=GetEnvironmentValue("USERPROFILE");
300#endif
301 }
302
303 if (home != (char *) NULL)
304 {
305 /* first check if $HOME exists */
306 (void) FormatLocaleString(path,MagickPathExtent,"%s",home);
307 status=GetPathAttributes(path,&attributes);
308 if (status == MagickFalse)
309 status=MagickCreateDirectory(path);
310
311 /* first check if $HOME/ImageMagick exists */
312 if (status != MagickFalse)
313 {
314 (void) FormatLocaleString(path,MagickPathExtent,
315 "%s%sImageMagick",home,DirectorySeparator);
316
317 status=GetPathAttributes(path,&attributes);
318 if (status == MagickFalse)
319 status=MagickCreateDirectory(path);
320 }
321
322 if (status != MagickFalse)
323 {
324 temp=(char*) AcquireCriticalMemory(strlen(path)+1);
325 (void) CopyMagickString(temp,path,strlen(path)+1);
326 }
327 home=DestroyString(home);
328 }
329 else
330 {
331 home=GetEnvironmentValue("HOME");
332 if (home != (char *) NULL)
333 {
334 /* first check if $HOME/.cache exists */
335 (void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache",
336 home,DirectorySeparator);
337 status=GetPathAttributes(path,&attributes);
338 if (status == MagickFalse)
339 status=MagickCreateDirectory(path);
340
341 /* first check if $HOME/.cache/ImageMagick exists */
342 if (status != MagickFalse)
343 {
344 (void) FormatLocaleString(path,MagickPathExtent,
345 "%s%s.cache%sImageMagick",home,DirectorySeparator,
346 DirectorySeparator);
347 status=GetPathAttributes(path,&attributes);
348 if (status == MagickFalse)
349 status=MagickCreateDirectory(path);
350 }
351
352 if (status != MagickFalse)
353 {
354 temp=(char*) AcquireCriticalMemory(strlen(path)+1);
355 (void) CopyMagickString(temp,path,strlen(path)+1);
356 }
357 home=DestroyString(home);
358 }
359 }
360 if (temp == (char *) NULL)
361 {
362 temp=AcquireString("?");
363 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
364 "Cannot use cache directory: \"%s\"",path);
365 }
366 else
367 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
368 "Using cache directory: \"%s\"",temp);
369 cache_directory=temp;
370 }
371 UnlockSemaphoreInfo(cache_directory_lock);
372 }
373 if (*cache_directory == '?')
374 return((const char *) NULL);
375 return(cache_directory);
376}
377
378static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
379{
380 MagickCLDevice
381 device;
382
383 size_t
384 i,
385 j;
386
387 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
388 "Selecting device for type: %d",(int) type);
389 for (i = 0; i < clEnv->number_devices; i++)
390 clEnv->devices[i]->enabled=MagickFalse;
391
392 for (i = 0; i < clEnv->number_devices; i++)
393 {
394 device=clEnv->devices[i];
395 if (device->type != type)
396 continue;
397
398 device->enabled=MagickTrue;
399 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
400 "Selected device: %s",device->name);
401 for (j = i+1; j < clEnv->number_devices; j++)
402 {
403 MagickCLDevice
404 other_device;
405
406 other_device=clEnv->devices[j];
407 if (IsSameOpenCLDevice(device,other_device))
408 other_device->enabled=MagickTrue;
409 }
410 }
411}
412
413static size_t StringSignature(const char* string)
414{
415 size_t
416 n,
417 i,
418 j,
419 signature,
420 stringLength;
421
422 union
423 {
424 const char* s;
425 const size_t* u;
426 } p;
427
428 stringLength=(size_t) strlen(string);
429 signature=stringLength;
430 n=stringLength/sizeof(size_t);
431 p.s=string;
432 for (i = 0; i < n; i++)
433 signature^=p.u[i];
434 if (n * sizeof(size_t) != stringLength)
435 {
436 char
437 padded[4];
438
439 j=n*sizeof(size_t);
440 for (i = 0; i < 4; i++, j++)
441 {
442 if (j < stringLength)
443 padded[i]=p.s[j];
444 else
445 padded[i]=0;
446 }
447 p.s=padded;
448 signature^=p.u[0];
449 }
450 return(signature);
451}
452
453static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
454{
455 ssize_t
456 i;
457
458 for (i=0; i < (ssize_t) info->event_count; i++)
459 openCL_library->clReleaseEvent(info->events[i]);
460 info->events=(cl_event *) RelinquishMagickMemory(info->events);
461 if (info->buffer != (cl_mem) NULL)
462 openCL_library->clReleaseMemObject(info->buffer);
463 RelinquishSemaphoreInfo(&info->events_semaphore);
464 ReleaseOpenCLDevice(info->device);
465 RelinquishMagickMemory(info);
466}
467
468/*
469 Provide call to OpenCL library methods
470*/
471
472MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
473 cl_mem_flags flags,size_t size,void *host_ptr)
474{
475 return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr,
476 (cl_int *) NULL));
477}
478
479MagickPrivate void ReleaseOpenCLKernel(cl_kernel kernel)
480{
481 (void) openCL_library->clReleaseKernel(kernel);
482}
483
484MagickPrivate void ReleaseOpenCLMemObject(cl_mem memobj)
485{
486 (void) openCL_library->clReleaseMemObject(memobj);
487}
488
489MagickPrivate void RetainOpenCLMemObject(cl_mem memobj)
490{
491 (void) openCL_library->clRetainMemObject(memobj);
492}
493
494MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,size_t arg_index,
495 size_t arg_size,const void *arg_value)
496{
497 return(openCL_library->clSetKernelArg(kernel,(cl_uint) arg_index,arg_size,
498 arg_value));
499}
500
501/*
502%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
503% %
504% %
505% %
506+ A c q u i r e M a g i c k C L C a c h e I n f o %
507% %
508% %
509% %
510%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
511%
512% AcquireMagickCLCacheInfo() acquires an OpenCL cache info structure.
513%
514% The format of the AcquireMagickCLCacheInfo method is:
515%
516% MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
517% Quantum *pixels,const MagickSizeType length)
518%
519% A description of each parameter follows:
520%
521% o device: the OpenCL device.
522%
523% o pixels: the pixel buffer of the image.
524%
525% o length: the length of the pixel buffer.
526%
527*/
528
529MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
530 Quantum *pixels,const MagickSizeType length)
531{
532 cl_int
533 status;
534
535 MagickCLCacheInfo
536 info;
537
538 info=(MagickCLCacheInfo) AcquireCriticalMemory(sizeof(*info));
539 (void) memset(info,0,sizeof(*info));
540 LockSemaphoreInfo(openCL_lock);
541 device->requested++;
542 UnlockSemaphoreInfo(openCL_lock);
543 info->device=device;
544 info->length=length;
545 info->pixels=pixels;
546 info->events_semaphore=AcquireSemaphoreInfo();
547 info->buffer=openCL_library->clCreateBuffer(device->context,
548 CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,
549 &status);
550 if (status == CL_SUCCESS)
551 return(info);
552 DestroyMagickCLCacheInfo(info);
553 return((MagickCLCacheInfo) NULL);
554}
555
556/*
557%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
558% %
559% %
560% %
561% A c q u i r e M a g i c k C L D e v i c e %
562% %
563% %
564% %
565%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
566%
567% AcquireMagickCLDevice() acquires an OpenCL device
568%
569% The format of the AcquireMagickCLDevice method is:
570%
571% MagickCLDevice AcquireMagickCLDevice()
572%
573*/
574
575static MagickCLDevice AcquireMagickCLDevice()
576{
577 MagickCLDevice
578 device;
579
580 device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device));
581 if (device != NULL)
582 {
583 (void) memset(device,0,sizeof(*device));
584 ActivateSemaphoreInfo(&device->lock);
585 device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
586 device->command_queues_index=-1;
587 device->enabled=MagickTrue;
588 }
589 return(device);
590}
591
592/*
593%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
594% %
595% %
596% %
597% A c q u i r e M a g i c k C L E n v %
598% %
599% %
600% %
601%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
602%
603% AcquireMagickCLEnv() allocates the MagickCLEnv structure
604%
605*/
606
607static MagickCLEnv AcquireMagickCLEnv(void)
608{
609 char
610 *option;
611
612 MagickCLEnv
613 clEnv;
614
615 clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv));
616 if (clEnv != (MagickCLEnv) NULL)
617 {
618 (void) memset(clEnv,0,sizeof(*clEnv));
619 ActivateSemaphoreInfo(&clEnv->lock);
620 clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
621 clEnv->enabled=MagickFalse;
622 option=GetEnvironmentValue("MAGICK_OCL_DEVICE");
623 if (option != (const char *) NULL)
624 {
625 if ((IsStringTrue(option) != MagickFalse) ||
626 (strcmp(option,"GPU") == 0) ||
627 (strcmp(option,"CPU") == 0))
628 clEnv->enabled=MagickTrue;
629 option=DestroyString(option);
630 }
631 }
632 return clEnv;
633}
634
635/*
636%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
637% %
638% %
639% %
640+ 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 %
641% %
642% %
643% %
644%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
645%
646% AcquireOpenCLCommandQueue() acquires an OpenCL command queue
647%
648% The format of the AcquireOpenCLCommandQueue method is:
649%
650% cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
651%
652% A description of each parameter follows:
653%
654% o device: the OpenCL device.
655%
656*/
657
658MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
659{
660 cl_command_queue
661 queue;
662
663 cl_command_queue_properties
664 properties;
665
666 assert(device != (MagickCLDevice) NULL);
667 LockSemaphoreInfo(device->lock);
668 if ((device->profile_kernels == MagickFalse) &&
669 (device->command_queues_index >= 0))
670 {
671 queue=device->command_queues[device->command_queues_index--];
672 UnlockSemaphoreInfo(device->lock);
673 }
674 else
675 {
676 UnlockSemaphoreInfo(device->lock);
677 properties=0;
678 if (device->profile_kernels != MagickFalse)
679 properties=CL_QUEUE_PROFILING_ENABLE;
680 queue=openCL_library->clCreateCommandQueue(device->context,
681 device->deviceID,properties,(cl_int *) NULL);
682 }
683 return(queue);
684}
685
686/*
687%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
688% %
689% %
690% %
691+ A c q u i r e O p e n C L K e r n e l %
692% %
693% %
694% %
695%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
696%
697% AcquireOpenCLKernel() acquires an OpenCL kernel
698%
699% The format of the AcquireOpenCLKernel method is:
700%
701% cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
702% MagickOpenCLProgram program, const char* kernelName)
703%
704% A description of each parameter follows:
705%
706% o clEnv: the OpenCL environment.
707%
708% o program: the OpenCL program module that the kernel belongs to.
709%
710% o kernelName: the name of the kernel
711%
712*/
713
714MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
715 const char *kernel_name)
716{
717 cl_kernel
718 kernel;
719
720 assert(device != (MagickCLDevice) NULL);
721 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Using kernel: %s",
722 kernel_name);
723 kernel=openCL_library->clCreateKernel(device->program,kernel_name,
724 (cl_int *) NULL);
725 return(kernel);
726}
727
728/*
729%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
730% %
731% %
732% %
733% A u t o S e l e c t O p e n C L D e v i c e s %
734% %
735% %
736% %
737%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
738%
739% AutoSelectOpenCLDevices() determines the best device based on the
740% information from the micro-benchmark.
741%
742% The format of the AutoSelectOpenCLDevices method is:
743%
744% void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception)
745%
746% A description of each parameter follows:
747%
748% o clEnv: the OpenCL environment.
749%
750% o exception: return any errors or warnings in this structure.
751%
752*/
753
754#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
755static MagickCLDeviceBenchmark* RelinquishDeviceBenchmark(
756 MagickCLDeviceBenchmark *device_benchmark)
757{
758 if (device_benchmark == (MagickCLDeviceBenchmark*) NULL)
759 return((MagickCLDeviceBenchmark *) NULL);
760
761 device_benchmark->platform_name=(char *) RelinquishMagickMemory(
762 device_benchmark->platform_name);
763 device_benchmark->vendor_name=(char *) RelinquishMagickMemory(
764 device_benchmark->vendor_name);
765 device_benchmark->name=(char *) RelinquishMagickMemory(
766 device_benchmark->name);
767 device_benchmark->version=(char *) RelinquishMagickMemory(
768 device_benchmark->version);
769 return((MagickCLDeviceBenchmark *) RelinquishMagickMemory(
770 device_benchmark));
771}
772
773static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
774{
775 char
776 keyword[MagickPathExtent],
777 *token;
778
779 const char
780 *q;
781
782 MagickCLDeviceBenchmark
783 *device_benchmark;
784
785 size_t
786 i,
787 extent;
788
789 if (xml == (char *) NULL)
790 return;
791 device_benchmark=(MagickCLDeviceBenchmark *) NULL;
792 token=AcquireString(xml);
793 extent=strlen(token)+MagickPathExtent;
794 for (q=(char *) xml; *q != '\0'; )
795 {
796 /*
797 Interpret XML.
798 */
799 if (SkipXMLComment(&q) == MagickFalse)
800 break;
801 (void) GetNextToken(q,&q,extent,token);
802 if (*token == '\0')
803 break;
804 (void) CopyMagickString(keyword,token,MagickPathExtent);
805 if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
806 {
807 if (SkipXMLDocType(&q) == MagickFalse)
808 break;
809 continue;
810 }
811 if (LocaleCompare(keyword,"<device") == 0)
812 {
813 /*
814 Device element.
815 */
816 device_benchmark=(MagickCLDeviceBenchmark *) AcquireQuantumMemory(1,
817 sizeof(*device_benchmark));
818 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
819 break;
820 (void) memset(device_benchmark,0,sizeof(*device_benchmark));
821 device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
822 continue;
823 }
824 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
825 continue;
826 if (LocaleCompare(keyword,"/>") == 0)
827 {
828 if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
829 {
830 if (LocaleCompare(device_benchmark->name,"CPU") == 0)
831 clEnv->cpu_score=device_benchmark->score;
832 else
833 {
834 MagickCLDevice
835 device;
836
837 /*
838 Set the score for all devices that match this device.
839 */
840 for (i = 0; i < clEnv->number_devices; i++)
841 {
842 device=clEnv->devices[i];
843 if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
844 device->score=device_benchmark->score;
845 }
846 }
847 }
848 device_benchmark=RelinquishDeviceBenchmark(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=RelinquishDeviceBenchmark(device_benchmark);
909}
910
911static MagickBooleanType CanWriteProfileToFile(const char *filename)
912{
913 FILE
914 *profileFile;
915
916 profileFile=fopen_utf8(filename,"ab");
917
918 if (profileFile == (FILE *) NULL)
919 {
920 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
921 "Unable to save profile to: \"%s\"",filename);
922 return(MagickFalse);
923 }
924
925 fclose(profileFile);
926 return(MagickTrue);
927}
928#endif
929
930static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
931{
932#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
933 char
934 filename[MagickPathExtent];
935
936 StringInfo
937 *option;
938
939 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
940 GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
941
942 /*
943 We don't run the benchmark when we can not write out a device profile. The
944 first GPU device will be used.
945 */
946 if (CanWriteProfileToFile(filename) == MagickFalse)
947#endif
948 {
949 size_t
950 i;
951
952 for (i = 0; i < clEnv->number_devices; i++)
953 clEnv->devices[i]->score=1.0;
954
955 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
956 return(MagickFalse);
957 }
958#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
959 option=ConfigureFileToStringInfo(filename);
960 LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(option));
961 option=DestroyStringInfo(option);
962 return(MagickTrue);
963#endif
964}
965
966static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
967{
968 char
969 *option;
970
971 double
972 best_score;
973
974 MagickBooleanType
975 benchmark;
976
977 size_t
978 i;
979
980 option=GetEnvironmentValue("MAGICK_OCL_DEVICE");
981 if (option != (const char *) NULL)
982 {
983 if (strcmp(option,"GPU") == 0)
984 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
985 else if (strcmp(option,"CPU") == 0)
986 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
987 option=DestroyString(option);
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 (void) 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_%.17g.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,(size_t) info->length,event_count,
1517 events,
1518 (cl_event *) NULL,(cl_int *) NULL);
1519 assert(pixels == info->pixels);
1520 ReleaseOpenCLCommandQueue(info->device,queue);
1521 events=(cl_event *) RelinquishMagickMemory(events);
1522 }
1523 return(RelinquishMagickCLCacheInfo(info,MagickFalse));
1524}
1525
1526/*
1527%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1528% %
1529% %
1530% %
1531+ D u m p O p e n C L P r o f i l e D a t a %
1532% %
1533% %
1534% %
1535%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1536%
1537% DumpOpenCLProfileData() dumps the kernel profile data.
1538%
1539% The format of the DumpProfileData method is:
1540%
1541% void DumpProfileData()
1542%
1543*/
1544
1545MagickPrivate void DumpOpenCLProfileData()
1546{
1547#define OpenCLLog(message) \
1548 fwrite(message,sizeof(char),strlen(message),log); \
1549 fwrite("\n",sizeof(char),1,log);
1550
1551 char
1552 buf[4096],
1553 filename[MagickPathExtent],
1554 indent[160];
1555
1556 FILE
1557 *log;
1558
1559 size_t
1560 i,
1561 j;
1562
1563 if (default_CLEnv == (MagickCLEnv) NULL)
1564 return;
1565
1566 for (i = 0; i < default_CLEnv->number_devices; i++)
1567 if (default_CLEnv->devices[i]->profile_kernels != MagickFalse)
1568 break;
1569 if (i == default_CLEnv->number_devices)
1570 return;
1571
1572 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1573 GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
1574
1575 log=fopen_utf8(filename,"wb");
1576 if (log == (FILE *) NULL)
1577 return;
1578 for (i = 0; i < default_CLEnv->number_devices; i++)
1579 {
1580 MagickCLDevice
1581 device;
1582
1583 device=default_CLEnv->devices[i];
1584 if ((device->profile_kernels == MagickFalse) ||
1585 (device->profile_records == (KernelProfileRecord *) NULL))
1586 continue;
1587
1588 OpenCLLog("====================================================");
1589 fprintf(log,"Device: %s\n",device->name);
1590 fprintf(log,"Version: %s\n",device->version);
1591 OpenCLLog("====================================================");
1592 OpenCLLog(" average calls min max");
1593 OpenCLLog(" ------- ----- --- ---");
1594 j=0;
1595 while (device->profile_records[j] != (KernelProfileRecord) NULL)
1596 {
1597 KernelProfileRecord
1598 profile;
1599
1600 profile=device->profile_records[j];
1601 (void) CopyMagickString(indent," ",
1602 sizeof(indent));
1603 (void) CopyMagickString(indent,profile->kernel_name,MagickMin(strlen(
1604 profile->kernel_name),strlen(indent)));
1605 (void) FormatLocaleString(buf,sizeof(buf),"%s %7d %7d %7d %7d",indent,
1606 (int) (profile->total/profile->count),(int) profile->count,
1607 (int) profile->min,(int) profile->max);
1608 OpenCLLog(buf);
1609 j++;
1610 }
1611 OpenCLLog("====================================================");
1612 fwrite("\n\n",sizeof(char),2,log);
1613 }
1614 fclose(log);
1615}
1616/*
1617%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1618% %
1619% %
1620% %
1621+ E n q u e u e O p e n C L K e r n e l %
1622% %
1623% %
1624% %
1625%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1626%
1627% EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
1628% events with the images.
1629%
1630% The format of the EnqueueOpenCLKernel method is:
1631%
1632% MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
1633% const size_t *global_work_offset,const size_t *global_work_size,
1634% const size_t *local_work_size,const Image *input_image,
1635% const Image *output_image,ExceptionInfo *exception)
1636%
1637% A description of each parameter follows:
1638%
1639% o kernel: the OpenCL kernel.
1640%
1641% o work_dim: the number of dimensions used to specify the global work-items
1642% and work-items in the work-group.
1643%
1644% o offset: can be used to specify an array of work_dim unsigned values
1645% that describe the offset used to calculate the global ID of a
1646% work-item.
1647%
1648% o gsize: points to an array of work_dim unsigned values that describe the
1649% number of global work-items in work_dim dimensions that will
1650% execute the kernel function.
1651%
1652% o lsize: points to an array of work_dim unsigned values that describe the
1653% number of work-items that make up a work-group that will execute
1654% the kernel specified by kernel.
1655%
1656% o input_image: the input image of the operation.
1657%
1658% o output_image: the output or secondary image of the operation.
1659%
1660% o exception: return any errors or warnings in this structure.
1661%
1662*/
1663
1664static MagickBooleanType RegisterCacheEvent(MagickCLCacheInfo info,
1665 cl_event event)
1666{
1667 assert(info != (MagickCLCacheInfo) NULL);
1668 assert(event != (cl_event) NULL);
1669 if (openCL_library->clRetainEvent(event) != CL_SUCCESS)
1670 {
1671 openCL_library->clWaitForEvents(1,&event);
1672 return(MagickFalse);
1673 }
1674 LockSemaphoreInfo(info->events_semaphore);
1675 if (info->events == (cl_event *) NULL)
1676 {
1677 info->events=(cl_event *) AcquireMagickMemory(sizeof(*info->events));
1678 info->event_count=1;
1679 }
1680 else
1681 info->events=(cl_event *) ResizeQuantumMemory(info->events,
1682 ++info->event_count,sizeof(*info->events));
1683 if (info->events == (cl_event *) NULL)
1684 {
1685 UnlockSemaphoreInfo(info->events_semaphore);
1686 return(MagickFalse);
1687 }
1688 info->events[info->event_count-1]=event;
1689 UnlockSemaphoreInfo(info->events_semaphore);
1690 return(MagickTrue);
1691}
1692
1693MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1694 cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
1695 const size_t *lsize,const Image *input_image,const Image *output_image,
1696 MagickBooleanType flush,ExceptionInfo *exception)
1697{
1698 CacheInfo
1699 *output_info,
1700 *input_info;
1701
1702 cl_event
1703 event,
1704 *events;
1705
1706 cl_int
1707 status;
1708
1709 cl_uint
1710 event_count;
1711
1712 assert(input_image != (const Image *) NULL);
1713 input_info=(CacheInfo *) input_image->cache;
1714 assert(input_info != (CacheInfo *) NULL);
1715 assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1716 output_info=(CacheInfo *) NULL;
1717 if (output_image == (const Image *) NULL)
1718 events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
1719 &event_count);
1720 else
1721 {
1722 output_info=(CacheInfo *) output_image->cache;
1723 assert(output_info != (CacheInfo *) NULL);
1724 assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1725 events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
1726 &event_count);
1727 }
1728 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1729 gsize,lsize,event_count,events,&event);
1730 /* This can fail due to memory issues and calling clFinish might help. */
1731 if ((status != CL_SUCCESS) && (event_count > 0))
1732 {
1733 openCL_library->clFinish(queue);
1734 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
1735 offset,gsize,lsize,event_count,events,&event);
1736 }
1737 events=(cl_event *) RelinquishMagickMemory(events);
1738 if (status != CL_SUCCESS)
1739 {
1740 (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1741 GetMagickModule(),ResourceLimitWarning,
1742 "clEnqueueNDRangeKernel failed.","'%s'",".");
1743 return(MagickFalse);
1744 }
1745 if (flush != MagickFalse)
1746 openCL_library->clFlush(queue);
1747 if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1748 {
1749 if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
1750 {
1751 if (output_info != (CacheInfo *) NULL)
1752 (void) RegisterCacheEvent(output_info->opencl,event);
1753 }
1754 }
1755 openCL_library->clReleaseEvent(event);
1756 return(MagickTrue);
1757}
1758
1759/*
1760%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1761% %
1762% %
1763% %
1764+ G e t C u r r e n t O p e n C L E n v %
1765% %
1766% %
1767% %
1768%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1769%
1770% GetCurrentOpenCLEnv() returns the current OpenCL env
1771%
1772% The format of the GetCurrentOpenCLEnv method is:
1773%
1774% MagickCLEnv GetCurrentOpenCLEnv()
1775%
1776*/
1777
1778MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
1779{
1780 if (default_CLEnv != (MagickCLEnv) NULL)
1781 {
1782 if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1783 (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1784 return((MagickCLEnv) NULL);
1785 else
1786 return(default_CLEnv);
1787 }
1788
1789 if (GetOpenCLCacheDirectory() == (char *) NULL)
1790 return((MagickCLEnv) NULL);
1791
1792 if (openCL_lock == (SemaphoreInfo *) NULL)
1793 ActivateSemaphoreInfo(&openCL_lock);
1794
1795 LockSemaphoreInfo(openCL_lock);
1796 if (default_CLEnv == (MagickCLEnv) NULL)
1797 default_CLEnv=AcquireMagickCLEnv();
1798 UnlockSemaphoreInfo(openCL_lock);
1799
1800 return(default_CLEnv);
1801}
1802
1803/*
1804%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1805% %
1806% %
1807% %
1808% 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 %
1809% %
1810% %
1811% %
1812%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1813%
1814% GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
1815% device. The score is determined by the duration of the micro benchmark so
1816% that means a lower score is better than a higher score.
1817%
1818% The format of the GetOpenCLDeviceBenchmarkScore method is:
1819%
1820% double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
1821%
1822% A description of each parameter follows:
1823%
1824% o device: the OpenCL device.
1825*/
1826
1827MagickExport double GetOpenCLDeviceBenchmarkScore(
1828 const MagickCLDevice device)
1829{
1830 if (device == (MagickCLDevice) NULL)
1831 return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1832 return(device->score);
1833}
1834
1835/*
1836%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1837% %
1838% %
1839% %
1840% G e t O p e n C L D e v i c e E n a b l e d %
1841% %
1842% %
1843% %
1844%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1845%
1846% GetOpenCLDeviceEnabled() returns true if the device is enabled.
1847%
1848% The format of the GetOpenCLDeviceEnabled method is:
1849%
1850% MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
1851%
1852% A description of each parameter follows:
1853%
1854% o device: the OpenCL device.
1855*/
1856
1857MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
1858 const MagickCLDevice device)
1859{
1860 if (device == (MagickCLDevice) NULL)
1861 return(MagickFalse);
1862 return(device->enabled);
1863}
1864
1865/*
1866%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1867% %
1868% %
1869% %
1870% G e t O p e n C L D e v i c e N a m e %
1871% %
1872% %
1873% %
1874%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1875%
1876% GetOpenCLDeviceName() returns the name of the device.
1877%
1878% The format of the GetOpenCLDeviceName method is:
1879%
1880% const char *GetOpenCLDeviceName(const MagickCLDevice device)
1881%
1882% A description of each parameter follows:
1883%
1884% o device: the OpenCL device.
1885*/
1886
1887MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
1888{
1889 if (device == (MagickCLDevice) NULL)
1890 return((const char *) NULL);
1891 return(device->name);
1892}
1893
1894/*
1895%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1896% %
1897% %
1898% %
1899% 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 %
1900% %
1901% %
1902% %
1903%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1904%
1905% GetOpenCLDeviceVendorName() returns the vendor name of the device.
1906%
1907% The format of the GetOpenCLDeviceVendorName method is:
1908%
1909% const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1910%
1911% A description of each parameter follows:
1912%
1913% o device: the OpenCL device.
1914*/
1915
1916MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1917{
1918 if (device == (MagickCLDevice) NULL)
1919 return((const char *) NULL);
1920 return(device->vendor_name);
1921}
1922
1923/*
1924%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1925% %
1926% %
1927% %
1928% G e t O p e n C L D e v i c e s %
1929% %
1930% %
1931% %
1932%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1933%
1934% GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
1935% value of length to the number of devices that are available.
1936%
1937% The format of the GetOpenCLDevices method is:
1938%
1939% const MagickCLDevice *GetOpenCLDevices(size_t *length,
1940% ExceptionInfo *exception)
1941%
1942% A description of each parameter follows:
1943%
1944% o length: the number of device.
1945%
1946% o exception: return any errors or warnings in this structure.
1947%
1948*/
1949
1950MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
1951 ExceptionInfo *exception)
1952{
1953 MagickCLEnv
1954 clEnv;
1955
1956 clEnv=GetCurrentOpenCLEnv();
1957 if (clEnv == (MagickCLEnv) NULL)
1958 {
1959 if (length != (size_t *) NULL)
1960 *length=0;
1961 return((MagickCLDevice *) NULL);
1962 }
1963 InitializeOpenCL(clEnv,exception);
1964 if (length != (size_t *) NULL)
1965 *length=clEnv->number_devices;
1966 return(clEnv->devices);
1967}
1968
1969/*
1970%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1971% %
1972% %
1973% %
1974% G e t O p e n C L D e v i c e T y p e %
1975% %
1976% %
1977% %
1978%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1979%
1980% GetOpenCLDeviceType() returns the type of the device.
1981%
1982% The format of the GetOpenCLDeviceType method is:
1983%
1984% MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
1985%
1986% A description of each parameter follows:
1987%
1988% o device: the OpenCL device.
1989*/
1990
1991MagickExport MagickCLDeviceType GetOpenCLDeviceType(
1992 const MagickCLDevice device)
1993{
1994 if (device == (MagickCLDevice) NULL)
1995 return(UndefinedCLDeviceType);
1996 if (device->type == CL_DEVICE_TYPE_GPU)
1997 return(GpuCLDeviceType);
1998 if (device->type == CL_DEVICE_TYPE_CPU)
1999 return(CpuCLDeviceType);
2000 return(UndefinedCLDeviceType);
2001}
2002
2003/*
2004%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2005% %
2006% %
2007% %
2008% G e t O p e n C L D e v i c e V e r s i o n %
2009% %
2010% %
2011% %
2012%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2013%
2014% GetOpenCLDeviceVersion() returns the version of the device.
2015%
2016% The format of the GetOpenCLDeviceName method is:
2017%
2018% const char *GetOpenCLDeviceVersion(MagickCLDevice device)
2019%
2020% A description of each parameter follows:
2021%
2022% o device: the OpenCL device.
2023*/
2024
2025MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
2026{
2027 if (device == (MagickCLDevice) NULL)
2028 return((const char *) NULL);
2029 return(device->version);
2030}
2031
2032/*
2033%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2034% %
2035% %
2036% %
2037% G e t O p e n C L E n a b l e d %
2038% %
2039% %
2040% %
2041%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2042%
2043% GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
2044%
2045% The format of the GetOpenCLEnabled method is:
2046%
2047% MagickBooleanType GetOpenCLEnabled()
2048%
2049*/
2050
2051MagickExport MagickBooleanType GetOpenCLEnabled(void)
2052{
2053 MagickCLEnv
2054 clEnv;
2055
2056 clEnv=GetCurrentOpenCLEnv();
2057 if (clEnv == (MagickCLEnv) NULL)
2058 return(MagickFalse);
2059 return(clEnv->enabled);
2060}
2061
2062/*
2063%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2064% %
2065% %
2066% %
2067% 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 %
2068% %
2069% %
2070% %
2071%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2072%
2073% GetOpenCLKernelProfileRecords() returns the profile records for the
2074% specified device and sets length to the number of profile records.
2075%
2076% The format of the GetOpenCLKernelProfileRecords method is:
2077%
2078% const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
2079%
2080% A description of each parameter follows:
2081%
2082% o length: the number of profiles records.
2083*/
2084
2085MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
2086 const MagickCLDevice device,size_t *length)
2087{
2088 if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
2089 (KernelProfileRecord *) NULL))
2090 {
2091 if (length != (size_t *) NULL)
2092 *length=0;
2093 return((const KernelProfileRecord *) NULL);
2094 }
2095 if (length != (size_t *) NULL)
2096 {
2097 *length=0;
2098 LockSemaphoreInfo(device->lock);
2099 while (device->profile_records[*length] != (KernelProfileRecord) NULL)
2100 *length=*length+1;
2101 UnlockSemaphoreInfo(device->lock);
2102 }
2103 return(device->profile_records);
2104}
2105
2106/*
2107%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2108% %
2109% %
2110% %
2111% H a s O p e n C L D e v i c e s %
2112% %
2113% %
2114% %
2115%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2116%
2117% HasOpenCLDevices() checks if the OpenCL environment has devices that are
2118% enabled and compiles the kernel for the device when necessary. False will be
2119% returned if no enabled devices could be found
2120%
2121% The format of the HasOpenCLDevices method is:
2122%
2123% MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2124% ExceptionInfo exception)
2125%
2126% A description of each parameter follows:
2127%
2128% o clEnv: the OpenCL environment.
2129%
2130% o exception: return any errors or warnings in this structure.
2131%
2132*/
2133
2134static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2135 ExceptionInfo *exception)
2136{
2137 char
2138 *accelerateKernelsBuffer,
2139 options[MagickPathExtent];
2140
2141 MagickBooleanType
2142 status;
2143
2144 size_t
2145 i;
2146
2147 size_t
2148 signature;
2149
2150 /* Check if there are enabled devices */
2151 for (i = 0; i < clEnv->number_devices; i++)
2152 {
2153 if ((clEnv->devices[i]->enabled != MagickFalse))
2154 break;
2155 }
2156 if (i == clEnv->number_devices)
2157 return(MagickFalse);
2158
2159 /* Check if we need to compile a kernel for one of the devices */
2160 status=MagickTrue;
2161 for (i = 0; i < clEnv->number_devices; i++)
2162 {
2163 if ((clEnv->devices[i]->enabled != MagickFalse) &&
2164 (clEnv->devices[i]->program == (cl_program) NULL))
2165 {
2166 status=MagickFalse;
2167 break;
2168 }
2169 }
2170 if (status != MagickFalse)
2171 return(MagickTrue);
2172
2173 /* Get additional options */
2174 (void) FormatLocaleString(options,MagickPathExtent,CLOptions,
2175 (float)QuantumRange,(float)CLCharQuantumScale,(float)MagickEpsilon,
2176 (float)MagickPI,(unsigned int)MaxMap,(unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2177
2178 signature=StringSignature(options);
2179 accelerateKernelsBuffer=(char*) AcquireQuantumMemory(1,
2180 strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2181 if (accelerateKernelsBuffer == (char*) NULL)
2182 return(MagickFalse);
2183 (void) FormatLocaleString(accelerateKernelsBuffer,strlen(accelerateKernels)+
2184 strlen(accelerateKernels2)+1,"%s%s",accelerateKernels,accelerateKernels2);
2185 signature^=StringSignature(accelerateKernelsBuffer);
2186
2187 status=MagickTrue;
2188 for (i = 0; i < clEnv->number_devices; i++)
2189 {
2190 MagickCLDevice
2191 device;
2192
2193 size_t
2194 device_signature;
2195
2196 device=clEnv->devices[i];
2197 if ((device->enabled == MagickFalse) ||
2198 (device->program != (cl_program) NULL))
2199 continue;
2200
2201 LockSemaphoreInfo(device->lock);
2202 if (device->program != (cl_program) NULL)
2203 {
2204 UnlockSemaphoreInfo(device->lock);
2205 continue;
2206 }
2207 device_signature=signature;
2208 device_signature^=StringSignature(device->platform_name);
2209 status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2210 device_signature,exception);
2211 UnlockSemaphoreInfo(device->lock);
2212 if (status == MagickFalse)
2213 break;
2214 }
2215 accelerateKernelsBuffer=(char *) RelinquishMagickMemory(
2216 accelerateKernelsBuffer);
2217 return(status);
2218}
2219
2220/*
2221%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2222% %
2223% %
2224% %
2225+ I n i t i a l i z e O p e n C L %
2226% %
2227% %
2228% %
2229%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2230%
2231% InitializeOpenCL() is used to initialize the OpenCL environment. This method
2232% makes sure the devices are properly initialized and benchmarked.
2233%
2234% The format of the InitializeOpenCL method is:
2235%
2236% MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2237%
2238% A description of each parameter follows:
2239%
2240% o exception: return any errors or warnings in this structure.
2241%
2242*/
2243
2244static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2245{
2246 char
2247 version[MagickPathExtent];
2248
2249 cl_uint
2250 num;
2251
2252 if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2253 MagickPathExtent,version,NULL) != CL_SUCCESS)
2254 return(0);
2255 if (strncmp(version,"OpenCL 1.0 ",11) == 0)
2256 return(0);
2257 if (clEnv->library->clGetDeviceIDs(platform,
2258 CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2259 return(0);
2260 return(num);
2261}
2262
2263static inline char *GetOpenCLPlatformString(cl_platform_id platform,
2264 cl_platform_info param_name)
2265{
2266 char
2267 *value;
2268
2269 size_t
2270 length;
2271
2272 openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
2273 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2274 openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
2275 return(value);
2276}
2277
2278static inline char *GetOpenCLDeviceString(cl_device_id device,
2279 cl_device_info param_name)
2280{
2281 char
2282 *value;
2283
2284 size_t
2285 length;
2286
2287 openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
2288 value=(char *) AcquireCriticalMemory(length*sizeof(*value));
2289 openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
2290 return(value);
2291}
2292
2293static void LoadOpenCLDevices(MagickCLEnv clEnv)
2294{
2295 cl_context_properties
2296 properties[3];
2297
2298 cl_device_id
2299 *devices;
2300
2301 cl_int
2302 status;
2303
2304 cl_platform_id
2305 *platforms;
2306
2307 cl_uint
2308 i,
2309 j,
2310 next,
2311 number_devices,
2312 number_platforms;
2313
2314 number_platforms=0;
2315 if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2316 return;
2317 if (number_platforms == 0)
2318 return;
2319 platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
2320 sizeof(cl_platform_id));
2321 if (platforms == (cl_platform_id *) NULL)
2322 return;
2323 if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2324 {
2325 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2326 return;
2327 }
2328 for (i = 0; i < number_platforms; i++)
2329 {
2330 number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2331 if (number_devices == 0)
2332 platforms[i]=(cl_platform_id) NULL;
2333 else
2334 clEnv->number_devices+=number_devices;
2335 }
2336 if (clEnv->number_devices == 0)
2337 {
2338 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2339 return;
2340 }
2341 clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2342 sizeof(MagickCLDevice));
2343 if (clEnv->devices == (MagickCLDevice *) NULL)
2344 {
2345 RelinquishMagickCLDevices(clEnv);
2346 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2347 return;
2348 }
2349 (void) memset(clEnv->devices,0,clEnv->number_devices*sizeof(MagickCLDevice));
2350 devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2351 sizeof(cl_device_id));
2352 if (devices == (cl_device_id *) NULL)
2353 {
2354 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2355 RelinquishMagickCLDevices(clEnv);
2356 return;
2357 }
2358 (void) memset(devices,0,clEnv->number_devices*sizeof(cl_device_id));
2359 clEnv->number_contexts=(size_t) number_platforms;
2360 clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2361 sizeof(cl_context));
2362 if (clEnv->contexts == (cl_context *) NULL)
2363 {
2364 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2365 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2366 RelinquishMagickCLDevices(clEnv);
2367 return;
2368 }
2369 (void) memset(clEnv->contexts,0,clEnv->number_contexts*sizeof(cl_context));
2370 next=0;
2371 for (i = 0; i < number_platforms; i++)
2372 {
2373 if (platforms[i] == (cl_platform_id) NULL)
2374 continue;
2375
2376 status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2377 CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2378 if (status != CL_SUCCESS)
2379 continue;
2380
2381 properties[0]=CL_CONTEXT_PLATFORM;
2382 properties[1]=(cl_context_properties) platforms[i];
2383 properties[2]=0;
2384 clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2385 devices,NULL,NULL,&status);
2386 if (status != CL_SUCCESS)
2387 continue;
2388
2389 for (j = 0; j < number_devices; j++,next++)
2390 {
2391 MagickCLDevice
2392 device;
2393
2394 device=AcquireMagickCLDevice();
2395 if (device == (MagickCLDevice) NULL)
2396 break;
2397
2398 device->context=clEnv->contexts[i];
2399 device->deviceID=devices[j];
2400
2401 device->platform_name=GetOpenCLPlatformString(platforms[i],
2402 CL_PLATFORM_NAME);
2403
2404 device->vendor_name=GetOpenCLPlatformString(platforms[i],
2405 CL_PLATFORM_VENDOR);
2406
2407 device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);
2408
2409 device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);
2410
2411 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2412 sizeof(cl_uint),&device->max_clock_frequency,NULL);
2413
2414 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2415 sizeof(cl_uint),&device->max_compute_units,NULL);
2416
2417 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2418 sizeof(cl_device_type),&device->type,NULL);
2419
2420 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2421 sizeof(cl_ulong),&device->local_memory_size,NULL);
2422
2423 clEnv->devices[next]=device;
2424 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
2425 "Found device: %s (%s)",device->name,device->platform_name);
2426 }
2427 }
2428 if (next != clEnv->number_devices)
2429 RelinquishMagickCLDevices(clEnv);
2430 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2431 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2432}
2433
2434MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2435 ExceptionInfo *exception)
2436{
2437 LockSemaphoreInfo(clEnv->lock);
2438 if (clEnv->initialized != MagickFalse)
2439 {
2440 UnlockSemaphoreInfo(clEnv->lock);
2441 return(HasOpenCLDevices(clEnv,exception));
2442 }
2443 if (LoadOpenCLLibrary() != MagickFalse)
2444 {
2445 clEnv->library=openCL_library;
2446 LoadOpenCLDevices(clEnv);
2447 if (clEnv->number_devices > 0)
2448 AutoSelectOpenCLDevices(clEnv);
2449 }
2450 clEnv->initialized=MagickTrue;
2451 UnlockSemaphoreInfo(clEnv->lock);
2452 return(HasOpenCLDevices(clEnv,exception));
2453}
2454
2455/*
2456%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2457% %
2458% %
2459% %
2460% L o a d O p e n C L L i b r a r y %
2461% %
2462% %
2463% %
2464%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2465%
2466% LoadOpenCLLibrary() load and binds the OpenCL library.
2467%
2468% The format of the LoadOpenCLLibrary method is:
2469%
2470% MagickBooleanType LoadOpenCLLibrary(void)
2471%
2472*/
2473
2474void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2475{
2476 if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2477 return (void *) NULL;
2478 return lt_dlsym(library,functionName);
2479}
2480
2481static MagickBooleanType BindOpenCLFunctions()
2482{
2483#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
2484#define BIND(X) openCL_library->X= &X;
2485#else
2486 (void) memset(openCL_library,0,sizeof(MagickLibrary));
2487#ifdef MAGICKCORE_WINDOWS_SUPPORT
2488 openCL_library->library=(void *)lt_dlopen("OpenCL.dll");
2489#else
2490 openCL_library->library=(void *)lt_dlopen("libOpenCL.so");
2491#endif
2492#define BIND(X) \
2493 if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2494 return(MagickFalse);
2495#endif
2496
2497 if (openCL_library->library == (void*) NULL)
2498 return(MagickFalse);
2499
2500 BIND(clGetPlatformIDs);
2501 BIND(clGetPlatformInfo);
2502
2503 BIND(clGetDeviceIDs);
2504 BIND(clGetDeviceInfo);
2505
2506 BIND(clCreateBuffer);
2507 BIND(clReleaseMemObject);
2508 BIND(clRetainMemObject);
2509
2510 BIND(clCreateContext);
2511 BIND(clReleaseContext);
2512
2513 BIND(clCreateCommandQueue);
2514 BIND(clReleaseCommandQueue);
2515 BIND(clFlush);
2516 BIND(clFinish);
2517
2518 BIND(clCreateProgramWithSource);
2519 BIND(clCreateProgramWithBinary);
2520 BIND(clReleaseProgram);
2521 BIND(clBuildProgram);
2522 BIND(clGetProgramBuildInfo);
2523 BIND(clGetProgramInfo);
2524
2525 BIND(clCreateKernel);
2526 BIND(clReleaseKernel);
2527 BIND(clSetKernelArg);
2528 BIND(clGetKernelInfo);
2529
2530 BIND(clEnqueueReadBuffer);
2531 BIND(clEnqueueMapBuffer);
2532 BIND(clEnqueueUnmapMemObject);
2533 BIND(clEnqueueNDRangeKernel);
2534
2535 BIND(clGetEventInfo);
2536 BIND(clWaitForEvents);
2537 BIND(clReleaseEvent);
2538 BIND(clRetainEvent);
2539 BIND(clSetEventCallback);
2540
2541 BIND(clGetEventProfilingInfo);
2542
2543 return(MagickTrue);
2544}
2545
2546static MagickBooleanType LoadOpenCLLibrary(void)
2547{
2548 openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2549 if (openCL_library == (MagickLibrary *) NULL)
2550 return(MagickFalse);
2551
2552 if (BindOpenCLFunctions() == MagickFalse)
2553 {
2554 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2555 return(MagickFalse);
2556 }
2557
2558 return(MagickTrue);
2559}
2560
2561/*
2562%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2563% %
2564% %
2565% %
2566+ O p e n C L T e r m i n u s %
2567% %
2568% %
2569% %
2570%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2571%
2572% OpenCLTerminus() destroys the OpenCL component.
2573%
2574% The format of the OpenCLTerminus method is:
2575%
2576% OpenCLTerminus(void)
2577%
2578*/
2579
2580MagickPrivate void OpenCLTerminus()
2581{
2582 DumpOpenCLProfileData();
2583 if (cache_directory != (char *) NULL)
2584 cache_directory=DestroyString(cache_directory);
2585 if (cache_directory_lock != (SemaphoreInfo *) NULL)
2586 RelinquishSemaphoreInfo(&cache_directory_lock);
2587 if (default_CLEnv != (MagickCLEnv) NULL)
2588 default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2589 if (openCL_lock != (SemaphoreInfo *) NULL)
2590 RelinquishSemaphoreInfo(&openCL_lock);
2591 if (openCL_library != (MagickLibrary *) NULL)
2592 {
2593 if (openCL_library->library != (void *) NULL)
2594 (void) lt_dlclose(openCL_library->library);
2595 openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2596 }
2597}
2598
2599/*
2600%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2601% %
2602% %
2603% %
2604+ 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 %
2605% %
2606% %
2607% %
2608%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2609%
2610% OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2611% configuration file. If an error occurs, MagickFalse is returned
2612% otherwise MagickTrue.
2613%
2614% The format of the OpenCLThrowMagickException method is:
2615%
2616% MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2617% const char *module,const char *function,const size_t line,
2618% const ExceptionType severity,const char *tag,const char *format,...)
2619%
2620% A description of each parameter follows:
2621%
2622% o exception: the exception info.
2623%
2624% o filename: the source module filename.
2625%
2626% o function: the function name.
2627%
2628% o line: the line number of the source module.
2629%
2630% o severity: Specifies the numeric error category.
2631%
2632% o tag: the locale tag.
2633%
2634% o format: the output format.
2635%
2636*/
2637
2638MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2639 MagickCLDevice device,ExceptionInfo *exception,const char *module,
2640 const char *function,const size_t line,const ExceptionType severity,
2641 const char *tag,const char *format,...)
2642{
2643 MagickBooleanType
2644 status;
2645
2646 assert(device != (MagickCLDevice) NULL);
2647 assert(exception != (ExceptionInfo *) NULL);
2648 assert(exception->signature == MagickCoreSignature);
2649 (void) exception;
2650 status=MagickTrue;
2651 if (severity != 0)
2652 {
2653 if (device->type == CL_DEVICE_TYPE_CPU)
2654 {
2655 /* Workaround for Intel OpenCL CPU runtime bug */
2656 /* Turn off OpenCL when a problem is detected! */
2657 if (strncmp(device->platform_name,"Intel",5) == 0)
2658 default_CLEnv->enabled=MagickFalse;
2659 }
2660 }
2661
2662#ifdef OPENCLLOG_ENABLED
2663 {
2664 va_list
2665 operands;
2666 va_start(operands,format);
2667 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2668 format,operands);
2669 va_end(operands);
2670 }
2671#else
2672 magick_unreferenced(module);
2673 magick_unreferenced(function);
2674 magick_unreferenced(line);
2675 magick_unreferenced(tag);
2676 magick_unreferenced(format);
2677#endif
2678
2679 return(status);
2680}
2681
2682/*
2683%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2684% %
2685% %
2686% %
2687+ R e c o r d P r o f i l e D a t a %
2688% %
2689% %
2690% %
2691%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2692%
2693% RecordProfileData() records profile data.
2694%
2695% The format of the RecordProfileData method is:
2696%
2697% void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2698% cl_event event)
2699%
2700% A description of each parameter follows:
2701%
2702% o device: the OpenCL device that did the operation.
2703%
2704% o event: the event that contains the profiling data.
2705%
2706*/
2707
2708MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2709 cl_kernel kernel,cl_event event)
2710{
2711 char
2712 *name;
2713
2714 cl_int
2715 status;
2716
2717 cl_ulong
2718 elapsed,
2719 end,
2720 start;
2721
2722 KernelProfileRecord
2723 profile_record;
2724
2725 size_t
2726 i,
2727 length;
2728
2729 if (device->profile_kernels == MagickFalse)
2730 return(MagickFalse);
2731 status=openCL_library->clWaitForEvents(1,&event);
2732 if (status != CL_SUCCESS)
2733 return(MagickFalse);
2734 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2735 &length);
2736 if (status != CL_SUCCESS)
2737 return(MagickTrue);
2738 name=(char *) AcquireQuantumMemory(length,sizeof(*name));
2739 if (name == (char *) NULL)
2740 return(MagickTrue);
2741 start=end=elapsed=0;
2742 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2743 name,(size_t *) NULL);
2744 status|=openCL_library->clGetEventProfilingInfo(event,
2745 CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2746 status|=openCL_library->clGetEventProfilingInfo(event,
2747 CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2748 if (status != CL_SUCCESS)
2749 {
2750 name=DestroyString(name);
2751 return(MagickTrue);
2752 }
2753 start/=1000; /* usecs */
2754 end/=1000;
2755 elapsed=end-start;
2756 LockSemaphoreInfo(device->lock);
2757 i=0;
2758 profile_record=(KernelProfileRecord) NULL;
2759 if (device->profile_records != (KernelProfileRecord *) NULL)
2760 {
2761 while (device->profile_records[i] != (KernelProfileRecord) NULL)
2762 {
2763 if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2764 {
2765 profile_record=device->profile_records[i];
2766 break;
2767 }
2768 i++;
2769 }
2770 }
2771 if (profile_record != (KernelProfileRecord) NULL)
2772 name=DestroyString(name);
2773 else
2774 {
2775 profile_record=(KernelProfileRecord) AcquireCriticalMemory(
2776 sizeof(*profile_record));
2777 (void) memset(profile_record,0,sizeof(*profile_record));
2778 profile_record->kernel_name=name;
2779 device->profile_records=(KernelProfileRecord *) ResizeQuantumMemory(
2780 device->profile_records,(i+2),sizeof(*device->profile_records));
2781 if (device->profile_records == (KernelProfileRecord *) NULL)
2782 {
2783 UnlockSemaphoreInfo(device->lock);
2784 profile_record=(KernelProfileRecord) RelinquishMagickMemory(
2785 profile_record);
2786 name=DestroyString(name);
2787 return(MagickFalse);
2788 }
2789 device->profile_records[i]=profile_record;
2790 device->profile_records[i+1]=(KernelProfileRecord) NULL;
2791 }
2792 if ((elapsed < profile_record->min) || (profile_record->count == 0))
2793 profile_record->min=(unsigned long) elapsed;
2794 if (elapsed > profile_record->max)
2795 profile_record->max=(unsigned long) elapsed;
2796 profile_record->total+=(unsigned long) elapsed;
2797 profile_record->count+=1;
2798 UnlockSemaphoreInfo(device->lock);
2799 return(MagickTrue);
2800}
2801
2802/*
2803%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2804% %
2805% %
2806% %
2807+ 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 %
2808% %
2809% %
2810% %
2811%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2812%
2813% ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2814%
2815% The format of the ReleaseOpenCLCommandQueue method is:
2816%
2817% void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2818% cl_command_queue queue)
2819%
2820% A description of each parameter follows:
2821%
2822% o device: the OpenCL device.
2823%
2824% o queue: the OpenCL queue to be released.
2825*/
2826
2827MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2828 cl_command_queue queue)
2829{
2830 if (queue == (cl_command_queue) NULL)
2831 return;
2832
2833 assert(device != (MagickCLDevice) NULL);
2834 LockSemaphoreInfo(device->lock);
2835 if ((device->profile_kernels != MagickFalse) ||
2836 (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2837 {
2838 UnlockSemaphoreInfo(device->lock);
2839 openCL_library->clFinish(queue);
2840 (void) openCL_library->clReleaseCommandQueue(queue);
2841 }
2842 else
2843 {
2844 openCL_library->clFlush(queue);
2845 device->command_queues[++device->command_queues_index]=queue;
2846 UnlockSemaphoreInfo(device->lock);
2847 }
2848}
2849
2850/*
2851%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2852% %
2853% %
2854% %
2855+ R e l e a s e M a g i c k C L D e v i c e %
2856% %
2857% %
2858% %
2859%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2860%
2861% ReleaseOpenCLDevice() returns the OpenCL device to the environment
2862%
2863% The format of the ReleaseOpenCLDevice method is:
2864%
2865% void ReleaseOpenCLDevice(MagickCLDevice device)
2866%
2867% A description of each parameter follows:
2868%
2869% o device: the OpenCL device to be released.
2870%
2871*/
2872
2873MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2874{
2875 assert(device != (MagickCLDevice) NULL);
2876 LockSemaphoreInfo(openCL_lock);
2877 device->requested--;
2878 UnlockSemaphoreInfo(openCL_lock);
2879}
2880
2881/*
2882%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2883% %
2884% %
2885% %
2886+ 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 %
2887% %
2888% %
2889% %
2890%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2891%
2892% RelinquishMagickCLCacheInfo() frees memory acquired with
2893% AcquireMagickCLCacheInfo()
2894%
2895% The format of the RelinquishMagickCLCacheInfo method is:
2896%
2897% MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2898% const MagickBooleanType relinquish_pixels)
2899%
2900% A description of each parameter follows:
2901%
2902% o info: the OpenCL cache info.
2903%
2904% o relinquish_pixels: the pixels will be relinquish when set to true.
2905%
2906*/
2907
2908static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2909 cl_event magick_unused(event),
2910 cl_int magick_unused(event_command_exec_status),void *user_data)
2911{
2912 MagickCLCacheInfo
2913 info;
2914
2915 Quantum
2916 *pixels;
2917
2918 ssize_t
2919 i;
2920
2921 magick_unreferenced(event);
2922 magick_unreferenced(event_command_exec_status);
2923 info=(MagickCLCacheInfo) user_data;
2924 for (i=(ssize_t)info->event_count-1; i >= 0; i--)
2925 {
2926 cl_int
2927 event_status;
2928
2929 cl_uint
2930 status;
2931
2932 status=openCL_library->clGetEventInfo(info->events[i],
2933 CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
2934 NULL);
2935 if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
2936 {
2937 openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
2938 &DestroyMagickCLCacheInfoAndPixels,info);
2939 return;
2940 }
2941 }
2942 pixels=info->pixels;
2943 RelinquishMagickResource(MemoryResource,info->length);
2944 DestroyMagickCLCacheInfo(info);
2945 (void) RelinquishAlignedMemory(pixels);
2946}
2947
2948MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2949 MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2950{
2951 if (info == (MagickCLCacheInfo) NULL)
2952 return((MagickCLCacheInfo) NULL);
2953 if (relinquish_pixels != MagickFalse)
2954 DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2955 else
2956 DestroyMagickCLCacheInfo(info);
2957 return((MagickCLCacheInfo) NULL);
2958}
2959
2960/*
2961%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2962% %
2963% %
2964% %
2965% R e l i n q u i s h M a g i c k C L D e v i c e %
2966% %
2967% %
2968% %
2969%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2970%
2971% RelinquishMagickCLDevice() releases the OpenCL device
2972%
2973% The format of the RelinquishMagickCLDevice method is:
2974%
2975% MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2976%
2977% A description of each parameter follows:
2978%
2979% o device: the OpenCL device to be released.
2980%
2981*/
2982
2983static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2984{
2985 if (device == (MagickCLDevice) NULL)
2986 return((MagickCLDevice) NULL);
2987
2988 device->platform_name=(char *) RelinquishMagickMemory(device->platform_name);
2989 device->vendor_name=(char *) RelinquishMagickMemory(device->vendor_name);
2990 device->name=(char *) RelinquishMagickMemory(device->name);
2991 device->version=(char *) RelinquishMagickMemory(device->version);
2992 if (device->program != (cl_program) NULL)
2993 (void) openCL_library->clReleaseProgram(device->program);
2994 while (device->command_queues_index >= 0)
2995 (void) openCL_library->clReleaseCommandQueue(
2996 device->command_queues[device->command_queues_index--]);
2997 RelinquishSemaphoreInfo(&device->lock);
2998 return((MagickCLDevice) RelinquishMagickMemory(device));
2999}
3000
3001/*
3002%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3003% %
3004% %
3005% %
3006% R e l i n q u i s h M a g i c k C L E n v %
3007% %
3008% %
3009% %
3010%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3011%
3012% RelinquishMagickCLEnv() releases the OpenCL environment
3013%
3014% The format of the RelinquishMagickCLEnv method is:
3015%
3016% MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
3017%
3018% A description of each parameter follows:
3019%
3020% o clEnv: the OpenCL environment to be released.
3021%
3022*/
3023
3024static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
3025{
3026 if (clEnv == (MagickCLEnv) NULL)
3027 return((MagickCLEnv) NULL);
3028
3029 RelinquishSemaphoreInfo(&clEnv->lock);
3030 RelinquishMagickCLDevices(clEnv);
3031 if (clEnv->contexts != (cl_context *) NULL)
3032 {
3033 ssize_t
3034 i;
3035
3036 for (i=0; i < (ssize_t) clEnv->number_contexts; i++)
3037 if (clEnv->contexts[i] != (cl_context) NULL)
3038 (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
3039 clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
3040 }
3041 return((MagickCLEnv) RelinquishMagickMemory(clEnv));
3042}
3043
3044/*
3045%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3046% %
3047% %
3048% %
3049+ R e q u e s t O p e n C L D e v i c e %
3050% %
3051% %
3052% %
3053%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3054%
3055% RequestOpenCLDevice() returns one of the enabled OpenCL devices.
3056%
3057% The format of the RequestOpenCLDevice method is:
3058%
3059% MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3060%
3061% A description of each parameter follows:
3062%
3063% o clEnv: the OpenCL environment.
3064*/
3065
3066MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3067{
3068 MagickCLDevice
3069 device;
3070
3071 double
3072 score,
3073 best_score;
3074
3075 size_t
3076 i;
3077
3078 if (clEnv == (MagickCLEnv) NULL)
3079 return((MagickCLDevice) NULL);
3080
3081 if (clEnv->number_devices == 1)
3082 {
3083 if (clEnv->devices[0]->enabled)
3084 return(clEnv->devices[0]);
3085 else
3086 return((MagickCLDevice) NULL);
3087 }
3088
3089 device=(MagickCLDevice) NULL;
3090 best_score=0.0;
3091 LockSemaphoreInfo(openCL_lock);
3092 for (i = 0; i < clEnv->number_devices; i++)
3093 {
3094 if (clEnv->devices[i]->enabled == MagickFalse)
3095 continue;
3096
3097 score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3098 clEnv->devices[i]->requested);
3099 if ((device == (MagickCLDevice) NULL) || (score < best_score))
3100 {
3101 device=clEnv->devices[i];
3102 best_score=score;
3103 }
3104 }
3105 if (device != (MagickCLDevice)NULL)
3106 device->requested++;
3107 UnlockSemaphoreInfo(openCL_lock);
3108
3109 return(device);
3110}
3111
3112/*
3113%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3114% %
3115% %
3116% %
3117% S e t O p e n C L D e v i c e E n a b l e d %
3118% %
3119% %
3120% %
3121%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3122%
3123% SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
3124%
3125% The format of the SetOpenCLDeviceEnabled method is:
3126%
3127% void SetOpenCLDeviceEnabled(MagickCLDevice device,
3128% MagickBooleanType value)
3129%
3130% A description of each parameter follows:
3131%
3132% o device: the OpenCL device.
3133%
3134% o value: determines if the device should be enabled or disabled.
3135*/
3136
3137MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
3138 const MagickBooleanType value)
3139{
3140 if (device == (MagickCLDevice) NULL)
3141 return;
3142 device->enabled=value;
3143}
3144
3145/*
3146%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3147% %
3148% %
3149% %
3150% 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 %
3151% %
3152% %
3153% %
3154%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3155%
3156% SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
3157% kernel profiling of a device.
3158%
3159% The format of the SetOpenCLKernelProfileEnabled method is:
3160%
3161% void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3162% MagickBooleanType value)
3163%
3164% A description of each parameter follows:
3165%
3166% o device: the OpenCL device.
3167%
3168% o value: determines if kernel profiling for the device should be enabled
3169% or disabled.
3170*/
3171
3172MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3173 const MagickBooleanType value)
3174{
3175 if (device == (MagickCLDevice) NULL)
3176 return;
3177 device->profile_kernels=value;
3178}
3179
3180/*
3181%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3182% %
3183% %
3184% %
3185% S e t O p e n C L E n a b l e d %
3186% %
3187% %
3188% %
3189%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3190%
3191% SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3192%
3193% The format of the SetOpenCLEnabled method is:
3194%
3195% void SetOpenCLEnabled(MagickBooleanType)
3196%
3197% A description of each parameter follows:
3198%
3199% o value: specify true to enable OpenCL acceleration
3200*/
3201
3202MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
3203{
3204 MagickCLEnv
3205 clEnv;
3206
3207 clEnv=GetCurrentOpenCLEnv();
3208 if (clEnv == (MagickCLEnv) NULL)
3209 return(MagickFalse);
3210 clEnv->enabled=value;
3211 return(clEnv->enabled);
3212}
3213
3214#else
3215
3216MagickExport double GetOpenCLDeviceBenchmarkScore(
3217 const MagickCLDevice magick_unused(device))
3218{
3219 magick_unreferenced(device);
3220 return(0.0);
3221}
3222
3223MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3224 const MagickCLDevice magick_unused(device))
3225{
3226 magick_unreferenced(device);
3227 return(MagickFalse);
3228}
3229
3230MagickExport const char *GetOpenCLDeviceName(
3231 const MagickCLDevice magick_unused(device))
3232{
3233 magick_unreferenced(device);
3234 return((const char *) NULL);
3235}
3236
3237MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
3238 ExceptionInfo *magick_unused(exception))
3239{
3240 magick_unreferenced(exception);
3241 if (length != (size_t *) NULL)
3242 *length=0;
3243 return((MagickCLDevice *) NULL);
3244}
3245
3246MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3247 const MagickCLDevice magick_unused(device))
3248{
3249 magick_unreferenced(device);
3250 return(UndefinedCLDeviceType);
3251}
3252
3253MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
3254 const MagickCLDevice magick_unused(device),size_t *length)
3255{
3256 magick_unreferenced(device);
3257 if (length != (size_t *) NULL)
3258 *length=0;
3259 return((const KernelProfileRecord *) NULL);
3260}
3261
3262MagickExport const char *GetOpenCLDeviceVersion(
3263 const MagickCLDevice magick_unused(device))
3264{
3265 magick_unreferenced(device);
3266 return((const char *) NULL);
3267}
3268
3269MagickExport MagickBooleanType GetOpenCLEnabled(void)
3270{
3271 return(MagickFalse);
3272}
3273
3274MagickExport void SetOpenCLDeviceEnabled(
3275 MagickCLDevice magick_unused(device),
3276 const MagickBooleanType magick_unused(value))
3277{
3278 magick_unreferenced(device);
3279 magick_unreferenced(value);
3280}
3281
3282MagickExport MagickBooleanType SetOpenCLEnabled(
3283 const MagickBooleanType magick_unused(value))
3284{
3285 magick_unreferenced(value);
3286 return(MagickFalse);
3287}
3288
3289MagickExport void SetOpenCLKernelProfileEnabled(
3290 MagickCLDevice magick_unused(device),
3291 const MagickBooleanType magick_unused(value))
3292{
3293 magick_unreferenced(device);
3294 magick_unreferenced(value);
3295}
3296#endif