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