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