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-2016 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 % http://www.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/artifact.h"
45 #include "MagickCore/cache.h"
46 #include "MagickCore/cache-private.h"
47 #include "MagickCore/color.h"
48 #include "MagickCore/compare.h"
49 #include "MagickCore/constitute.h"
50 #include "MagickCore/configure.h"
51 #include "MagickCore/distort.h"
52 #include "MagickCore/draw.h"
53 #include "MagickCore/effect.h"
54 #include "MagickCore/exception.h"
55 #include "MagickCore/exception-private.h"
56 #include "MagickCore/fx.h"
57 #include "MagickCore/gem.h"
58 #include "MagickCore/geometry.h"
59 #include "MagickCore/image.h"
60 #include "MagickCore/image-private.h"
61 #include "MagickCore/layer.h"
62 #include "MagickCore/mime-private.h"
63 #include "MagickCore/memory_.h"
64 #include "MagickCore/monitor.h"
65 #include "MagickCore/montage.h"
66 #include "MagickCore/morphology.h"
67 #include "MagickCore/nt-base.h"
68 #include "MagickCore/nt-base-private.h"
69 #include "MagickCore/opencl.h"
70 #include "MagickCore/opencl-private.h"
71 #include "MagickCore/option.h"
72 #include "MagickCore/policy.h"
73 #include "MagickCore/property.h"
74 #include "MagickCore/quantize.h"
75 #include "MagickCore/quantum.h"
76 #include "MagickCore/random_.h"
77 #include "MagickCore/random-private.h"
78 #include "MagickCore/resample.h"
79 #include "MagickCore/resource_.h"
80 #include "MagickCore/splay-tree.h"
81 #include "MagickCore/semaphore.h"
82 #include "MagickCore/statistic.h"
83 #include "MagickCore/string_.h"
84 #include "MagickCore/string-private.h"
85 #include "MagickCore/token.h"
86 #include "MagickCore/utility.h"
87 #include "MagickCore/utility-private.h"
88
89 #if defined(MAGICKCORE_OPENCL_SUPPORT)
90
91 #ifndef MAGICKCORE_WINDOWS_SUPPORT
92 #include <dlfcn.h>
93 #endif
94
95 #ifdef MAGICKCORE_HAVE_OPENCL_CL_H
96 #define MAGICKCORE_OPENCL_MACOSX 1
97 #endif
98
99 /*
100 Define declarations.
101 */
102 #define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"
103
104 /*
105 Typedef declarations.
106 */
107 typedef struct
108 {
109 long long freq;
110 long long clocks;
111 long long start;
112 } AccelerateTimer;
113
114 typedef struct
115 {
116 char
117 *name,
118 *platform_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
133 static MagickBooleanType
134 HasOpenCLDevices(MagickCLEnv,ExceptionInfo *),
135 LoadOpenCLLibrary(void);
136
137 static MagickCLDevice
138 RelinquishMagickCLDevice(MagickCLDevice);
139
140 static MagickCLEnv
141 RelinquishMagickCLEnv(MagickCLEnv);
142
143 static void
144 BenchmarkOpenCLDevices(MagickCLEnv);
145
146 extern const char
147 *accelerateKernels, *accelerateKernels2;
148
149 /* OpenCL library */
150 MagickLibrary
151 *openCL_library;
152
153 /* Default OpenCL environment */
154 MagickCLEnv
155 default_CLEnv;
156 MagickThreadType
157 test_thread_id=0;
158 SemaphoreInfo
159 *openCL_lock;
160
161 /* Cached location of the OpenCL cache files */
162 char
163 *cache_directory;
164 SemaphoreInfo
165 *cache_directory_lock;
166
IsSameOpenCLDevice(MagickCLDevice a,MagickCLDevice b)167 static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
168 MagickCLDevice b)
169 {
170 if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
171 (LocaleCompare(a->name,b->name) == 0) &&
172 (LocaleCompare(a->version,b->version) == 0) &&
173 (a->max_clock_frequency == b->max_clock_frequency) &&
174 (a->max_compute_units == b->max_compute_units))
175 return(MagickTrue);
176
177 return(MagickFalse);
178 }
179
IsBenchmarkedOpenCLDevice(MagickCLDevice a,MagickCLDeviceBenchmark * b)180 static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
181 MagickCLDeviceBenchmark *b)
182 {
183 if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
184 (LocaleCompare(a->name,b->name) == 0) &&
185 (LocaleCompare(a->version,b->version) == 0) &&
186 (a->max_clock_frequency == b->max_clock_frequency) &&
187 (a->max_compute_units == b->max_compute_units))
188 return(MagickTrue);
189
190 return(MagickFalse);
191 }
192
RelinquishMagickCLDevices(MagickCLEnv clEnv)193 static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
194 {
195 size_t
196 i;
197
198 if (clEnv->devices != (MagickCLDevice *) NULL)
199 {
200 for (i = 0; i < clEnv->number_devices; i++)
201 clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
202 clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
203 }
204 clEnv->number_devices=0;
205 }
206
MagickCreateDirectory(const char * path)207 static inline MagickBooleanType MagickCreateDirectory(const char *path)
208 {
209 int
210 status;
211
212 #ifdef MAGICKCORE_WINDOWS_SUPPORT
213 status=mkdir(path);
214 #else
215 status=mkdir(path, 0777);
216 #endif
217 return(status == 0 ? MagickTrue : MagickFalse);
218 }
219
InitAccelerateTimer(AccelerateTimer * timer)220 static inline void InitAccelerateTimer(AccelerateTimer *timer)
221 {
222 #ifdef _WIN32
223 QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
224 #else
225 timer->freq=(long long)1.0E3;
226 #endif
227 timer->clocks=0;
228 timer->start=0;
229 }
230
ReadAccelerateTimer(AccelerateTimer * timer)231 static inline double ReadAccelerateTimer(AccelerateTimer *timer)
232 {
233 return (double)timer->clocks/(double)timer->freq;
234 }
235
StartAccelerateTimer(AccelerateTimer * timer)236 static inline void StartAccelerateTimer(AccelerateTimer* timer)
237 {
238 #ifdef _WIN32
239 QueryPerformanceCounter((LARGE_INTEGER*)&timer->start);
240 #else
241 struct timeval
242 s;
243 gettimeofday(&s,0);
244 timer->start=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
245 (long long)1.0E3;
246 #endif
247 }
248
StopAccelerateTimer(AccelerateTimer * timer)249 static inline void StopAccelerateTimer(AccelerateTimer *timer)
250 {
251 long long
252 n;
253
254 n=0;
255 #ifdef _WIN32
256 QueryPerformanceCounter((LARGE_INTEGER*)&(n));
257 #else
258 struct timeval
259 s;
260 gettimeofday(&s,0);
261 n=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
262 (long long)1.0E3;
263 #endif
264 n-=timer->start;
265 timer->start=0;
266 timer->clocks+=n;
267 }
268
GetOpenCLCacheDirectory()269 static const char *GetOpenCLCacheDirectory()
270 {
271 if (cache_directory == (char *) NULL)
272 {
273 if (cache_directory_lock == (SemaphoreInfo *) NULL)
274 ActivateSemaphoreInfo(&cache_directory_lock);
275 LockSemaphoreInfo(cache_directory_lock);
276 if (cache_directory == (char *) NULL)
277 {
278 char
279 *home,
280 path[MagickPathExtent],
281 *temp;
282
283 MagickBooleanType
284 status;
285
286 struct stat
287 attributes;
288
289 temp=(char *) NULL;
290 home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
291 if (home == (char *) NULL)
292 {
293 home=GetEnvironmentValue("XDG_CACHE_HOME");
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 }
301
302 if (home != (char *) NULL)
303 {
304 /* first check if $HOME exists */
305 (void) FormatLocaleString(path,MagickPathExtent,"%s",home);
306 status=GetPathAttributes(path,&attributes);
307 if (status == MagickFalse)
308 status=MagickCreateDirectory(path);
309
310 /* first check if $HOME/ImageMagick exists */
311 if (status != MagickFalse)
312 {
313 (void) FormatLocaleString(path,MagickPathExtent,
314 "%s%sImageMagick",home,DirectorySeparator);
315
316 status=GetPathAttributes(path,&attributes);
317 if (status == MagickFalse)
318 status=MagickCreateDirectory(path);
319 }
320
321 if (status != MagickFalse)
322 {
323 temp=(char*) AcquireMagickMemory(strlen(path)+1);
324 CopyMagickString(temp,path,strlen(path)+1);
325 }
326 home=DestroyString(home);
327 }
328 else
329 {
330 home=GetEnvironmentValue("HOME");
331 if (home != (char *) NULL)
332 {
333 /* first check if $HOME/.cache exists */
334 (void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache",
335 home,DirectorySeparator);
336 status=GetPathAttributes(path,&attributes);
337 if (status == MagickFalse)
338 status=MagickCreateDirectory(path);
339
340 /* first check if $HOME/.cache/ImageMagick exists */
341 if (status != MagickFalse)
342 {
343 (void) FormatLocaleString(path,MagickPathExtent,
344 "%s%s.cache%sImageMagick",home,DirectorySeparator,
345 DirectorySeparator);
346 status=GetPathAttributes(path,&attributes);
347 if (status == MagickFalse)
348 status=MagickCreateDirectory(path);
349 }
350
351 if (status != MagickFalse)
352 {
353 temp=(char*) AcquireMagickMemory(strlen(path)+1);
354 CopyMagickString(temp,path,strlen(path)+1);
355 }
356 home=DestroyString(home);
357 }
358 }
359 if (temp == (char *) NULL)
360 temp=AcquireString("?");
361 cache_directory=temp;
362 }
363 UnlockSemaphoreInfo(cache_directory_lock);
364 }
365 if (*cache_directory == '?')
366 return((const char *) NULL);
367 return(cache_directory);
368 }
369
SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)370 static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
371 {
372 MagickCLDevice
373 device;
374
375 size_t
376 i,
377 j;
378
379 for (i = 0; i < clEnv->number_devices; i++)
380 clEnv->devices[i]->enabled=MagickFalse;
381
382 for (i = 0; i < clEnv->number_devices; i++)
383 {
384 device=clEnv->devices[i];
385 if (device->type != type)
386 continue;
387
388 device->enabled=MagickTrue;
389 for (j = i+1; j < clEnv->number_devices; j++)
390 {
391 MagickCLDevice
392 other_device;
393
394 other_device=clEnv->devices[j];
395 if (IsSameOpenCLDevice(device,other_device))
396 other_device->enabled=MagickTrue;
397 }
398 }
399 }
400
StringSignature(const char * string)401 static size_t StringSignature(const char* string)
402 {
403 size_t
404 n,
405 i,
406 j,
407 signature,
408 stringLength;
409
410 union
411 {
412 const char* s;
413 const size_t* u;
414 } p;
415
416 stringLength=(size_t) strlen(string);
417 signature=stringLength;
418 n=stringLength/sizeof(size_t);
419 p.s=string;
420 for (i = 0; i < n; i++)
421 signature^=p.u[i];
422 if (n * sizeof(size_t) != stringLength)
423 {
424 char
425 padded[4];
426
427 j=n*sizeof(size_t);
428 for (i = 0; i < 4; i++, j++)
429 {
430 if (j < stringLength)
431 padded[i]=p.s[j];
432 else
433 padded[i]=0;
434 }
435 p.s=padded;
436 signature^=p.u[0];
437 }
438 return(signature);
439 }
440
441 /*
442 Provide call to OpenCL library methods
443 */
444
CreateOpenCLBuffer(MagickCLDevice device,cl_mem_flags flags,size_t size,void * host_ptr)445 MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
446 cl_mem_flags flags,size_t size,void *host_ptr)
447 {
448 return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr,
449 (cl_int *) NULL));
450 }
451
ReleaseOpenCLKernel(cl_kernel kernel)452 MagickPrivate void ReleaseOpenCLKernel(cl_kernel kernel)
453 {
454 (void) openCL_library->clReleaseKernel(kernel);
455 }
456
ReleaseOpenCLMemObject(cl_mem memobj)457 MagickPrivate void ReleaseOpenCLMemObject(cl_mem memobj)
458 {
459 (void) openCL_library->clReleaseMemObject(memobj);
460 }
461
SetOpenCLKernelArg(cl_kernel kernel,cl_uint arg_index,size_t arg_size,const void * arg_value)462 MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,cl_uint arg_index,
463 size_t arg_size,const void *arg_value)
464 {
465 return(openCL_library->clSetKernelArg(kernel,arg_index,arg_size,arg_value));
466 }
467
468 /*
469 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
470 % %
471 % %
472 % %
473 + A c q u i r e M a g i c k C L C a c h e I n f o %
474 % %
475 % %
476 % %
477 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
478 %
479 % AcquireMagickCLCacheInfo() acquires an OpenCL cache info structure.
480 %
481 % The format of the AcquireMagickCLCacheInfo method is:
482 %
483 % MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
484 % Quantum *pixels,const MagickSizeType length)
485 %
486 % A description of each parameter follows:
487 %
488 % o device: the OpenCL device.
489 %
490 % o pixels: the pixel buffer of the image.
491 %
492 % o length: the length of the pixel buffer.
493 %
494 */
495
AcquireMagickCLCacheInfo(MagickCLDevice device,Quantum * pixels,const MagickSizeType length)496 MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
497 Quantum *pixels,const MagickSizeType length)
498 {
499 cl_int
500 status;
501
502 MagickCLCacheInfo
503 info;
504
505 info=(MagickCLCacheInfo) AcquireMagickMemory(sizeof(*info));
506 if (info == (MagickCLCacheInfo) NULL)
507 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
508 (void) ResetMagickMemory(info,0,sizeof(*info));
509 LockSemaphoreInfo(openCL_lock);
510 device->requested++;
511 UnlockSemaphoreInfo(openCL_lock);
512 info->device=device;
513 info->length=length;
514 info->pixels=pixels;
515 info->buffer=openCL_library->clCreateBuffer(device->context,
516 CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,
517 &status);
518 if (status == CL_SUCCESS)
519 return(info);
520 LockSemaphoreInfo(openCL_lock);
521 device->requested--;
522 UnlockSemaphoreInfo(openCL_lock);
523 return((MagickCLCacheInfo) RelinquishMagickMemory(info));
524 }
525
526 /*
527 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
528 % %
529 % %
530 % %
531 % A c q u i r e M a g i c k C L D e v i c e %
532 % %
533 % %
534 % %
535 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
536 %
537 % AcquireMagickCLDevice() acquires an OpenCL device
538 %
539 % The format of the AcquireMagickCLDevice method is:
540 %
541 % MagickCLDevice AcquireMagickCLDevice()
542 %
543 */
544
AcquireMagickCLDevice()545 static MagickCLDevice AcquireMagickCLDevice()
546 {
547 MagickCLDevice
548 device;
549
550 device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device));
551 if (device != NULL)
552 {
553 (void) ResetMagickMemory(device,0,sizeof(*device));
554 ActivateSemaphoreInfo(&device->lock);
555 device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
556 device->command_queues_index=-1;
557 device->enabled=MagickTrue;
558 }
559 return(device);
560 }
561
562 /*
563 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
564 % %
565 % %
566 % %
567 % A c q u i r e M a g i c k C L E n v %
568 % %
569 % %
570 % %
571 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
572 %
573 % AcquireMagickCLEnv() allocates the MagickCLEnv structure
574 %
575 */
576
AcquireMagickCLEnv(void)577 static MagickCLEnv AcquireMagickCLEnv(void)
578 {
579 const char
580 *option;
581
582 MagickCLEnv
583 clEnv;
584
585 clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv));
586 if (clEnv != (MagickCLEnv) NULL)
587 {
588 (void) ResetMagickMemory(clEnv,0,sizeof(*clEnv));
589 ActivateSemaphoreInfo(&clEnv->lock);
590 clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
591 clEnv->enabled=MagickTrue;
592 option=getenv("MAGICK_OCL_DEVICE");
593 if ((option != (const char *) NULL) && (strcmp(option,"OFF") == 0))
594 clEnv->enabled=MagickFalse;
595 }
596 return clEnv;
597 }
598
599 /*
600 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
601 % %
602 % %
603 % %
604 + 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 %
605 % %
606 % %
607 % %
608 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
609 %
610 % AcquireOpenCLCommandQueue() acquires an OpenCL command queue
611 %
612 % The format of the AcquireOpenCLCommandQueue method is:
613 %
614 % cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
615 %
616 % A description of each parameter follows:
617 %
618 % o device: the OpenCL device.
619 %
620 */
621
AcquireOpenCLCommandQueue(MagickCLDevice device)622 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
623 {
624 cl_command_queue
625 queue;
626
627 cl_command_queue_properties
628 properties;
629
630 assert(device != (MagickCLDevice) NULL);
631 LockSemaphoreInfo(device->lock);
632 if ((device->profile_kernels == MagickFalse) &&
633 (device->command_queues_index >= 0))
634 {
635 queue=device->command_queues[device->command_queues_index--];
636 UnlockSemaphoreInfo(device->lock);
637 }
638 else
639 {
640 UnlockSemaphoreInfo(device->lock);
641 properties=(cl_command_queue_properties) NULL;
642 if (device->profile_kernels != MagickFalse)
643 properties=CL_QUEUE_PROFILING_ENABLE;
644 queue=openCL_library->clCreateCommandQueue(device->context,
645 device->deviceID,properties,(cl_int *) NULL);
646 }
647 return(queue);
648 }
649
650 /*
651 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
652 % %
653 % %
654 % %
655 + A c q u i r e O p e n C L K e r n e l %
656 % %
657 % %
658 % %
659 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
660 %
661 % AcquireOpenCLKernel() acquires an OpenCL kernel
662 %
663 % The format of the AcquireOpenCLKernel method is:
664 %
665 % cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
666 % MagickOpenCLProgram program, const char* kernelName)
667 %
668 % A description of each parameter follows:
669 %
670 % o clEnv: the OpenCL environment.
671 %
672 % o program: the OpenCL program module that the kernel belongs to.
673 %
674 % o kernelName: the name of the kernel
675 %
676 */
677
AcquireOpenCLKernel(MagickCLDevice device,const char * kernel_name)678 MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
679 const char *kernel_name)
680 {
681 cl_kernel
682 kernel;
683
684 assert(device != (MagickCLDevice) NULL);
685 kernel=openCL_library->clCreateKernel(device->program,kernel_name,
686 (cl_int *) NULL);
687 return(kernel);
688 }
689
690 /*
691 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
692 % %
693 % %
694 % %
695 % A u t o S e l e c t O p e n C L D e v i c e s %
696 % %
697 % %
698 % %
699 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
700 %
701 % AutoSelectOpenCLDevices() determines the best device based on the
702 % information from the micro-benchmark.
703 %
704 % The format of the AutoSelectOpenCLDevices method is:
705 %
706 % void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception)
707 %
708 % A description of each parameter follows:
709 %
710 % o clEnv: the OpenCL environment.
711 %
712 % o exception: return any errors or warnings in this structure.
713 %
714 */
715
LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char * xml)716 static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
717 {
718 char
719 keyword[MagickPathExtent],
720 *token;
721
722 const char
723 *q;
724
725 MagickCLDeviceBenchmark
726 *device_benchmark;
727
728 MagickStatusType
729 status;
730
731 size_t
732 i,
733 extent;
734
735 if (xml == (char *) NULL)
736 return;
737 status=MagickTrue;
738 device_benchmark=(MagickCLDeviceBenchmark *) NULL;
739 token=AcquireString(xml);
740 extent=strlen(token)+MagickPathExtent;
741 for (q=(char *) xml; *q != '\0'; )
742 {
743 /*
744 Interpret XML.
745 */
746 GetNextToken(q,&q,extent,token);
747 if (*token == '\0')
748 break;
749 (void) CopyMagickString(keyword,token,MagickPathExtent);
750 if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
751 {
752 /*
753 Doctype element.
754 */
755 while ((LocaleNCompare(q,"]>",2) != 0) && (*q != '\0'))
756 GetNextToken(q,&q,extent,token);
757 continue;
758 }
759 if (LocaleNCompare(keyword,"<!--",4) == 0)
760 {
761 /*
762 Comment element.
763 */
764 while ((LocaleNCompare(q,"->",2) != 0) && (*q != '\0'))
765 GetNextToken(q,&q,extent,token);
766 continue;
767 }
768 if (LocaleCompare(keyword,"<device") == 0)
769 {
770 /*
771 Device element.
772 */
773 device_benchmark=(MagickCLDeviceBenchmark *) AcquireMagickMemory(
774 sizeof(*device_benchmark));
775 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
776 break;
777 (void) ResetMagickMemory(device_benchmark,0,sizeof(*device_benchmark));
778 device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
779 continue;
780 }
781 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
782 continue;
783 if (LocaleCompare(keyword,"/>") == 0)
784 {
785 if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
786 {
787 if (LocaleCompare(device_benchmark->name, "CPU") == 0)
788 clEnv->cpu_score=device_benchmark->score;
789 else
790 {
791 MagickCLDevice
792 device;
793
794 /*
795 Set the score for all devices that match this device.
796 */
797 for (i = 0; i < clEnv->number_devices; i++)
798 {
799 device=clEnv->devices[i];
800 if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
801 device->score=device_benchmark->score;
802 }
803 }
804 }
805
806 device_benchmark->platform_name=RelinquishMagickMemory(
807 device_benchmark->platform_name);
808 device_benchmark->name=RelinquishMagickMemory(device_benchmark->name);
809 device_benchmark->version=RelinquishMagickMemory(
810 device_benchmark->version);
811 device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
812 device_benchmark);
813 continue;
814 }
815 GetNextToken(q,(const char **) NULL,extent,token);
816 if (*token != '=')
817 continue;
818 GetNextToken(q,&q,extent,token);
819 GetNextToken(q,&q,extent,token);
820 switch (*keyword)
821 {
822 case 'M':
823 case 'm':
824 {
825 if (LocaleCompare((char *) keyword,"maxClockFrequency") == 0)
826 {
827 device_benchmark->max_clock_frequency=StringToInteger(token);
828 break;
829 }
830 if (LocaleCompare((char *) keyword,"maxComputeUnits") == 0)
831 {
832 device_benchmark->max_compute_units=StringToInteger(token);
833 break;
834 }
835 break;
836 }
837 case 'N':
838 case 'n':
839 {
840 if (LocaleCompare((char *) keyword,"name") == 0)
841 device_benchmark->name=ConstantString(token);
842 break;
843 }
844 case 'P':
845 case 'p':
846 {
847 if (LocaleCompare((char *) keyword,"platform") == 0)
848 device_benchmark->platform_name=ConstantString(token);
849 break;
850 }
851 case 'S':
852 case 's':
853 {
854 if (LocaleCompare((char *) keyword,"score") == 0)
855 device_benchmark->score=StringToDouble(token,(char **) NULL);
856 break;
857 }
858 case 'V':
859 case 'v':
860 {
861 if (LocaleCompare((char *) keyword,"version") == 0)
862 device_benchmark->version=ConstantString(token);
863 break;
864 }
865 default:
866 break;
867 }
868 }
869 token=(char *) RelinquishMagickMemory(token);
870 device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
871 device_benchmark);
872 }
873
CanWriteProfileToFile(const char * filename)874 static MagickBooleanType CanWriteProfileToFile(const char *filename)
875 {
876 FILE
877 *profileFile;
878
879 profileFile=fopen(filename,"ab");
880
881 if (profileFile == (FILE *)NULL)
882 return(MagickFalse);
883
884 fclose(profileFile);
885 return(MagickTrue);
886 }
887
LoadOpenCLBenchmarks(MagickCLEnv clEnv,ExceptionInfo * exception)888 static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv,
889 ExceptionInfo *exception)
890 {
891 char
892 filename[MagickPathExtent];
893
894 const StringInfo
895 *option;
896
897 LinkedListInfo
898 *options;
899
900 size_t
901 i;
902
903 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
904 GetOpenCLCacheDirectory(),DirectorySeparator,
905 IMAGEMAGICK_PROFILE_FILE);
906
907 /*
908 We don't run the benchmark when we can not write out a device profile. The
909 first GPU device will be used.
910 */
911 #if !defined(MAGICKCORE_ZERO_CONFIGURATION_SUPPORT)
912 if (CanWriteProfileToFile(filename) == MagickFalse)
913 #endif
914 {
915 for (i = 0; i < clEnv->number_devices; i++)
916 clEnv->devices[i]->score=1.0;
917
918 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
919 return(MagickFalse);
920 }
921
922 options=GetConfigureOptions(filename,exception);
923 option=(const StringInfo *) GetNextValueInLinkedList(options);
924 while (option != (const StringInfo *) NULL)
925 {
926 LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(
927 option));
928 option=(const StringInfo *) GetNextValueInLinkedList(options);
929 }
930 options=DestroyConfigureOptions(options);
931 return(MagickTrue);
932 }
933
AutoSelectOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo * exception)934 static void AutoSelectOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
935 {
936 const char
937 *option;
938
939 double
940 best_score;
941
942 MagickBooleanType
943 benchmark;
944
945 size_t
946 i;
947
948 option=getenv("MAGICK_OCL_DEVICE");
949 if (option != (const char *) NULL)
950 {
951 if (strcmp(option,"GPU") == 0)
952 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
953 else if (strcmp(option,"CPU") == 0)
954 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
955 else if (strcmp(option,"OFF") == 0)
956 {
957 for (i = 0; i < clEnv->number_devices; i++)
958 clEnv->devices[i]->enabled=MagickFalse;
959 clEnv->enabled=MagickFalse;
960 }
961 }
962
963 if (LoadOpenCLBenchmarks(clEnv,exception) == MagickFalse)
964 return;
965
966 benchmark=MagickFalse;
967 if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
968 benchmark=MagickTrue;
969 else
970 {
971 for (i = 0; i < clEnv->number_devices; i++)
972 {
973 if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
974 {
975 benchmark=MagickTrue;
976 break;
977 }
978 }
979 }
980
981 if (benchmark != MagickFalse)
982 BenchmarkOpenCLDevices(clEnv);
983
984 best_score=clEnv->cpu_score;
985 for (i = 0; i < clEnv->number_devices; i++)
986 best_score=MagickMin(clEnv->devices[i]->score,best_score);
987
988 for (i = 0; i < clEnv->number_devices; i++)
989 {
990 if (clEnv->devices[i]->score != best_score)
991 clEnv->devices[i]->enabled=MagickFalse;
992 }
993 }
994
995 /*
996 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
997 % %
998 % %
999 % %
1000 % B e n c h m a r k O p e n C L D e v i c e s %
1001 % %
1002 % %
1003 % %
1004 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1005 %
1006 % BenchmarkOpenCLDevices() benchmarks the OpenCL devices and the CPU to help
1007 % the automatic selection of the best device.
1008 %
1009 % The format of the BenchmarkOpenCLDevices method is:
1010 %
1011 % void BenchmarkOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
1012 %
1013 % A description of each parameter follows:
1014 %
1015 % o clEnv: the OpenCL environment.
1016 %
1017 % o exception: return any errors or warnings
1018 */
1019
RunOpenCLBenchmark(MagickBooleanType is_cpu)1020 static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
1021 {
1022 AccelerateTimer
1023 timer;
1024
1025 ExceptionInfo
1026 *exception;
1027
1028 Image
1029 *inputImage;
1030
1031 ImageInfo
1032 *imageInfo;
1033
1034 size_t
1035 i;
1036
1037 exception=AcquireExceptionInfo();
1038 imageInfo=AcquireImageInfo();
1039 CloneString(&imageInfo->size,"2048x1536");
1040 CopyMagickString(imageInfo->filename,"xc:none",MagickPathExtent);
1041 inputImage=ReadImage(imageInfo,exception);
1042
1043 InitAccelerateTimer(&timer);
1044
1045 for (i=0; i<=2; i++)
1046 {
1047 Image
1048 *bluredImage,
1049 *resizedImage,
1050 *unsharpedImage;
1051
1052 if (i > 0)
1053 StartAccelerateTimer(&timer);
1054
1055 bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
1056 unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
1057 exception);
1058 resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
1059 exception);
1060
1061 /*
1062 We need this to get a proper performance benchmark, the operations
1063 are executed asynchronous.
1064 */
1065 if (is_cpu == MagickFalse)
1066 {
1067 CacheInfo
1068 *cache_info;
1069
1070 cache_info=(CacheInfo *) resizedImage->cache;
1071 if (cache_info->opencl != (MagickCLCacheInfo) NULL)
1072 openCL_library->clWaitForEvents(cache_info->opencl->event_count,
1073 cache_info->opencl->events);
1074 }
1075
1076 if (i > 0)
1077 StopAccelerateTimer(&timer);
1078
1079 if (bluredImage != (Image *) NULL)
1080 DestroyImage(bluredImage);
1081 if (unsharpedImage != (Image *) NULL)
1082 DestroyImage(unsharpedImage);
1083 if (resizedImage != (Image *) NULL)
1084 DestroyImage(resizedImage);
1085 }
1086 DestroyImage(inputImage);
1087 return(ReadAccelerateTimer(&timer));
1088 }
1089
RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,MagickCLDevice device)1090 static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
1091 MagickCLDevice device)
1092 {
1093 testEnv->devices[0]=device;
1094 default_CLEnv=testEnv;
1095 device->score=RunOpenCLBenchmark(MagickFalse);
1096 default_CLEnv=clEnv;
1097 testEnv->devices[0]=(MagickCLDevice) NULL;
1098 }
1099
CacheOpenCLBenchmarks(MagickCLEnv clEnv)1100 static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
1101 {
1102 char
1103 filename[MagickPathExtent];
1104
1105 FILE
1106 *cache_file;
1107
1108 MagickCLDevice
1109 device;
1110
1111 size_t
1112 i,
1113 j;
1114
1115 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1116 GetOpenCLCacheDirectory(),DirectorySeparator,
1117 IMAGEMAGICK_PROFILE_FILE);
1118
1119 cache_file=fopen_utf8(filename,"wb");
1120 if (cache_file == (FILE *) NULL)
1121 return;
1122 fwrite("<devices>\n",sizeof(char),10,cache_file);
1123 fprintf(cache_file," <device name=\"CPU\" score=\"%.4g\"/>\n",
1124 clEnv->cpu_score);
1125 for (i = 0; i < clEnv->number_devices; i++)
1126 {
1127 MagickBooleanType
1128 duplicate;
1129
1130 device=clEnv->devices[i];
1131 duplicate=MagickFalse;
1132 for (j = 0; j < i; j++)
1133 {
1134 if (IsSameOpenCLDevice(clEnv->devices[j],device))
1135 {
1136 duplicate=MagickTrue;
1137 break;
1138 }
1139 }
1140
1141 if (duplicate)
1142 continue;
1143
1144 if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1145 fprintf(cache_file," <device platform=\"%s\" name=\"%s\" version=\"%s\"\
1146 maxClockFrequency=\"%d\" maxComputeUnits=\"%d\" score=\"%.4g\"/>\n",
1147 device->platform_name,device->name,device->version,
1148 (int)device->max_clock_frequency,(int)device->max_compute_units,
1149 device->score);
1150 }
1151 fwrite("</devices>",sizeof(char),10,cache_file);
1152
1153 fclose(cache_file);
1154 }
1155
BenchmarkOpenCLDevices(MagickCLEnv clEnv)1156 static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
1157 {
1158 MagickCLDevice
1159 device;
1160
1161 MagickCLEnv
1162 testEnv;
1163
1164 size_t
1165 i,
1166 j;
1167
1168 testEnv=AcquireMagickCLEnv();
1169 testEnv->library=openCL_library;
1170 testEnv->devices=(MagickCLDevice *) AcquireMagickMemory(
1171 sizeof(MagickCLDevice));
1172 testEnv->number_devices=1;
1173 testEnv->benchmark_thread_id=GetMagickThreadId();
1174 testEnv->initialized=MagickTrue;
1175
1176 for (i = 0; i < clEnv->number_devices; i++)
1177 clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
1178
1179 for (i = 0; i < clEnv->number_devices; i++)
1180 {
1181 device=clEnv->devices[i];
1182 if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1183 RunDeviceBenckmark(clEnv,testEnv,device);
1184
1185 /* Set the score on all the other devices that are the same */
1186 for (j = i+1; j < clEnv->number_devices; j++)
1187 {
1188 MagickCLDevice
1189 other_device;
1190
1191 other_device=clEnv->devices[j];
1192 if (IsSameOpenCLDevice(device,other_device))
1193 other_device->score=device->score;
1194 }
1195 }
1196
1197 testEnv->enabled=MagickFalse;
1198 default_CLEnv=testEnv;
1199 clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
1200 default_CLEnv=clEnv;
1201
1202 testEnv=RelinquishMagickCLEnv(testEnv);
1203 CacheOpenCLBenchmarks(clEnv);
1204 }
1205
1206 /*
1207 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1208 % %
1209 % %
1210 % %
1211 % C o m p i l e O p e n C L K e r n e l %
1212 % %
1213 % %
1214 % %
1215 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1216 %
1217 % CompileOpenCLKernel() compiles the kernel for the specified device. The
1218 % kernel will be cached on disk to reduce the compilation time.
1219 %
1220 % The format of the CompileOpenCLKernel method is:
1221 %
1222 % MagickBooleanType AcquireOpenCLKernel(MagickCLDevice clEnv,
1223 % unsigned int signature,const char *kernel,const char *options,
1224 % ExceptionInfo *exception)
1225 %
1226 % A description of each parameter follows:
1227 %
1228 % o device: the OpenCL device.
1229 %
1230 % o kernel: the source code of the kernel.
1231 %
1232 % o options: options for the compiler.
1233 %
1234 % o signature: a number to uniquely identify the kernel
1235 %
1236 % o exception: return any errors or warnings in this structure.
1237 %
1238 */
1239
CacheOpenCLKernel(MagickCLDevice device,char * filename,ExceptionInfo * exception)1240 static void CacheOpenCLKernel(MagickCLDevice device,char *filename,
1241 ExceptionInfo *exception)
1242 {
1243 cl_uint
1244 status;
1245
1246 size_t
1247 binaryProgramSize;
1248
1249 unsigned char
1250 *binaryProgram;
1251
1252 status=openCL_library->clGetProgramInfo(device->program,
1253 CL_PROGRAM_BINARY_SIZES,sizeof(size_t),&binaryProgramSize,NULL);
1254 if (status != CL_SUCCESS)
1255 return;
1256
1257 binaryProgram=(unsigned char*) AcquireMagickMemory(binaryProgramSize);
1258 status=openCL_library->clGetProgramInfo(device->program,
1259 CL_PROGRAM_BINARIES,sizeof(unsigned char*),&binaryProgram,NULL);
1260 if (status == CL_SUCCESS)
1261 (void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
1262 binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1263 }
1264
LoadCachedOpenCLKernel(MagickCLDevice device,const char * filename)1265 static MagickBooleanType LoadCachedOpenCLKernel(MagickCLDevice device,
1266 const char *filename)
1267 {
1268 cl_int
1269 binaryStatus,
1270 status;
1271
1272 ExceptionInfo
1273 *exception;
1274
1275 size_t
1276 length;
1277
1278 unsigned char
1279 *binaryProgram;
1280
1281 exception=AcquireExceptionInfo();
1282 binaryProgram=(unsigned char *) FileToBlob(filename,~0UL,&length,exception);
1283 exception=DestroyExceptionInfo(exception);
1284 if (binaryProgram == (unsigned char *) NULL)
1285 return(MagickFalse);
1286 device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
1287 &device->deviceID,&length,(const unsigned char**)&binaryProgram,
1288 &binaryStatus,&status);
1289 binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1290 return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
1291 MagickTrue);
1292 }
1293
LogOpenCLBuildFailure(MagickCLDevice device,const char * kernel,ExceptionInfo * exception)1294 static void LogOpenCLBuildFailure(MagickCLDevice device,const char *kernel,
1295 ExceptionInfo *exception)
1296 {
1297 char
1298 filename[MagickPathExtent],
1299 *log;
1300
1301 size_t
1302 logSize;
1303
1304 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1305 GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.cl");
1306
1307 (void) remove_utf8(filename);
1308 (void) BlobToFile(filename,kernel,strlen(kernel),exception);
1309
1310 openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1311 CL_PROGRAM_BUILD_LOG,0,NULL,&logSize);
1312 log=(char*)AcquireMagickMemory(logSize);
1313 openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1314 CL_PROGRAM_BUILD_LOG,logSize,log,&logSize);
1315
1316 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1317 GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.log");
1318
1319 (void) remove_utf8(filename);
1320 (void) BlobToFile(filename,log,logSize,exception);
1321 }
1322
CompileOpenCLKernel(MagickCLDevice device,const char * kernel,const char * options,size_t signature,ExceptionInfo * exception)1323 static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
1324 const char *kernel,const char *options,size_t signature,
1325 ExceptionInfo *exception)
1326 {
1327 char
1328 deviceName[MagickPathExtent],
1329 filename[MagickPathExtent],
1330 *ptr;
1331
1332 cl_int
1333 status;
1334
1335 MagickBooleanType
1336 loaded;
1337
1338 size_t
1339 length;
1340
1341 (void) CopyMagickString(deviceName,device->name,MagickPathExtent);
1342 ptr=deviceName;
1343 /* Strip out illegal characters for file names */
1344 while (*ptr != '\0')
1345 {
1346 if ((*ptr == ' ') || (*ptr == '\\') || (*ptr == '/') || (*ptr == ':') ||
1347 (*ptr == '*') || (*ptr == '?') || (*ptr == '"') || (*ptr == '<') ||
1348 (*ptr == '>' || *ptr == '|'))
1349 *ptr = '_';
1350 ptr++;
1351 }
1352 (void) FormatLocaleString(filename,MagickPathExtent,
1353 "%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
1354 DirectorySeparator,"magick_opencl",deviceName,signature,
1355 (double) sizeof(char*)*8);
1356 loaded=LoadCachedOpenCLKernel(device,filename);
1357 if (loaded == MagickFalse)
1358 {
1359 /* Binary CL program unavailable, compile the program from source */
1360 length=strlen(kernel);
1361 device->program=openCL_library->clCreateProgramWithSource(
1362 device->context,1,&kernel,&length,&status);
1363 if (status != CL_SUCCESS)
1364 return(MagickFalse);
1365 }
1366
1367 status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
1368 options,NULL,NULL);
1369 if (status != CL_SUCCESS)
1370 {
1371 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
1372 "clBuildProgram failed.","(%d)",(int)status);
1373 LogOpenCLBuildFailure(device,kernel,exception);
1374 return(MagickFalse);
1375 }
1376
1377 /* Save the binary to a file to avoid re-compilation of the kernels */
1378 if (loaded == MagickFalse)
1379 CacheOpenCLKernel(device,filename,exception);
1380
1381 return(MagickTrue);
1382 }
1383
1384 /*
1385 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1386 % %
1387 % %
1388 % %
1389 + C o p y M a g i c k C L C a c h e I n f o %
1390 % %
1391 % %
1392 % %
1393 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1394 %
1395 % CopyMagickCLCacheInfo() copies the memory from the device into host memory.
1396 %
1397 % The format of the CopyMagickCLCacheInfo method is:
1398 %
1399 % void CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1400 %
1401 % A description of each parameter follows:
1402 %
1403 % o info: the OpenCL cache info.
1404 %
1405 */
CopyMagickCLCacheInfo(MagickCLCacheInfo info)1406 MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1407 {
1408 cl_command_queue
1409 queue;
1410
1411 Quantum
1412 *pixels;
1413
1414 if (info == (MagickCLCacheInfo) NULL)
1415 return((MagickCLCacheInfo) NULL);
1416 if (info->event_count > 0)
1417 {
1418 queue=AcquireOpenCLCommandQueue(info->device);
1419 pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
1420 CL_MAP_READ | CL_MAP_WRITE,0,info->length,info->event_count,
1421 info->events,(cl_event *) NULL,(cl_int *) NULL);
1422 assert(pixels == info->pixels);
1423 ReleaseOpenCLCommandQueue(info->device,queue);
1424 }
1425 return(RelinquishMagickCLCacheInfo(info,MagickFalse));
1426 }
1427
1428 /*
1429 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1430 % %
1431 % %
1432 % %
1433 + D u m p O p e n C L P r o f i l e D a t a %
1434 % %
1435 % %
1436 % %
1437 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1438 %
1439 % DumpOpenCLProfileData() dumps the kernel profile data.
1440 %
1441 % The format of the DumpProfileData method is:
1442 %
1443 % void DumpProfileData()
1444 %
1445 */
1446
DumpOpenCLProfileData()1447 MagickPrivate void DumpOpenCLProfileData()
1448 {
1449 #define OpenCLLog(message) \
1450 fwrite(message,sizeof(char),strlen(message),log); \
1451 fwrite("\n",sizeof(char),1,log);
1452
1453 char
1454 buf[4096],
1455 filename[MagickPathExtent],
1456 indent[160];
1457
1458 FILE
1459 *log;
1460
1461 MagickCLEnv
1462 clEnv;
1463
1464 size_t
1465 i,
1466 j;
1467
1468 clEnv=GetCurrentOpenCLEnv();
1469 if (clEnv == (MagickCLEnv) NULL)
1470 return;
1471
1472 for (i = 0; i < clEnv->number_devices; i++)
1473 if (clEnv->devices[i]->profile_kernels != MagickFalse)
1474 break;
1475 if (i == clEnv->number_devices)
1476 return;
1477
1478 (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1479 GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
1480
1481 log=fopen_utf8(filename,"wb");
1482
1483 for (i = 0; i < clEnv->number_devices; i++)
1484 {
1485 MagickCLDevice
1486 device;
1487
1488 device=clEnv->devices[i];
1489 if ((device->profile_kernels == MagickFalse) ||
1490 (device->profile_records == (KernelProfileRecord *) NULL))
1491 continue;
1492
1493 OpenCLLog("====================================================");
1494 fprintf(log,"Device: %s\n",device->name);
1495 fprintf(log,"Version: %s\n",device->version);
1496 OpenCLLog("====================================================");
1497 OpenCLLog(" average calls min max");
1498 OpenCLLog(" ------- ----- --- ---");
1499 j=0;
1500 while (device->profile_records[j] != (KernelProfileRecord) NULL)
1501 {
1502 KernelProfileRecord
1503 profile;
1504
1505 profile=device->profile_records[j];
1506 strcpy(indent," ");
1507 strncpy(indent,profile->kernel_name,MagickMin(strlen(
1508 profile->kernel_name),strlen(indent)-1));
1509 sprintf(buf,"%s %7d %7d %7d %7d",indent,(int) (profile->total/
1510 profile->count),(int) profile->count,(int) profile->min,
1511 (int) profile->max);
1512 OpenCLLog(buf);
1513 j++;
1514 }
1515 OpenCLLog("====================================================");
1516 fwrite("\n\n",sizeof(char),2,log);
1517 }
1518 fclose(log);
1519 }
1520 /*
1521 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1522 % %
1523 % %
1524 % %
1525 + E n q u e u e O p e n C L K e r n e l %
1526 % %
1527 % %
1528 % %
1529 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1530 %
1531 % EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
1532 % events with the images.
1533 %
1534 % The format of the EnqueueOpenCLKernel method is:
1535 %
1536 % MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
1537 % const size_t *global_work_offset,const size_t *global_work_size,
1538 % const size_t *local_work_size,const Image *input_image,
1539 % const Image *output_image,ExceptionInfo *exception)
1540 %
1541 % A description of each parameter follows:
1542 %
1543 % o kernel: the OpenCL kernel.
1544 %
1545 % o work_dim: the number of dimensions used to specify the global work-items
1546 % and work-items in the work-group.
1547 %
1548 % o offset: can be used to specify an array of work_dim unsigned values
1549 % that describe the offset used to calculate the global ID of a
1550 % work-item.
1551 %
1552 % o gsize: points to an array of work_dim unsigned values that describe the
1553 % number of global work-items in work_dim dimensions that will
1554 % execute the kernel function.
1555 %
1556 % o lsize: points to an array of work_dim unsigned values that describe the
1557 % number of work-items that make up a work-group that will execute
1558 % the kernel specified by kernel.
1559 %
1560 % o input_image: the input image of the operation.
1561 %
1562 % o output_image: the output or secondairy image of the operation.
1563 %
1564 % o exception: return any errors or warnings in this structure.
1565 %
1566 */
1567
RegisterCacheEvent(MagickCLCacheInfo info,cl_event event)1568 static void RegisterCacheEvent(MagickCLCacheInfo info,cl_event event)
1569 {
1570 assert(info != (MagickCLCacheInfo) NULL);
1571 assert(event != (cl_event) NULL);
1572 if (info->events == (cl_event *) NULL)
1573 {
1574 info->events=AcquireMagickMemory(sizeof(*info->events));
1575 info->event_count=1;
1576 }
1577 else
1578 info->events=ResizeQuantumMemory(info->events,++info->event_count,
1579 sizeof(*info->events));
1580 if (info->events == (cl_event *) NULL)
1581 ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
1582 info->events[info->event_count-1]=event;
1583 openCL_library->clRetainEvent(event);
1584 }
1585
EnqueueOpenCLKernel(cl_command_queue queue,cl_kernel kernel,cl_uint work_dim,const size_t * offset,const size_t * gsize,const size_t * lsize,const Image * input_image,const Image * output_image,ExceptionInfo * exception)1586 MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1587 cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
1588 const size_t *lsize,const Image *input_image,const Image *output_image,
1589 ExceptionInfo *exception)
1590 {
1591 CacheInfo
1592 *output_info,
1593 *input_info;
1594
1595 cl_event
1596 event,
1597 *events;
1598
1599 cl_int
1600 status;
1601
1602 cl_uint
1603 event_count;
1604
1605 assert(input_image != (const Image *) NULL);
1606 input_info=(CacheInfo *) input_image->cache;
1607 assert(input_info != (CacheInfo *) NULL);
1608 assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1609 event_count=input_info->opencl->event_count;
1610 events=input_info->opencl->events;
1611 output_info=(CacheInfo *) NULL;
1612 if (output_image != (const Image *) NULL)
1613 {
1614 output_info=(CacheInfo *) output_image->cache;
1615 assert(output_info != (CacheInfo *) NULL);
1616 assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1617 if (output_info->opencl->event_count > 0)
1618 {
1619 ssize_t
1620 i;
1621
1622 event_count+=output_info->opencl->event_count;
1623 events=AcquireQuantumMemory(event_count,sizeof(*events));
1624 if (events == (cl_event *) NULL)
1625 return(MagickFalse);
1626 for (i=0; i < (ssize_t) event_count; i++)
1627 {
1628 if (i < (ssize_t) input_info->opencl->event_count)
1629 events[i]=input_info->opencl->events[i];
1630 else
1631 events[i]=output_info->opencl->events[i-
1632 input_info->opencl->event_count];
1633 }
1634 }
1635 }
1636 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1637 gsize,lsize,event_count,events,&event);
1638 if ((output_info != (CacheInfo *) NULL) &&
1639 (output_info->opencl->event_count > 0))
1640 events=(cl_event *) RelinquishMagickMemory(events);
1641 if (status != CL_SUCCESS)
1642 {
1643 (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1644 GetMagickModule(),ResourceLimitWarning,
1645 "clEnqueueNDRangeKernel failed.","'%s'",".");
1646 return(MagickFalse);
1647 }
1648 if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1649 {
1650 RegisterCacheEvent(input_info->opencl,event);
1651 if (output_info != (CacheInfo *) NULL)
1652 RegisterCacheEvent(output_info->opencl,event);
1653 }
1654 openCL_library->clReleaseEvent(event);
1655 return(MagickTrue);
1656 }
1657
1658 /*
1659 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1660 % %
1661 % %
1662 % %
1663 + G e t C u r r u n t O p e n C L E n v %
1664 % %
1665 % %
1666 % %
1667 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1668 %
1669 % GetCurrentOpenCLEnv() returns the current OpenCL env
1670 %
1671 % The format of the GetCurrentOpenCLEnv method is:
1672 %
1673 % MagickCLEnv GetCurrentOpenCLEnv()
1674 %
1675 */
1676
GetCurrentOpenCLEnv(void)1677 MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
1678 {
1679 if (default_CLEnv != (MagickCLEnv) NULL)
1680 {
1681 if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1682 (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1683 return((MagickCLEnv) NULL);
1684 else
1685 return(default_CLEnv);
1686 }
1687
1688 if (GetOpenCLCacheDirectory() == (char *) NULL)
1689 return((MagickCLEnv) NULL);
1690
1691 if (openCL_lock == (SemaphoreInfo *) NULL)
1692 ActivateSemaphoreInfo(&openCL_lock);
1693
1694 LockSemaphoreInfo(openCL_lock);
1695 if (default_CLEnv == (MagickCLEnv) NULL)
1696 default_CLEnv=AcquireMagickCLEnv();
1697 UnlockSemaphoreInfo(openCL_lock);
1698
1699 return(default_CLEnv);
1700 }
1701
1702 /*
1703 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1704 % %
1705 % %
1706 % %
1707 % 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 %
1708 % %
1709 % %
1710 % %
1711 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1712 %
1713 % GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
1714 % device. The score is determined by the duration of the micro benchmark so
1715 % that means a lower score is better than a higher score.
1716 %
1717 % The format of the GetOpenCLDeviceBenchmarkScore method is:
1718 %
1719 % double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
1720 %
1721 % A description of each parameter follows:
1722 %
1723 % o device: the OpenCL device.
1724 */
1725
GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)1726 MagickExport double GetOpenCLDeviceBenchmarkScore(
1727 const MagickCLDevice device)
1728 {
1729 if (device == (MagickCLDevice) NULL)
1730 return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1731 return(device->score);
1732 }
1733
1734 /*
1735 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1736 % %
1737 % %
1738 % %
1739 % G e t O p e n C L D e v i c e E n a b l e d %
1740 % %
1741 % %
1742 % %
1743 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1744 %
1745 % GetOpenCLDeviceEnabled() returns true if the device is enabled.
1746 %
1747 % The format of the GetOpenCLDeviceEnabled method is:
1748 %
1749 % MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
1750 %
1751 % A description of each parameter follows:
1752 %
1753 % o device: the OpenCL device.
1754 */
1755
GetOpenCLDeviceEnabled(const MagickCLDevice device)1756 MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
1757 const MagickCLDevice device)
1758 {
1759 if (device == (MagickCLDevice) NULL)
1760 return(MagickFalse);
1761 return(device->enabled);
1762 }
1763
1764 /*
1765 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1766 % %
1767 % %
1768 % %
1769 % G e t O p e n C L D e v i c e N a m e %
1770 % %
1771 % %
1772 % %
1773 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1774 %
1775 % GetOpenCLDeviceName() returns the name of the device.
1776 %
1777 % The format of the GetOpenCLDeviceName method is:
1778 %
1779 % const char *GetOpenCLDeviceName(const MagickCLDevice device)
1780 %
1781 % A description of each parameter follows:
1782 %
1783 % o device: the OpenCL device.
1784 */
1785
GetOpenCLDeviceName(const MagickCLDevice device)1786 MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
1787 {
1788 if (device == (MagickCLDevice) NULL)
1789 return((const char *) NULL);
1790 return(device->name);
1791 }
1792
1793 /*
1794 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1795 % %
1796 % %
1797 % %
1798 % G e t O p e n C L D e v i c e s %
1799 % %
1800 % %
1801 % %
1802 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1803 %
1804 % GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
1805 % value of length to the number of devices that are available.
1806 %
1807 % The format of the GetOpenCLDevices method is:
1808 %
1809 % const MagickCLDevice *GetOpenCLDevices(size_t *length,
1810 % ExceptionInfo *exception)
1811 %
1812 % A description of each parameter follows:
1813 %
1814 % o length: the number of device.
1815 %
1816 % o exception: return any errors or warnings in this structure.
1817 %
1818 */
1819
GetOpenCLDevices(size_t * length,ExceptionInfo * exception)1820 MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
1821 ExceptionInfo *exception)
1822 {
1823 MagickCLEnv
1824 clEnv;
1825
1826 clEnv=GetCurrentOpenCLEnv();
1827 if (clEnv == (MagickCLEnv) NULL)
1828 {
1829 if (length != (size_t *) NULL)
1830 *length=0;
1831 return((MagickCLDevice *) NULL);
1832 }
1833 InitializeOpenCL(clEnv,exception);
1834 if (length != (size_t *) NULL)
1835 *length=clEnv->number_devices;
1836 return(clEnv->devices);
1837 }
1838
1839 /*
1840 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1841 % %
1842 % %
1843 % %
1844 % G e t O p e n C L D e v i c e T y p e %
1845 % %
1846 % %
1847 % %
1848 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1849 %
1850 % GetOpenCLDeviceType() returns the type of the device.
1851 %
1852 % The format of the GetOpenCLDeviceType method is:
1853 %
1854 % MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
1855 %
1856 % A description of each parameter follows:
1857 %
1858 % o device: the OpenCL device.
1859 */
1860
GetOpenCLDeviceType(const MagickCLDevice device)1861 MagickExport MagickCLDeviceType GetOpenCLDeviceType(
1862 const MagickCLDevice device)
1863 {
1864 if (device == (MagickCLDevice) NULL)
1865 return(UndefinedCLDeviceType);
1866 if (device->type == CL_DEVICE_TYPE_GPU)
1867 return(GpuCLDeviceType);
1868 if (device->type == CL_DEVICE_TYPE_CPU)
1869 return(CpuCLDeviceType);
1870 return(UndefinedCLDeviceType);
1871 }
1872
1873 /*
1874 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1875 % %
1876 % %
1877 % %
1878 % G e t O p e n C L D e v i c e V e r s i o n %
1879 % %
1880 % %
1881 % %
1882 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1883 %
1884 % GetOpenCLDeviceVersion() returns the version of the device.
1885 %
1886 % The format of the GetOpenCLDeviceName method is:
1887 %
1888 % const char *GetOpenCLDeviceVersion(MagickCLDevice device)
1889 %
1890 % A description of each parameter follows:
1891 %
1892 % o device: the OpenCL device.
1893 */
1894
GetOpenCLDeviceVersion(const MagickCLDevice device)1895 MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
1896 {
1897 if (device == (MagickCLDevice) NULL)
1898 return((const char *) NULL);
1899 return(device->version);
1900 }
1901
1902 /*
1903 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1904 % %
1905 % %
1906 % %
1907 % G e t O p e n C L E n a b l e d %
1908 % %
1909 % %
1910 % %
1911 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1912 %
1913 % GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
1914 %
1915 % The format of the GetOpenCLEnabled method is:
1916 %
1917 % MagickBooleanType GetOpenCLEnabled()
1918 %
1919 */
1920
GetOpenCLEnabled(void)1921 MagickExport MagickBooleanType GetOpenCLEnabled(void)
1922 {
1923 MagickCLEnv
1924 clEnv;
1925
1926 clEnv=GetCurrentOpenCLEnv();
1927 if (clEnv == (MagickCLEnv) NULL)
1928 return(MagickFalse);
1929 return(clEnv->enabled);
1930 }
1931
1932 /*
1933 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1934 % %
1935 % %
1936 % %
1937 % 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 %
1938 % %
1939 % %
1940 % %
1941 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1942 %
1943 % GetOpenCLKernelProfileRecords() returns the profile records for the
1944 % specified device and sets length to the number of profile records.
1945 %
1946 % The format of the GetOpenCLKernelProfileRecords method is:
1947 %
1948 % const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
1949 %
1950 % A description of each parameter follows:
1951 %
1952 % o length: the number of profiles records.
1953 */
1954
GetOpenCLKernelProfileRecords(const MagickCLDevice device,size_t * length)1955 MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
1956 const MagickCLDevice device,size_t *length)
1957 {
1958 if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
1959 (KernelProfileRecord *) NULL))
1960 {
1961 if (length != (size_t *) NULL)
1962 *length=0;
1963 return((const KernelProfileRecord *) NULL);
1964 }
1965 if (length != (size_t *) NULL)
1966 {
1967 *length=0;
1968 LockSemaphoreInfo(device->lock);
1969 while (device->profile_records[*length] != (KernelProfileRecord) NULL)
1970 *length=*length+1;
1971 UnlockSemaphoreInfo(device->lock);
1972 }
1973 return(device->profile_records);
1974 }
1975
1976 /*
1977 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1978 % %
1979 % %
1980 % %
1981 % H a s O p e n C L D e v i c e s %
1982 % %
1983 % %
1984 % %
1985 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1986 %
1987 % HasOpenCLDevices() checks if the OpenCL environment has devices that are
1988 % enabled and compiles the kernel for the device when necessary. False will be
1989 % returned if no enabled devices could be found
1990 %
1991 % The format of the HasOpenCLDevices method is:
1992 %
1993 % MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
1994 % ExceptionInfo exception)
1995 %
1996 % A description of each parameter follows:
1997 %
1998 % o clEnv: the OpenCL environment.
1999 %
2000 % o exception: return any errors or warnings in this structure.
2001 %
2002 */
2003
HasOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo * exception)2004 static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2005 ExceptionInfo *exception)
2006 {
2007 char
2008 *accelerateKernelsBuffer,
2009 options[MagickPathExtent];
2010
2011 MagickStatusType
2012 status;
2013
2014 size_t
2015 i;
2016
2017 size_t
2018 signature;
2019
2020 /* Check if there are enabled devices */
2021 for (i = 0; i < clEnv->number_devices; i++)
2022 {
2023 if ((clEnv->devices[i]->enabled != MagickFalse))
2024 break;
2025 }
2026 if (i == clEnv->number_devices)
2027 return(MagickFalse);
2028
2029 /* Check if we need to compile a kernel for one of the devices */
2030 status=MagickTrue;
2031 for (i = 0; i < clEnv->number_devices; i++)
2032 {
2033 if ((clEnv->devices[i]->enabled != MagickFalse) &&
2034 (clEnv->devices[i]->program == (cl_program) NULL))
2035 {
2036 status=MagickFalse;
2037 break;
2038 }
2039 }
2040 if (status != MagickFalse)
2041 return(MagickTrue);
2042
2043 /* Get additional options */
2044 (void) FormatLocaleString(options,MaxTextExtent,CLOptions,
2045 (float)QuantumRange,(float)QuantumScale,(float)CLCharQuantumScale,
2046 (float)MagickEpsilon,(float)MagickPI,(unsigned int)MaxMap,
2047 (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2048
2049 signature=StringSignature(options);
2050 accelerateKernelsBuffer=(char*) AcquireMagickMemory(
2051 strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2052 if (accelerateKernelsBuffer == (char*) NULL)
2053 return(MagickFalse);
2054 sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
2055 signature^=StringSignature(accelerateKernelsBuffer);
2056
2057 status=MagickTrue;
2058 for (i = 0; i < clEnv->number_devices; i++)
2059 {
2060 MagickCLDevice
2061 device;
2062
2063 size_t
2064 device_signature;
2065
2066 device=clEnv->devices[i];
2067 if ((device->enabled == MagickFalse) ||
2068 (device->program != (cl_program) NULL))
2069 continue;
2070
2071 LockSemaphoreInfo(device->lock);
2072 if (device->program != (cl_program) NULL)
2073 {
2074 UnlockSemaphoreInfo(device->lock);
2075 continue;
2076 }
2077 device_signature=signature;
2078 device_signature^=StringSignature(device->platform_name);
2079 status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2080 device_signature,exception);
2081 UnlockSemaphoreInfo(device->lock);
2082 if (status == MagickFalse)
2083 break;
2084 }
2085 accelerateKernelsBuffer=RelinquishMagickMemory(accelerateKernelsBuffer);
2086 return(status);
2087 }
2088
2089 /*
2090 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2091 % %
2092 % %
2093 % %
2094 + I n i t i a l i z e O p e n C L %
2095 % %
2096 % %
2097 % %
2098 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2099 %
2100 % InitializeOpenCL() is used to initialize the OpenCL environment. This method
2101 % makes sure the devices are propertly initialized and benchmarked.
2102 %
2103 % The format of the InitializeOpenCL method is:
2104 %
2105 % MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2106 %
2107 % A description of each parameter follows:
2108 %
2109 % o exception: return any errors or warnings in this structure.
2110 %
2111 */
2112
GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)2113 static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2114 {
2115 char
2116 version[MagickPathExtent];
2117
2118 cl_uint
2119 num;
2120
2121 if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2122 MagickPathExtent,version,NULL) != CL_SUCCESS)
2123 return(0);
2124 if (strncmp(version, "OpenCL 1.0 ", 11) == 0)
2125 return(0);
2126 if (clEnv->library->clGetDeviceIDs(platform,
2127 CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2128 return(0);
2129 return(num);
2130 }
2131
LoadOpenCLDevices(MagickCLEnv clEnv)2132 static void LoadOpenCLDevices(MagickCLEnv clEnv)
2133 {
2134 cl_context_properties
2135 properties[3];
2136
2137 cl_device_id
2138 *devices;
2139
2140 cl_int
2141 status;
2142
2143 cl_platform_id
2144 *platforms;
2145
2146 cl_uint
2147 i,
2148 j,
2149 next,
2150 number_devices,
2151 number_platforms;
2152
2153 size_t
2154 length;
2155
2156 number_platforms=0;
2157 if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2158 return;
2159 if (number_platforms == 0)
2160 return;
2161 platforms=(cl_platform_id *) AcquireMagickMemory(number_platforms*
2162 sizeof(cl_platform_id));
2163 if (platforms == (cl_platform_id *) NULL)
2164 return;
2165 if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2166 {
2167 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2168 return;
2169 }
2170 for (i = 0; i < number_platforms; i++)
2171 {
2172 number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2173 if (number_devices == 0)
2174 platforms[i]=(cl_platform_id) NULL;
2175 else
2176 clEnv->number_devices+=number_devices;
2177 }
2178 if (clEnv->number_devices == 0)
2179 {
2180 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2181 return;
2182 }
2183 clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2184 sizeof(MagickCLDevice));
2185 if (clEnv->devices == (MagickCLDevice *) NULL)
2186 {
2187 RelinquishMagickCLDevices(clEnv);
2188 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2189 return;
2190 }
2191 (void) ResetMagickMemory(clEnv->devices,0,clEnv->number_devices*
2192 sizeof(MagickCLDevice));
2193 devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2194 sizeof(cl_device_id));
2195 if (devices == (cl_device_id *) NULL)
2196 {
2197 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2198 RelinquishMagickCLDevices(clEnv);
2199 return;
2200 }
2201 clEnv->number_contexts=(size_t) number_platforms;
2202 clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2203 sizeof(cl_context));
2204 if (clEnv->contexts == (cl_context *) NULL)
2205 {
2206 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2207 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2208 RelinquishMagickCLDevices(clEnv);
2209 return;
2210 }
2211 next=0;
2212 for (i = 0; i < number_platforms; i++)
2213 {
2214 if (platforms[i] == (cl_platform_id) NULL)
2215 continue;
2216
2217 status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2218 CL_DEVICE_TYPE_GPU,clEnv->number_devices,devices,&number_devices);
2219 if (status != CL_SUCCESS)
2220 continue;
2221
2222 properties[0]=CL_CONTEXT_PLATFORM;
2223 properties[1]=(cl_context_properties) platforms[i];
2224 properties[2]=0;
2225 clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2226 devices,NULL,NULL,&status);
2227 if (status != CL_SUCCESS)
2228 continue;
2229
2230 for (j = 0; j < number_devices; j++,next++)
2231 {
2232 MagickCLDevice
2233 device;
2234
2235 device=AcquireMagickCLDevice();
2236 if (device == (MagickCLDevice) NULL)
2237 break;
2238
2239 device->context=clEnv->contexts[i];
2240 device->deviceID=devices[j];
2241
2242 openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,0,NULL,
2243 &length);
2244 device->platform_name=AcquireQuantumMemory(length,
2245 sizeof(*device->platform_name));
2246 openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,length,
2247 device->platform_name,NULL);
2248
2249 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_NAME,0,NULL,
2250 &length);
2251 device->name=AcquireQuantumMemory(length,sizeof(*device->name));
2252 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_NAME,length,
2253 device->name,NULL);
2254
2255 openCL_library->clGetDeviceInfo(devices[j],CL_DRIVER_VERSION,0,NULL,
2256 &length);
2257 device->version=AcquireQuantumMemory(length,sizeof(*device->version));
2258 openCL_library->clGetDeviceInfo(devices[j],CL_DRIVER_VERSION,length,
2259 device->version,NULL);
2260
2261 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2262 sizeof(cl_uint),&device->max_clock_frequency,NULL);
2263
2264 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2265 sizeof(cl_uint),&device->max_compute_units,NULL);
2266
2267 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2268 sizeof(cl_device_type),&device->type,NULL);
2269
2270 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2271 sizeof(cl_ulong),&device->local_memory_size,NULL);
2272
2273 clEnv->devices[next]=device;
2274 }
2275 }
2276 if (next != clEnv->number_devices)
2277 RelinquishMagickCLDevices(clEnv);
2278 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2279 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2280 }
2281
InitializeOpenCL(MagickCLEnv clEnv,ExceptionInfo * exception)2282 MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2283 ExceptionInfo *exception)
2284 {
2285 LockSemaphoreInfo(clEnv->lock);
2286 if (clEnv->initialized != MagickFalse)
2287 {
2288 UnlockSemaphoreInfo(clEnv->lock);
2289 return(HasOpenCLDevices(clEnv,exception));
2290 }
2291 if (LoadOpenCLLibrary() != MagickFalse)
2292 {
2293 clEnv->library=openCL_library;
2294 LoadOpenCLDevices(clEnv);
2295 if (clEnv->number_devices > 0)
2296 AutoSelectOpenCLDevices(clEnv,exception);
2297 }
2298 clEnv->initialized=MagickTrue;
2299 UnlockSemaphoreInfo(clEnv->lock);
2300 return(HasOpenCLDevices(clEnv,exception));
2301 }
2302
2303 /*
2304 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2305 % %
2306 % %
2307 % %
2308 % L o a d O p e n C L L i b r a r y %
2309 % %
2310 % %
2311 % %
2312 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2313 %
2314 % LoadOpenCLLibrary() load and binds the OpenCL library.
2315 %
2316 % The format of the LoadOpenCLLibrary method is:
2317 %
2318 % MagickBooleanType LoadOpenCLLibrary(void)
2319 %
2320 */
2321
OsLibraryGetFunctionAddress(void * library,const char * functionName)2322 void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2323 {
2324 if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2325 return (void *) NULL;
2326 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2327 return (void *) GetProcAddress((HMODULE)library,functionName);
2328 #else
2329 return (void *) dlsym(library,functionName);
2330 #endif
2331 }
2332
BindOpenCLFunctions()2333 static MagickBooleanType BindOpenCLFunctions()
2334 {
2335 void
2336 *library;
2337
2338 #ifdef MAGICKCORE_OPENCL_MACOSX
2339 #define BIND(X) openCL_library->X= &X;
2340 #else
2341 (void) ResetMagickMemory(openCL_library,0,sizeof(MagickLibrary));
2342 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2343 library=(void *)LoadLibraryA("OpenCL.dll");
2344 #else
2345 library=(void *)dlopen("libOpenCL.so", RTLD_NOW);
2346 #endif
2347
2348 #define BIND(X) \
2349 if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(library,#X)) == NULL) \
2350 return(MagickFalse);
2351 #endif
2352
2353 BIND(clGetPlatformIDs);
2354 BIND(clGetPlatformInfo);
2355
2356 BIND(clGetDeviceIDs);
2357 BIND(clGetDeviceInfo);
2358
2359 BIND(clCreateBuffer);
2360 BIND(clReleaseMemObject);
2361
2362 BIND(clCreateContext);
2363 BIND(clReleaseContext);
2364
2365 BIND(clCreateCommandQueue);
2366 BIND(clReleaseCommandQueue);
2367 BIND(clFlush);
2368 BIND(clFinish);
2369
2370 BIND(clCreateProgramWithSource);
2371 BIND(clCreateProgramWithBinary);
2372 BIND(clReleaseProgram);
2373 BIND(clBuildProgram);
2374 BIND(clGetProgramBuildInfo);
2375 BIND(clGetProgramInfo);
2376
2377 BIND(clCreateKernel);
2378 BIND(clReleaseKernel);
2379 BIND(clSetKernelArg);
2380 BIND(clGetKernelInfo);
2381
2382 BIND(clEnqueueReadBuffer);
2383 BIND(clEnqueueMapBuffer);
2384 BIND(clEnqueueUnmapMemObject);
2385 BIND(clEnqueueNDRangeKernel);
2386
2387 BIND(clWaitForEvents);
2388 BIND(clReleaseEvent);
2389 BIND(clRetainEvent);
2390 BIND(clSetEventCallback);
2391
2392 BIND(clGetEventProfilingInfo);
2393
2394 return(MagickTrue);
2395 }
2396
LoadOpenCLLibrary(void)2397 static MagickBooleanType LoadOpenCLLibrary(void)
2398 {
2399 openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2400 if (openCL_library == (MagickLibrary *) NULL)
2401 return(MagickFalse);
2402
2403 if (BindOpenCLFunctions() == MagickFalse)
2404 {
2405 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2406 return(MagickFalse);
2407 }
2408
2409 return(MagickTrue);
2410 }
2411
2412 /*
2413 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2414 % %
2415 % %
2416 % %
2417 + O p e n C L T e r m i n u s %
2418 % %
2419 % %
2420 % %
2421 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2422 %
2423 % AnnotateComponentTerminus() destroys the annotate component.
2424 %
2425 % The format of the AnnotateComponentTerminus method is:
2426 %
2427 % AnnotateComponentTerminus(void)
2428 %
2429 */
2430
OpenCLTerminus()2431 MagickPrivate void OpenCLTerminus()
2432 {
2433 DumpOpenCLProfileData();
2434 if (cache_directory != (char *) NULL)
2435 cache_directory=DestroyString(cache_directory);
2436 if (cache_directory_lock != (SemaphoreInfo *) NULL)
2437 RelinquishSemaphoreInfo(&cache_directory_lock);
2438 if (default_CLEnv != (MagickCLEnv) NULL)
2439 default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2440 if (openCL_lock != (SemaphoreInfo *) NULL)
2441 RelinquishSemaphoreInfo(&openCL_lock);
2442 if (openCL_library != (MagickLibrary *) NULL)
2443 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2444 }
2445
2446 /*
2447 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2448 % %
2449 % %
2450 % %
2451 + 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 %
2452 % %
2453 % %
2454 % %
2455 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2456 %
2457 % OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2458 % configuration file. If an error occurs, MagickFalse is returned
2459 % otherwise MagickTrue.
2460 %
2461 % The format of the OpenCLThrowMagickException method is:
2462 %
2463 % MagickBooleanType ThrowFileException(ExceptionInfo *exception,
2464 % const char *module,const char *function,const size_t line,
2465 % const ExceptionType severity,const char *tag,const char *format,...)
2466 %
2467 % A description of each parameter follows:
2468 %
2469 % o exception: the exception info.
2470 %
2471 % o filename: the source module filename.
2472 %
2473 % o function: the function name.
2474 %
2475 % o line: the line number of the source module.
2476 %
2477 % o severity: Specifies the numeric error category.
2478 %
2479 % o tag: the locale tag.
2480 %
2481 % o format: the output format.
2482 %
2483 */
2484
OpenCLThrowMagickException(MagickCLDevice device,ExceptionInfo * exception,const char * module,const char * function,const size_t line,const ExceptionType severity,const char * tag,const char * format,...)2485 MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2486 MagickCLDevice device,ExceptionInfo *exception,const char *module,
2487 const char *function,const size_t line,const ExceptionType severity,
2488 const char *tag,const char *format,...)
2489 {
2490 MagickBooleanType
2491 status;
2492
2493 assert(device != (MagickCLDevice) NULL);
2494 assert(exception != (ExceptionInfo *) NULL);
2495 assert(exception->signature == MagickCoreSignature);
2496
2497 status=MagickTrue;
2498 if (severity != 0)
2499 {
2500 if (device->type == CL_DEVICE_TYPE_CPU)
2501 {
2502 /* Workaround for Intel OpenCL CPU runtime bug */
2503 /* Turn off OpenCL when a problem is detected! */
2504 if (strncmp(device->platform_name, "Intel",5) == 0)
2505 default_CLEnv->enabled=MagickFalse;
2506 }
2507 }
2508
2509 #ifdef OPENCLLOG_ENABLED
2510 {
2511 va_list
2512 operands;
2513 va_start(operands,format);
2514 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2515 format,operands);
2516 va_end(operands);
2517 }
2518 #else
2519 magick_unreferenced(module);
2520 magick_unreferenced(function);
2521 magick_unreferenced(line);
2522 magick_unreferenced(tag);
2523 magick_unreferenced(format);
2524 #endif
2525
2526 return(status);
2527 }
2528
2529 /*
2530 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2531 % %
2532 % %
2533 % %
2534 + R e c o r d P r o f i l e D a t a %
2535 % %
2536 % %
2537 % %
2538 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2539 %
2540 % RecordProfileData() records profile data.
2541 %
2542 % The format of the RecordProfileData method is:
2543 %
2544 % void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2545 % cl_event event)
2546 %
2547 % A description of each parameter follows:
2548 %
2549 % o device: the OpenCL device that did the operation.
2550 %
2551 % o event: the event that contains the profiling data.
2552 %
2553 */
2554
RecordProfileData(MagickCLDevice device,cl_kernel kernel,cl_event event)2555 MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2556 cl_kernel kernel,cl_event event)
2557 {
2558 char
2559 *name;
2560
2561 cl_int
2562 status;
2563
2564 cl_ulong
2565 elapsed,
2566 end,
2567 start;
2568
2569 KernelProfileRecord
2570 profile_record;
2571
2572 size_t
2573 i,
2574 length;
2575
2576 if (device->profile_kernels == MagickFalse)
2577 return(MagickFalse);
2578 status=openCL_library->clWaitForEvents(1,&event);
2579 if (status != CL_SUCCESS)
2580 return(MagickFalse);
2581 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2582 &length);
2583 if (status != CL_SUCCESS)
2584 return(MagickTrue);
2585 name=AcquireQuantumMemory(length,sizeof(*name));
2586 if (name == (char *) NULL)
2587 return(MagickTrue);
2588 start=end=elapsed=0;
2589 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2590 name,(size_t *) NULL);
2591 status|=openCL_library->clGetEventProfilingInfo(event,
2592 CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2593 status|=openCL_library->clGetEventProfilingInfo(event,
2594 CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2595 if (status != CL_SUCCESS)
2596 {
2597 name=DestroyString(name);
2598 return(MagickTrue);
2599 }
2600 start/=1000; // usecs
2601 end/=1000; // usecs
2602 elapsed=end-start;
2603 LockSemaphoreInfo(device->lock);
2604 i=0;
2605 profile_record=(KernelProfileRecord) NULL;
2606 if (device->profile_records != (KernelProfileRecord *) NULL)
2607 {
2608 while (device->profile_records[i] != (KernelProfileRecord) NULL)
2609 {
2610 if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2611 {
2612 profile_record=device->profile_records[i];
2613 break;
2614 }
2615 i++;
2616 }
2617 }
2618 if (profile_record != (KernelProfileRecord) NULL)
2619 name=DestroyString(name);
2620 else
2621 {
2622 profile_record=AcquireMagickMemory(sizeof(*profile_record));
2623 (void) ResetMagickMemory(profile_record,0,sizeof(*profile_record));
2624 profile_record->kernel_name=name;
2625 device->profile_records=ResizeMagickMemory(device->profile_records,(i+2)*
2626 sizeof(*device->profile_records));
2627 device->profile_records[i]=profile_record;
2628 device->profile_records[i+1]=(KernelProfileRecord) NULL;
2629 }
2630 if ((elapsed < profile_record->min) || (profile_record->count == 0))
2631 profile_record->min=elapsed;
2632 if (elapsed > profile_record->max)
2633 profile_record->max=elapsed;
2634 profile_record->total+=elapsed;
2635 profile_record->count+=1;
2636 UnlockSemaphoreInfo(device->lock);
2637 return(MagickTrue);
2638 }
2639
2640 /*
2641 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2642 % %
2643 % %
2644 % %
2645 + 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 %
2646 % %
2647 % %
2648 % %
2649 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2650 %
2651 % ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2652 %
2653 % The format of the ReleaseOpenCLCommandQueue method is:
2654 %
2655 % void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2656 % cl_command_queue queue)
2657 %
2658 % A description of each parameter follows:
2659 %
2660 % o device: the OpenCL device.
2661 %
2662 % o queue: the OpenCL queue to be released.
2663 */
2664
ReleaseOpenCLCommandQueue(MagickCLDevice device,cl_command_queue queue)2665 MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2666 cl_command_queue queue)
2667 {
2668 if (queue == (cl_command_queue) NULL)
2669 return;
2670
2671 assert(device != (MagickCLDevice) NULL);
2672 LockSemaphoreInfo(device->lock);
2673 if ((device->profile_kernels != MagickFalse) ||
2674 (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2675 {
2676 UnlockSemaphoreInfo(device->lock);
2677 openCL_library->clFinish(queue);
2678 (void) openCL_library->clReleaseCommandQueue(queue);
2679 }
2680 else
2681 {
2682 openCL_library->clFlush(queue);
2683 device->command_queues[++device->command_queues_index]=queue;
2684 UnlockSemaphoreInfo(device->lock);
2685 }
2686 }
2687
2688 /*
2689 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2690 % %
2691 % %
2692 % %
2693 + R e l e a s e M a g i c k C L D e v i c e %
2694 % %
2695 % %
2696 % %
2697 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2698 %
2699 % ReleaseOpenCLDevice() returns the OpenCL device to the environment
2700 %
2701 % The format of the ReleaseOpenCLDevice method is:
2702 %
2703 % void ReleaseOpenCLDevice(MagickCLDevice device)
2704 %
2705 % A description of each parameter follows:
2706 %
2707 % o device: the OpenCL device to be released.
2708 %
2709 */
2710
ReleaseOpenCLDevice(MagickCLDevice device)2711 MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2712 {
2713 assert(device != (MagickCLDevice) NULL);
2714 LockSemaphoreInfo(openCL_lock);
2715 device->requested--;
2716 UnlockSemaphoreInfo(openCL_lock);
2717 }
2718
2719 /*
2720 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2721 % %
2722 % %
2723 % %
2724 + 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 %
2725 % %
2726 % %
2727 % %
2728 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2729 %
2730 % RelinquishMagickCLCacheInfo() frees memory acquired with
2731 % AcquireMagickCLCacheInfo()
2732 %
2733 % The format of the RelinquishMagickCLCacheInfo method is:
2734 %
2735 % MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2736 % const MagickBooleanType relinquish_pixels)
2737 %
2738 % A description of each parameter follows:
2739 %
2740 % o info: the OpenCL cache info.
2741 %
2742 % o relinquish_pixels: the pixels will be relinquish when set to true.
2743 %
2744 */
DestroyMagickCLCacheInfo(MagickCLCacheInfo info)2745 static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
2746 {
2747 ssize_t
2748 i;
2749
2750 for (i=0; i < (ssize_t) info->event_count; i++)
2751 openCL_library->clReleaseEvent(info->events[i]);
2752 info->events=(cl_event *) RelinquishMagickMemory(info->events);
2753 if (info->buffer != (cl_mem) NULL)
2754 openCL_library->clReleaseMemObject(info->buffer);
2755 ReleaseOpenCLDevice(info->device);
2756 RelinquishMagickMemory(info);
2757 }
2758
DestroyMagickCLCacheInfoAndPixels(cl_event magick_unused (event),cl_int magick_unused (event_command_exec_status),void * user_data)2759 static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2760 cl_event magick_unused(event),
2761 cl_int magick_unused(event_command_exec_status),void *user_data)
2762 {
2763 MagickCLCacheInfo
2764 info;
2765
2766 magick_unreferenced(event);
2767 magick_unreferenced(event_command_exec_status);
2768 info=(MagickCLCacheInfo) user_data;
2769 (void) RelinquishAlignedMemory(info->pixels);
2770 RelinquishMagickResource(MemoryResource,info->length);
2771 DestroyMagickCLCacheInfo(info);
2772 }
2773
RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)2774 MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2775 MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2776 {
2777 if (info == (MagickCLCacheInfo) NULL)
2778 return((MagickCLCacheInfo) NULL);
2779 if (relinquish_pixels != MagickFalse)
2780 {
2781 if (info->event_count > 0)
2782 openCL_library->clSetEventCallback(info->events[info->event_count-1],
2783 CL_COMPLETE,&DestroyMagickCLCacheInfoAndPixels,info);
2784 else
2785 DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2786 }
2787 else
2788 DestroyMagickCLCacheInfo(info);
2789 return((MagickCLCacheInfo) NULL);
2790 }
2791
2792 /*
2793 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2794 % %
2795 % %
2796 % %
2797 % R e l i n q u i s h M a g i c k C L D e v i c e %
2798 % %
2799 % %
2800 % %
2801 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2802 %
2803 % RelinquishMagickCLDevice() releases the OpenCL device
2804 %
2805 % The format of the RelinquishMagickCLDevice method is:
2806 %
2807 % MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2808 %
2809 % A description of each parameter follows:
2810 %
2811 % o device: the OpenCL device to be released.
2812 %
2813 */
2814
RelinquishMagickCLDevice(MagickCLDevice device)2815 static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2816 {
2817 if (device == (MagickCLDevice) NULL)
2818 return((MagickCLDevice) NULL);
2819
2820 device->platform_name=RelinquishMagickMemory(device->platform_name);
2821 device->name=RelinquishMagickMemory(device->name);
2822 device->version=RelinquishMagickMemory(device->version);
2823 if (device->program != (cl_program) NULL)
2824 (void) openCL_library->clReleaseProgram(device->program);
2825 while (device->command_queues_index >= 0)
2826 (void) openCL_library->clReleaseCommandQueue(
2827 device->command_queues[device->command_queues_index--]);
2828 RelinquishSemaphoreInfo(&device->lock);
2829 return((MagickCLDevice) RelinquishMagickMemory(device));
2830 }
2831
2832 /*
2833 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2834 % %
2835 % %
2836 % %
2837 % R e l i n q u i s h M a g i c k C L E n v %
2838 % %
2839 % %
2840 % %
2841 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2842 %
2843 % RelinquishMagickCLEnv() releases the OpenCL environment
2844 %
2845 % The format of the RelinquishMagickCLEnv method is:
2846 %
2847 % MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
2848 %
2849 % A description of each parameter follows:
2850 %
2851 % o clEnv: the OpenCL environment to be released.
2852 %
2853 */
2854
RelinquishMagickCLEnv(MagickCLEnv clEnv)2855 static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
2856 {
2857 if (clEnv == (MagickCLEnv) NULL)
2858 return((MagickCLEnv) NULL);
2859
2860 RelinquishSemaphoreInfo(&clEnv->lock);
2861 RelinquishMagickCLDevices(clEnv);
2862 if (clEnv->contexts != (cl_context *) NULL)
2863 {
2864 ssize_t
2865 i;
2866
2867 for (i=0; i < clEnv->number_contexts; i++)
2868 (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
2869 clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
2870 }
2871 return((MagickCLEnv) RelinquishMagickMemory(clEnv));
2872 }
2873
2874 /*
2875 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2876 % %
2877 % %
2878 % %
2879 + R e q u e s t O p e n C L D e v i c e %
2880 % %
2881 % %
2882 % %
2883 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2884 %
2885 % RequestOpenCLDevice() returns one of the enabled OpenCL devices.
2886 %
2887 % The format of the RequestOpenCLDevice method is:
2888 %
2889 % MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
2890 %
2891 % A description of each parameter follows:
2892 %
2893 % o clEnv: the OpenCL environment.
2894 */
2895
RequestOpenCLDevice(MagickCLEnv clEnv)2896 MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
2897 {
2898 MagickCLDevice
2899 device;
2900
2901 double
2902 score,
2903 best_score;
2904
2905 size_t
2906 i;
2907
2908 if (clEnv == (MagickCLEnv) NULL)
2909 return((MagickCLDevice) NULL);
2910
2911 if (clEnv->number_devices == 1)
2912 {
2913 if (clEnv->devices[0]->enabled)
2914 return(clEnv->devices[0]);
2915 else
2916 return((MagickCLDevice) NULL);
2917 }
2918
2919 device=(MagickCLDevice) NULL;
2920 best_score=0.0;
2921 LockSemaphoreInfo(openCL_lock);
2922 for (i = 0; i < clEnv->number_devices; i++)
2923 {
2924 if (clEnv->devices[i]->enabled == MagickFalse)
2925 continue;
2926
2927 score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
2928 clEnv->devices[i]->requested);
2929 if ((device == (MagickCLDevice) NULL) || (score < best_score))
2930 {
2931 device=clEnv->devices[i];
2932 best_score=score;
2933 }
2934 }
2935 if (device != (MagickCLDevice)NULL)
2936 device->requested++;
2937 UnlockSemaphoreInfo(openCL_lock);
2938
2939 return(device);
2940 }
2941
2942 /*
2943 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2944 % %
2945 % %
2946 % %
2947 % S e t O p e n C L D e v i c e E n a b l e d %
2948 % %
2949 % %
2950 % %
2951 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2952 %
2953 % SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
2954 %
2955 % The format of the SetOpenCLDeviceEnabled method is:
2956 %
2957 % void SetOpenCLDeviceEnabled(MagickCLDevice device,
2958 % MagickBooleanType value)
2959 %
2960 % A description of each parameter follows:
2961 %
2962 % o device: the OpenCL device.
2963 %
2964 % o value: determines if the device should be enabled or disabled.
2965 */
2966
SetOpenCLDeviceEnabled(MagickCLDevice device,const MagickBooleanType value)2967 MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
2968 const MagickBooleanType value)
2969 {
2970 if (device == (MagickCLDevice) NULL)
2971 return;
2972 device->enabled=value;
2973 }
2974
2975 /*
2976 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2977 % %
2978 % %
2979 % %
2980 % 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 %
2981 % %
2982 % %
2983 % %
2984 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2985 %
2986 % SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
2987 % kernel profiling of a device.
2988 %
2989 % The format of the SetOpenCLKernelProfileEnabled method is:
2990 %
2991 % void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
2992 % MagickBooleanType value)
2993 %
2994 % A description of each parameter follows:
2995 %
2996 % o device: the OpenCL device.
2997 %
2998 % o value: determines if kernel profiling for the device should be enabled
2999 % or disabled.
3000 */
3001
SetOpenCLKernelProfileEnabled(MagickCLDevice device,const MagickBooleanType value)3002 MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3003 const MagickBooleanType value)
3004 {
3005 if (device == (MagickCLDevice) NULL)
3006 return;
3007 device->profile_kernels=value;
3008 }
3009
3010 /*
3011 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3012 % %
3013 % %
3014 % %
3015 % S e t O p e n C L E n a b l e d %
3016 % %
3017 % %
3018 % %
3019 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3020 %
3021 % SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3022 %
3023 % The format of the SetOpenCLEnabled method is:
3024 %
3025 % void SetOpenCLEnabled(MagickBooleanType)
3026 %
3027 % A description of each parameter follows:
3028 %
3029 % o value: specify true to enable OpenCL acceleration
3030 */
3031
SetOpenCLEnabled(const MagickBooleanType value)3032 MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
3033 {
3034 MagickCLEnv
3035 clEnv;
3036
3037 clEnv=GetCurrentOpenCLEnv();
3038 if (clEnv == (MagickCLEnv) NULL)
3039 return(MagickFalse);
3040 clEnv->enabled=value;
3041 return(clEnv->enabled);
3042 }
3043
3044 #else
3045
GetOpenCLDeviceBenchmarkScore(const MagickCLDevice magick_unused (device))3046 MagickExport double GetOpenCLDeviceBenchmarkScore(
3047 const MagickCLDevice magick_unused(device))
3048 {
3049 magick_unreferenced(device);
3050 return(0.0);
3051 }
3052
GetOpenCLDeviceEnabled(const MagickCLDevice magick_unused (device))3053 MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3054 const MagickCLDevice magick_unused(device))
3055 {
3056 magick_unreferenced(device);
3057 return(MagickFalse);
3058 }
3059
GetOpenCLDeviceName(const MagickCLDevice magick_unused (device))3060 MagickExport const char *GetOpenCLDeviceName(
3061 const MagickCLDevice magick_unused(device))
3062 {
3063 magick_unreferenced(device);
3064 return((const char *) NULL);
3065 }
3066
GetOpenCLDevices(size_t * length,ExceptionInfo * magick_unused (exception))3067 MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
3068 ExceptionInfo *magick_unused(exception))
3069 {
3070 magick_unreferenced(exception);
3071 if (length != (size_t *) NULL)
3072 *length=0;
3073 return((MagickCLDevice *) NULL);
3074 }
3075
GetOpenCLDeviceType(const MagickCLDevice magick_unused (device))3076 MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3077 const MagickCLDevice magick_unused(device))
3078 {
3079 magick_unreferenced(device);
3080 return(UndefinedCLDeviceType);
3081 }
3082
GetOpenCLKernelProfileRecords(const MagickCLDevice magick_unused (device),size_t * length)3083 MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
3084 const MagickCLDevice magick_unused(device),size_t *length)
3085 {
3086 magick_unreferenced(device);
3087 if (length != (size_t *) NULL)
3088 *length=0;
3089 return((const KernelProfileRecord *) NULL);
3090 }
3091
GetOpenCLDeviceVersion(const MagickCLDevice magick_unused (device))3092 MagickExport const char *GetOpenCLDeviceVersion(
3093 const MagickCLDevice magick_unused(device))
3094 {
3095 magick_unreferenced(device);
3096 return((const char *) NULL);
3097 }
3098
GetOpenCLEnabled(void)3099 MagickExport MagickBooleanType GetOpenCLEnabled(void)
3100 {
3101 return(MagickFalse);
3102 }
3103
SetOpenCLDeviceEnabled(MagickCLDevice magick_unused (device),const MagickBooleanType magick_unused (value))3104 MagickExport void SetOpenCLDeviceEnabled(
3105 MagickCLDevice magick_unused(device),
3106 const MagickBooleanType magick_unused(value))
3107 {
3108 magick_unreferenced(device);
3109 magick_unreferenced(value);
3110 }
3111
SetOpenCLEnabled(const MagickBooleanType magick_unused (value))3112 MagickExport MagickBooleanType SetOpenCLEnabled(
3113 const MagickBooleanType magick_unused(value))
3114 {
3115 magick_unreferenced(value);
3116 return(MagickFalse);
3117 }
3118
SetOpenCLKernelProfileEnabled(MagickCLDevice magick_unused (device),const MagickBooleanType magick_unused (value))3119 MagickExport void SetOpenCLKernelProfileEnabled(
3120 MagickCLDevice magick_unused(device),
3121 const MagickBooleanType magick_unused(value))
3122 {
3123 magick_unreferenced(device);
3124 magick_unreferenced(value);
3125 }
3126 #endif