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