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/script/license.php %
27% %
28% Unless required by applicable law or agreed to in writing, software %
29% distributed under the License is distributed on an "AS IS" BASIS, %
30% WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. %
31% See the License for the specific language governing permissions and %
32% limitations under the License. %
33% %
34%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
35%
36%
37%
38*/
39␌
40/*
41 Include declarations.
42*/
43#include "magick/studio.h"
44#include "magick/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 for (i = 0; i < num_devices; i++)
914 {
915 binary_program[i]=(unsigned char *) AcquireQuantumMemory(
916 MagickMax(*(program_sizes+i),1),sizeof(**binary_program));
917 if (binary_program[i] == (unsigned char *) NULL)
918 {
919 status=CL_OUT_OF_HOST_MEMORY;
920 break;
921 }
922 }
923 if (status == CL_SUCCESS)
924 status=clEnv->library->clGetProgramInfo(clEnv->programs[prog],
925 CL_PROGRAM_BINARIES,binary_program_size,binary_program,NULL);
926 if (status == CL_SUCCESS)
927 {
928 for (i = 0; i < num_devices; i++)
929 {
930 int
931 file;
932
933 size_t
934 program_size;
935
936 program_size=*(program_sizes+i);
937 if (program_size < 1)
938 continue;
939 file=open_utf8(filename,O_WRONLY | O_CREAT | O_BINARY,S_MODE);
940 if (file != -1)
941 {
942 write(file,binary_program[i],program_size);
943 file=close(file);
944 }
945 else
946 (void) ThrowMagickException(exception,GetMagickModule(),
947 DelegateWarning,"Saving kernel failed.","`%s'",filename);
948 break;
949 }
950 }
951 for (i = 0; i < num_devices; i++)
952 binary_program[i]=(unsigned char *) RelinquishMagickMemory(
953 binary_program[i]);
954 binary_program=(unsigned char **) RelinquishMagickMemory(binary_program);
955 }
956 program_sizes=(size_t *) RelinquishMagickMemory(program_sizes);
957}
958
959static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
960{
961 MagickBooleanType loadSuccessful;
962 unsigned char* binaryProgram;
963 char* binaryFileName;
964 FILE* fileHandle;
965
966#ifdef MAGICKCORE_CLPERFMARKER
967 clBeginPerfMarkerAMD(__FUNCTION__,"");
968#endif
969
970 binaryProgram = NULL;
971 binaryFileName = NULL;
972 fileHandle = NULL;
973 loadSuccessful = MagickFalse;
974
975 binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
976 fileHandle = fopen(binaryFileName, "rb");
977 if (fileHandle != NULL)
978 {
979 int b_error;
980 size_t length;
981 cl_int clStatus;
982 cl_int clBinaryStatus;
983
984 b_error = 0 ;
985 length = 0;
986 b_error |= fseek( fileHandle, 0, SEEK_END ) < 0;
987 b_error |= ( length = ftell( fileHandle ) ) <= 0;
988 b_error |= fseek( fileHandle, 0, SEEK_SET ) < 0;
989 if( b_error )
990 goto cleanup;
991
992 binaryProgram = (unsigned char*)AcquireMagickMemory(length);
993 if (binaryProgram == NULL)
994 goto cleanup;
995
996 memset(binaryProgram, 0, length);
997 b_error |= fread(binaryProgram, 1, length, fileHandle) != length;
998
999 clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
1000 if (clStatus != CL_SUCCESS
1001 || clBinaryStatus != CL_SUCCESS)
1002 goto cleanup;
1003
1004 loadSuccessful = MagickTrue;
1005 }
1006
1007cleanup:
1008 if (fileHandle != NULL)
1009 fclose(fileHandle);
1010 if (binaryFileName != NULL)
1011 RelinquishMagickMemory(binaryFileName);
1012 if (binaryProgram != NULL)
1013 RelinquishMagickMemory(binaryProgram);
1014
1015#ifdef MAGICKCORE_CLPERFMARKER
1016 clEndPerfMarkerAMD();
1017#endif
1018
1019 return loadSuccessful;
1020}
1021
1022static unsigned int stringSignature(const char* string)
1023{
1024 unsigned int stringLength;
1025 unsigned int n,i,j;
1026 unsigned int signature;
1027 union
1028 {
1029 const char* s;
1030 const unsigned int* u;
1031 }p;
1032
1033#ifdef MAGICKCORE_CLPERFMARKER
1034 clBeginPerfMarkerAMD(__FUNCTION__,"");
1035#endif
1036
1037 stringLength = (unsigned int) strlen(string);
1038 signature = stringLength;
1039 n = stringLength/sizeof(unsigned int);
1040 p.s = string;
1041 for (i = 0; i < n; i++)
1042 {
1043 signature^=p.u[i];
1044 }
1045 if (n * sizeof(unsigned int) != stringLength)
1046 {
1047 char padded[4];
1048 j = n * sizeof(unsigned int);
1049 for (i = 0; i < 4; i++,j++)
1050 {
1051 if (j < stringLength)
1052 padded[i] = p.s[j];
1053 else
1054 padded[i] = 0;
1055 }
1056 p.s = padded;
1057 signature^=p.u[0];
1058 }
1059
1060#ifdef MAGICKCORE_CLPERFMARKER
1061 clEndPerfMarkerAMD();
1062#endif
1063
1064 return signature;
1065}
1066
1067/* OpenCL kernels for accelerate.c */
1068extern const char *accelerateKernels, *accelerateKernels2;
1069
1070static MagickBooleanType CompileOpenCLKernels(MagickCLEnv clEnv, ExceptionInfo* exception)
1071{
1072 MagickBooleanType status = MagickFalse;
1073 cl_int clStatus;
1074 unsigned int i;
1075 char* accelerateKernelsBuffer = NULL;
1076
1077 /* The index of the program strings in this array has to match the value of the enum MagickOpenCLProgram */
1078 const char* MagickOpenCLProgramStrings[MAGICK_OPENCL_NUM_PROGRAMS];
1079
1080 char options[MaxTextExtent];
1081 unsigned int optionsSignature;
1082
1083#ifdef MAGICKCORE_CLPERFMARKER
1084 clBeginPerfMarkerAMD(__FUNCTION__,"");
1085#endif
1086
1087 /* Get additional options */
1088 (void) FormatLocaleString(options, MaxTextExtent, CLOptions, (float)QuantumRange,
1089 (float)QuantumScale, (float)CLCharQuantumScale, (float)MagickEpsilon, (float)MagickPI, (unsigned int)MaxMap, (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
1090
1091 /*
1092 if (getenv("MAGICK_OCL_DEF"))
1093 {
1094 strcat(options," ");
1095 strcat(options,getenv("MAGICK_OCL_DEF"));
1096 }
1097 */
1098
1099 /*
1100 if (getenv("MAGICK_OCL_BUILD"))
1101 printf("options: %s\n", options);
1102 */
1103
1104 optionsSignature = stringSignature(options);
1105
1106 /* get all the OpenCL program strings here */
1107 accelerateKernelsBuffer = (char*) AcquireQuantumMemory(1,strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
1108 FormatLocaleString(accelerateKernelsBuffer,strlen(accelerateKernels)+
1109 strlen(accelerateKernels2)+1,"%s%s",accelerateKernels,accelerateKernels2);
1110 MagickOpenCLProgramStrings[MAGICK_OPENCL_ACCELERATE] = accelerateKernelsBuffer;
1111
1112 for (i = 0; i < MAGICK_OPENCL_NUM_PROGRAMS; i++)
1113 {
1114 MagickBooleanType loadSuccessful = MagickFalse;
1115 unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature;
1116
1117 /* try to load the binary first */
1118 if (clEnv->disableProgramCache != MagickTrue
1119 && !getenv("MAGICK_OCL_REC"))
1120 loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature);
1121
1122 if (loadSuccessful == MagickFalse)
1123 {
1124 /* Binary CL program unavailable, compile the program from source */
1125 size_t programLength = strlen(MagickOpenCLProgramStrings[i]);
1126 clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
1127 if (clStatus!=CL_SUCCESS)
1128 {
1129 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1130 "clCreateProgramWithSource failed.", "(%d)", (int)clStatus);
1131
1132 goto cleanup;
1133 }
1134 }
1135
1136 clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
1137 if (clStatus!=CL_SUCCESS)
1138 {
1139 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1140 "clBuildProgram failed.", "(%d)", (int)clStatus);
1141
1142 if (loadSuccessful == MagickFalse)
1143 {
1144 char path[MaxTextExtent];
1145 FILE* fileHandle;
1146
1147 /* dump the source into a file */
1148 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
1149 ,GetOpenCLCachedFilesDirectory()
1150 ,DirectorySeparator,"magick_badcl.cl");
1151 fileHandle = fopen(path, "wb");
1152 if (fileHandle != NULL)
1153 {
1154 fwrite(MagickOpenCLProgramStrings[i], sizeof(char), strlen(MagickOpenCLProgramStrings[i]), fileHandle);
1155 fclose(fileHandle);
1156 }
1157
1158 /* dump the build log */
1159 {
1160 char* log;
1161 size_t logSize;
1162 clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
1163 log = (char*)AcquireCriticalMemory(logSize);
1164 clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);
1165
1166 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
1167 ,GetOpenCLCachedFilesDirectory()
1168 ,DirectorySeparator,"magick_badcl_build.log");
1169 fileHandle = fopen(path, "wb");
1170 if (fileHandle != NULL)
1171 {
1172 const char* buildOptionsTitle = "build options: ";
1173 fwrite(buildOptionsTitle, sizeof(char), strlen(buildOptionsTitle), fileHandle);
1174 fwrite(options, sizeof(char), strlen(options), fileHandle);
1175 fwrite("\n",sizeof(char), 1, fileHandle);
1176 fwrite(log, sizeof(char), logSize, fileHandle);
1177 fclose(fileHandle);
1178 }
1179 RelinquishMagickMemory(log);
1180 }
1181 }
1182 goto cleanup;
1183 }
1184
1185 if (loadSuccessful == MagickFalse)
1186 {
1187 /* Save the binary to a file to avoid re-compilation of the kernels in the future */
1188 saveBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception);
1189 }
1190
1191 }
1192 status = MagickTrue;
1193
1194cleanup:
1195
1196 if (accelerateKernelsBuffer!=NULL) RelinquishMagickMemory(accelerateKernelsBuffer);
1197
1198#ifdef MAGICKCORE_CLPERFMARKER
1199 clEndPerfMarkerAMD();
1200#endif
1201
1202 return status;
1203}
1204
1205static MagickBooleanType InitOpenCLPlatformDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
1206 int i,j;
1207 cl_int status;
1208 cl_uint numPlatforms = 0;
1209 cl_platform_id *platforms = NULL;
1210 char* MAGICK_OCL_DEVICE = NULL;
1211 MagickBooleanType OpenCLAvailable = MagickFalse;
1212
1213#ifdef MAGICKCORE_CLPERFMARKER
1214 clBeginPerfMarkerAMD(__FUNCTION__,"");
1215#endif
1216
1217 /* check if there's an environment variable overriding the device selection */
1218 MAGICK_OCL_DEVICE = getenv("MAGICK_OCL_DEVICE");
1219 if (MAGICK_OCL_DEVICE == (char *) NULL)
1220 return(MagickFalse);
1221 if (strcmp(MAGICK_OCL_DEVICE, "CPU") == 0)
1222 clEnv->deviceType = CL_DEVICE_TYPE_CPU;
1223 else if (strcmp(MAGICK_OCL_DEVICE, "GPU") == 0)
1224 clEnv->deviceType = CL_DEVICE_TYPE_GPU;
1225 else if (IsStringTrue(MAGICK_OCL_DEVICE) != MagickFalse)
1226 {
1227 if (clEnv->deviceType == 0)
1228 clEnv->deviceType = CL_DEVICE_TYPE_ALL;
1229 }
1230 else
1231 return(MagickFalse);
1232
1233 if (clEnv->device != NULL)
1234 {
1235 status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
1236 if (status != CL_SUCCESS) {
1237 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1238 "Failed to get OpenCL platform from the selected device.", "(%d)", status);
1239 }
1240 goto cleanup;
1241 }
1242 else if (clEnv->platform != NULL)
1243 {
1244 numPlatforms = 1;
1245 platforms = (cl_platform_id *) AcquireQuantumMemory(1,numPlatforms * sizeof(cl_platform_id));
1246 if (platforms == (cl_platform_id *) NULL)
1247 {
1248 (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1249 "AcquireMagickMemory failed.",".");
1250 goto cleanup;
1251 }
1252 platforms[0] = clEnv->platform;
1253 }
1254 else
1255 {
1256 clEnv->device = NULL;
1257
1258 /* Get the number of OpenCL platforms available */
1259 status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms);
1260 if (status != CL_SUCCESS)
1261 {
1262 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1263 "clGetplatformIDs failed.", "(%d)", status);
1264 goto cleanup;
1265 }
1266
1267 /* No OpenCL available, just leave */
1268 if (numPlatforms == 0) {
1269 goto cleanup;
1270 }
1271
1272 platforms = (cl_platform_id *) AcquireQuantumMemory(1,numPlatforms * sizeof(cl_platform_id));
1273 if (platforms == (cl_platform_id *) NULL)
1274 {
1275 (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1276 "AcquireMagickMemory failed.",".");
1277 goto cleanup;
1278 }
1279
1280 status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL);
1281 if (status != CL_SUCCESS)
1282 {
1283 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1284 "clGetPlatformIDs failed.", "(%d)", status);
1285 goto cleanup;
1286 }
1287 }
1288
1289 /* Device selection */
1290 clEnv->device = NULL;
1291 for (j = 0; j < 2; j++)
1292 {
1293
1294 cl_device_type deviceType;
1295 if (clEnv->deviceType == CL_DEVICE_TYPE_ALL)
1296 {
1297 if (j == 0)
1298 deviceType = CL_DEVICE_TYPE_GPU;
1299 else
1300 deviceType = CL_DEVICE_TYPE_CPU;
1301 }
1302 else if (j == 1)
1303 {
1304 break;
1305 }
1306 else
1307 deviceType = clEnv->deviceType;
1308
1309 for (i = 0; i < numPlatforms; i++)
1310 {
1311 char version[MaxTextExtent];
1312 cl_uint numDevices;
1313 status = clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_VERSION, MaxTextExtent, version, NULL);
1314 if (status != CL_SUCCESS)
1315 {
1316 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1317 "clGetPlatformInfo failed.", "(%d)", status);
1318 goto cleanup;
1319 }
1320 if (strncmp(version,"OpenCL 1.0 ",11) == 0)
1321 continue;
1322 status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
1323 if (status != CL_SUCCESS)
1324 {
1325 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1326 "clGetDeviceIDs failed.", "(%d)", status);
1327 goto cleanup;
1328 }
1329 if (clEnv->device != NULL)
1330 {
1331 clEnv->platform = platforms[i];
1332 goto cleanup;
1333 }
1334 }
1335 }
1336
1337cleanup:
1338 if (platforms!=NULL)
1339 RelinquishMagickMemory(platforms);
1340
1341 OpenCLAvailable = (clEnv->platform!=NULL
1342 && clEnv->device!=NULL)?MagickTrue:MagickFalse;
1343
1344#ifdef MAGICKCORE_CLPERFMARKER
1345 clEndPerfMarkerAMD();
1346#endif
1347
1348 return OpenCLAvailable;
1349}
1350
1351static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) {
1352 if (clEnv->OpenCLInitialized != MagickFalse
1353 && clEnv->platform != NULL
1354 && clEnv->device != NULL) {
1355 clEnv->OpenCLDisabled = MagickFalse;
1356 return MagickTrue;
1357 }
1358 clEnv->OpenCLDisabled = MagickTrue;
1359 return MagickFalse;
1360}
1361
1362
1363static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception);
1364/*
1365%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1366% %
1367% %
1368% %
1369+ I n i t O p e n C L E n v %
1370% %
1371% %
1372% %
1373%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1374%
1375% InitOpenCLEnv() initialize the OpenCL environment
1376%
1377% The format of the RelinquishMagickOpenCLEnv method is:
1378%
1379% MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception)
1380%
1381% A description of each parameter follows:
1382%
1383% o clEnv: OpenCL environment structure
1384%
1385% o exception: return any errors or warnings.
1386%
1387*/
1388
1389static void RelinquishCommandQueues(MagickCLEnv clEnv)
1390{
1391 if (clEnv == (MagickCLEnv) NULL)
1392 return;
1393
1394 LockSemaphoreInfo(clEnv->commandQueuesLock);
1395 while (clEnv->commandQueuesPos >= 0)
1396 clEnv->library->clReleaseCommandQueue(
1397 clEnv->commandQueues[clEnv->commandQueuesPos--]);
1398 UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1399}
1400
1401MagickExport
1402MagickBooleanType InitOpenCLEnvInternal(MagickCLEnv clEnv, ExceptionInfo* exception) {
1403 MagickBooleanType status = MagickTrue;
1404 cl_int clStatus;
1405 cl_context_properties cps[3];
1406
1407#ifdef MAGICKCORE_CLPERFMARKER
1408 {
1409 int status = clInitializePerfMarkerAMD();
1410 if (status == AP_SUCCESS) {
1411 /* printf("PerfMarker successfully initialized\n"); */
1412 }
1413 }
1414#endif
1415 clEnv->OpenCLInitialized = MagickTrue;
1416
1417 /* check and init the global lib */
1418 OpenCLLib=GetOpenCLLib();
1419 if (OpenCLLib)
1420 {
1421 clEnv->library=OpenCLLib;
1422 }
1423 else
1424 {
1425 /* turn off opencl */
1426 MagickBooleanType flag;
1427 flag = MagickTrue;
1428 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
1429 , sizeof(MagickBooleanType), &flag, exception);
1430 }
1431
1432 if (clEnv->OpenCLDisabled != MagickFalse)
1433 goto cleanup;
1434
1435 clEnv->OpenCLDisabled = MagickTrue;
1436 /* setup the OpenCL platform and device */
1437 status = InitOpenCLPlatformDevice(clEnv, exception);
1438 if (status == MagickFalse) {
1439 /* No OpenCL device available */
1440 goto cleanup;
1441 }
1442
1443 /* create an OpenCL context */
1444 cps[0] = CL_CONTEXT_PLATFORM;
1445 cps[1] = (cl_context_properties)clEnv->platform;
1446 cps[2] = 0;
1447 clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
1448 if (clStatus != CL_SUCCESS)
1449 {
1450 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1451 "clCreateContext failed.", "(%d)", clStatus);
1452 status = MagickFalse;
1453 goto cleanup;
1454 }
1455
1456 RelinquishCommandQueues(clEnv);
1457
1458 status = CompileOpenCLKernels(clEnv, exception);
1459 if (status == MagickFalse) {
1460 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1461 "clCreateCommandQueue failed.", "(%d)", status);
1462
1463 goto cleanup;
1464 }
1465
1466 status = EnableOpenCLInternal(clEnv);
1467
1468cleanup:
1469 return status;
1470}
1471
1472
1473MagickExport
1474MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) {
1475 MagickBooleanType status = MagickFalse;
1476
1477 if ((clEnv == NULL) || (getenv("MAGICK_OCL_DEVICE") == (const char *) NULL))
1478 return MagickFalse;
1479
1480#ifdef MAGICKCORE_CLPERFMARKER
1481 clBeginPerfMarkerAMD(__FUNCTION__,"");
1482#endif
1483
1484 LockSemaphoreInfo(clEnv->lock);
1485 if (clEnv->OpenCLInitialized == MagickFalse) {
1486 if (clEnv->device==NULL && clEnv->OpenCLDisabled == MagickFalse)
1487 status = autoSelectDevice(clEnv, exception);
1488 else
1489 status = InitOpenCLEnvInternal(clEnv, exception);
1490 }
1491 UnlockSemaphoreInfo(clEnv->lock);
1492
1493#ifdef MAGICKCORE_CLPERFMARKER
1494 clEndPerfMarkerAMD();
1495#endif
1496 return status;
1497}
1498
1499
1500/*
1501%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1502% %
1503% %
1504% %
1505+ 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 %
1506% %
1507% %
1508% %
1509%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1510%
1511% AcquireOpenCLCommandQueue() acquires an OpenCL command queue
1512%
1513% The format of the AcquireOpenCLCommandQueue method is:
1514%
1515% cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1516%
1517% A description of each parameter follows:
1518%
1519% o clEnv: the OpenCL environment.
1520%
1521*/
1522
1523MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1524{
1525 cl_command_queue
1526 queue;
1527
1528 cl_command_queue_properties
1529 properties;
1530
1531 if (clEnv == (MagickCLEnv) NULL)
1532 return (cl_command_queue) NULL;
1533 LockSemaphoreInfo(clEnv->commandQueuesLock);
1534 if (clEnv->commandQueuesPos >= 0) {
1535 queue=clEnv->commandQueues[clEnv->commandQueuesPos--];
1536 UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1537 }
1538 else {
1539 UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1540 properties=0;
1541#if PROFILE_OCL_KERNELS
1542 properties=CL_QUEUE_PROFILING_ENABLE;
1543#endif
1544 queue=clEnv->library->clCreateCommandQueue(clEnv->context,clEnv->device,
1545 properties,NULL);
1546 }
1547 return(queue);
1548}
1549
1550/*
1551%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1552% %
1553% %
1554% %
1555+ 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 %
1556% %
1557% %
1558% %
1559%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1560%
1561% RelinquishOpenCLCommandQueue() releases the OpenCL command queue
1562%
1563% The format of the RelinquishOpenCLCommandQueue method is:
1564%
1565% MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
1566% cl_command_queue queue)
1567%
1568% A description of each parameter follows:
1569%
1570% o clEnv: the OpenCL environment.
1571%
1572% o queue: the OpenCL queue to be released.
1573%
1574%
1575*/
1576
1577MagickPrivate MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
1578 cl_command_queue queue)
1579{
1580 MagickBooleanType
1581 status;
1582
1583 if (clEnv == NULL)
1584 return(MagickFalse);
1585
1586 LockSemaphoreInfo(clEnv->commandQueuesLock);
1587
1588 if (clEnv->commandQueuesPos >= MAX_COMMAND_QUEUES-1)
1589 {
1590 clEnv->library->clFinish(queue);
1591 status=(clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ?
1592 MagickTrue : MagickFalse;
1593 }
1594 else
1595 {
1596 clEnv->library->clFlush(queue);
1597 clEnv->commandQueues[++clEnv->commandQueuesPos]=queue;
1598 status=MagickTrue;
1599 }
1600
1601 UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1602
1603 return(status);
1604}
1605
1606/*
1607%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1608% %
1609% %
1610% %
1611+ A c q u i r e O p e n C L K e r n e l %
1612% %
1613% %
1614% %
1615%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1616%
1617% AcquireOpenCLKernel() acquires an OpenCL kernel
1618%
1619% The format of the AcquireOpenCLKernel method is:
1620%
1621% cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
1622% MagickOpenCLProgram program, const char* kernelName)
1623%
1624% A description of each parameter follows:
1625%
1626% o clEnv: the OpenCL environment.
1627%
1628% o program: the OpenCL program module that the kernel belongs to.
1629%
1630% o kernelName: the name of the kernel
1631%
1632*/
1633
1634MagickPrivate
1635 cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv, MagickOpenCLProgram program, const char* kernelName)
1636{
1637 cl_int clStatus;
1638 cl_kernel kernel = NULL;
1639 if (clEnv != NULL && kernelName!=NULL)
1640 {
1641 kernel = clEnv->library->clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
1642 }
1643 return kernel;
1644}
1645
1646
1647/*
1648%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1649% %
1650% %
1651% %
1652+ R e l i n q u i s h O p e n C L K e r n e l %
1653% %
1654% %
1655% %
1656%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1657%
1658% RelinquishOpenCLKernel() releases an OpenCL kernel
1659%
1660% The format of the RelinquishOpenCLKernel method is:
1661%
1662% MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv,
1663% cl_kernel kernel)
1664%
1665% A description of each parameter follows:
1666%
1667% o clEnv: the OpenCL environment.
1668%
1669% o kernel: the OpenCL kernel object to be released.
1670%
1671%
1672*/
1673
1674MagickPrivate
1675 MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel)
1676{
1677 MagickBooleanType status = MagickFalse;
1678 if (clEnv != NULL && kernel != NULL)
1679 {
1680 status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
1681 }
1682 return status;
1683}
1684
1685/*
1686%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1687% %
1688% %
1689% %
1690+ 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 %
1691% %
1692% %
1693% %
1694%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1695%
1696% GetOpenCLDeviceLocalMemorySize() returns local memory size of the device
1697%
1698% The format of the GetOpenCLDeviceLocalMemorySize method is:
1699%
1700% unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1701%
1702% A description of each parameter follows:
1703%
1704% o clEnv: the OpenCL environment.
1705%
1706%
1707*/
1708
1709MagickPrivate
1710 unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1711{
1712 cl_ulong localMemorySize;
1713 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
1714 return (unsigned long)localMemorySize;
1715}
1716
1717MagickPrivate
1718 unsigned long GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv clEnv)
1719{
1720 cl_ulong maxMemAllocSize;
1721 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
1722 return (unsigned long)maxMemAllocSize;
1723}
1724
1725
1726/*
1727 Beginning of the OpenCL device selection infrastructure
1728*/
1729
1730
1731typedef enum {
1732 DS_SUCCESS = 0
1733 ,DS_INVALID_PROFILE = 1000
1734 ,DS_MEMORY_ERROR
1735 ,DS_INVALID_PERF_EVALUATOR_TYPE
1736 ,DS_INVALID_PERF_EVALUATOR
1737 ,DS_PERF_EVALUATOR_ERROR
1738 ,DS_FILE_ERROR
1739 ,DS_UNKNOWN_DEVICE_TYPE
1740 ,DS_PROFILE_FILE_ERROR
1741 ,DS_SCORE_SERIALIZER_ERROR
1742 ,DS_SCORE_DESERIALIZER_ERROR
1743} ds_status;
1744
1745/* device type */
1746typedef enum {
1747 DS_DEVICE_NATIVE_CPU = 0
1748 ,DS_DEVICE_OPENCL_DEVICE
1749} ds_device_type;
1750
1751
1752typedef struct {
1753 ds_device_type type;
1754 cl_device_type oclDeviceType;
1755 cl_device_id oclDeviceID;
1756 char* oclDeviceName;
1757 char* oclDriverVersion;
1758 cl_uint oclMaxClockFrequency;
1759 cl_uint oclMaxComputeUnits;
1760 void* score; /* a pointer to the score data, the content/format is application defined */
1761} ds_device;
1762
1763typedef struct {
1764 unsigned int numDevices;
1765 ds_device* devices;
1766 const char* version;
1767} ds_profile;
1768
1769/* deallocate memory used by score */
1770typedef ds_status (*ds_score_release)(void* score);
1771
1772static ds_status releaseDeviceResource(ds_device* device, ds_score_release sr) {
1773 ds_status status = DS_SUCCESS;
1774 if (device) {
1775 if (device->oclDeviceName) RelinquishMagickMemory(device->oclDeviceName);
1776 if (device->oclDriverVersion) RelinquishMagickMemory(device->oclDriverVersion);
1777 if (device->score) status = sr(device->score);
1778 }
1779 return status;
1780}
1781
1782static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) {
1783 ds_status status = DS_SUCCESS;
1784 if (profile!=NULL) {
1785 if (profile->devices!=NULL && sr!=NULL) {
1786 unsigned int i;
1787 for (i = 0; i < profile->numDevices; i++) {
1788 status = releaseDeviceResource(profile->devices+i,sr);
1789 if (status != DS_SUCCESS)
1790 break;
1791 }
1792 RelinquishMagickMemory(profile->devices);
1793 }
1794 RelinquishMagickMemory(profile);
1795 }
1796 return status;
1797}
1798
1799
1800static ds_status initDSProfile(ds_profile** p, const char* version) {
1801 int numDevices = 0;
1802 cl_uint numPlatforms = 0;
1803 cl_platform_id* platforms = NULL;
1804 cl_device_id* devices = NULL;
1805 ds_status status = DS_SUCCESS;
1806 ds_profile* profile = NULL;
1807 unsigned int next = 0;
1808 unsigned int i;
1809
1810 if (p == NULL)
1811 return DS_INVALID_PROFILE;
1812
1813 profile = (ds_profile*) AcquireMagickMemory(sizeof(ds_profile));
1814 if (profile == NULL)
1815 return DS_MEMORY_ERROR;
1816
1817 memset(profile, 0, sizeof(ds_profile));
1818
1819 OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
1820 if (numPlatforms > 0) {
1821 platforms = (cl_platform_id*) AcquireQuantumMemory(numPlatforms,sizeof(cl_platform_id));
1822 if (platforms == NULL) {
1823 status = DS_MEMORY_ERROR;
1824 goto cleanup;
1825 }
1826 OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
1827 for (i = 0; i < (unsigned int)numPlatforms; i++) {
1828 cl_uint num;
1829 if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
1830 numDevices+=num;
1831 }
1832 }
1833
1834 profile->numDevices = numDevices+1; /* +1 to numDevices to include the native CPU */
1835
1836 profile->devices = (ds_device*) AcquireQuantumMemory(profile->numDevices,sizeof(ds_device));
1837 if (profile->devices == NULL) {
1838 profile->numDevices = 0;
1839 status = DS_MEMORY_ERROR;
1840 goto cleanup;
1841 }
1842 memset(profile->devices, 0, profile->numDevices*sizeof(ds_device));
1843
1844 if (numDevices > 0) {
1845 devices = (cl_device_id*) AcquireQuantumMemory(numDevices,sizeof(cl_device_id));
1846 if (devices == NULL) {
1847 status = DS_MEMORY_ERROR;
1848 goto cleanup;
1849 }
1850 for (i = 0; i < (unsigned int)numPlatforms; i++) {
1851 cl_uint num;
1852
1853 int d;
1854 for (d = 0; d < 2; d++) {
1855 unsigned int j;
1856 cl_device_type deviceType;
1857 switch(d) {
1858 case 0:
1859 deviceType = CL_DEVICE_TYPE_GPU;
1860 break;
1861 case 1:
1862 deviceType = CL_DEVICE_TYPE_CPU;
1863 break;
1864 default:
1865 continue;
1866 break;
1867 }
1868 if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
1869 continue;
1870 for (j = 0; j < num; j++, next++) {
1871 size_t length;
1872
1873 profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
1874 profile->devices[next].oclDeviceID = devices[j];
1875
1876 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1877 , 0, NULL, &length);
1878 profile->devices[next].oclDeviceName = (char*) AcquireQuantumMemory(length,sizeof(char));
1879 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1880 , length, profile->devices[next].oclDeviceName, NULL);
1881
1882 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1883 , 0, NULL, &length);
1884 profile->devices[next].oclDriverVersion = (char*) AcquireQuantumMemory(length,sizeof(char));
1885 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1886 , length, profile->devices[next].oclDriverVersion, NULL);
1887
1888 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
1889 , sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);
1890
1891 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
1892 , sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);
1893
1894 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_TYPE
1895 , sizeof(cl_device_type), &profile->devices[next].oclDeviceType, NULL);
1896 }
1897 }
1898 }
1899 }
1900
1901 profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
1902 profile->version = version;
1903
1904cleanup:
1905 if (platforms) RelinquishMagickMemory(platforms);
1906 if (devices) RelinquishMagickMemory(devices);
1907 if (status == DS_SUCCESS) {
1908 *p = profile;
1909 }
1910 else {
1911 if (profile) {
1912 if (profile->devices)
1913 RelinquishMagickMemory(profile->devices);
1914 RelinquishMagickMemory(profile);
1915 }
1916 }
1917 return status;
1918}
1919
1920/* Pointer to a function that calculates the score of a device (ex: device->score)
1921 update the data size of score. The encoding and the format of the score data
1922 is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
1923 */
1924typedef ds_status (*ds_perf_evaluator)(ds_device* device, void* data);
1925
1926typedef enum {
1927 DS_EVALUATE_ALL
1928 ,DS_EVALUATE_NEW_ONLY
1929} ds_evaluation_type;
1930
1931static ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type
1932 ,ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates) {
1933 ds_status status = DS_SUCCESS;
1934 unsigned int i;
1935 unsigned int updates = 0;
1936
1937 if (profile == NULL) {
1938 return DS_INVALID_PROFILE;
1939 }
1940 if (evaluator == NULL) {
1941 return DS_INVALID_PERF_EVALUATOR;
1942 }
1943
1944 for (i = 0; i < profile->numDevices; i++) {
1945 ds_status evaluatorStatus;
1946
1947 switch (type) {
1948 case DS_EVALUATE_NEW_ONLY:
1949 if (profile->devices[i].score != NULL)
1950 break;
1951 /* else fall through */
1952 case DS_EVALUATE_ALL:
1953 evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
1954 if (evaluatorStatus != DS_SUCCESS) {
1955 status = evaluatorStatus;
1956 return status;
1957 }
1958 updates++;
1959 break;
1960 default:
1961 return DS_INVALID_PERF_EVALUATOR_TYPE;
1962 break;
1963 };
1964 }
1965 if (numUpdates)
1966 *numUpdates = updates;
1967 return status;
1968}
1969
1970
1971#define DS_TAG_VERSION "<version>"
1972#define DS_TAG_VERSION_END "</version>"
1973#define DS_TAG_DEVICE "<device>"
1974#define DS_TAG_DEVICE_END "</device>"
1975#define DS_TAG_SCORE "<score>"
1976#define DS_TAG_SCORE_END "</score>"
1977#define DS_TAG_DEVICE_TYPE "<type>"
1978#define DS_TAG_DEVICE_TYPE_END "</type>"
1979#define DS_TAG_DEVICE_NAME "<name>"
1980#define DS_TAG_DEVICE_NAME_END "</name>"
1981#define DS_TAG_DEVICE_DRIVER_VERSION "<driver>"
1982#define DS_TAG_DEVICE_DRIVER_VERSION_END "</driver>"
1983#define DS_TAG_DEVICE_MAX_COMPUTE_UNITS "<max cu>"
1984#define DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END "</max cu>"
1985#define DS_TAG_DEVICE_MAX_CLOCK_FREQ "<max clock>"
1986#define DS_TAG_DEVICE_MAX_CLOCK_FREQ_END "</max clock>"
1987
1988#define DS_DEVICE_NATIVE_CPU_STRING "native_cpu"
1989
1990
1991
1992typedef ds_status (*ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
1993static ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file) {
1994 ds_status status = DS_SUCCESS;
1995 FILE* profileFile = NULL;
1996
1997
1998 if (profile == NULL)
1999 return DS_INVALID_PROFILE;
2000
2001 profileFile = fopen(file, "wb");
2002 if (profileFile==NULL) {
2003 status = DS_FILE_ERROR;
2004 }
2005 else {
2006 unsigned int i;
2007
2008 /* write version string */
2009 fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
2010 fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
2011 fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
2012 fwrite("\n", sizeof(char), 1, profileFile);
2013
2014 for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) {
2015 void* serializedScore;
2016 unsigned int serializedScoreSize;
2017
2018 fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);
2019
2020 fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
2021 fwrite(&profile->devices[i].type,sizeof(ds_device_type),1, profileFile);
2022 fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);
2023
2024 switch(profile->devices[i].type) {
2025 case DS_DEVICE_NATIVE_CPU:
2026 {
2027 /* There's no need to emit a device name for the native CPU device. */
2028 /*
2029 fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
2030 fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
2031 fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
2032 */
2033 }
2034 break;
2035 case DS_DEVICE_OPENCL_DEVICE:
2036 {
2037 char tmp[16];
2038
2039 fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
2040 fwrite(profile->devices[i].oclDeviceName,sizeof(char),strlen(profile->devices[i].oclDeviceName), profileFile);
2041 fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
2042
2043 fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
2044 fwrite(profile->devices[i].oclDriverVersion,sizeof(char),strlen(profile->devices[i].oclDriverVersion), profileFile);
2045 fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);
2046
2047 fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS), profileFile);
2048 (void) FormatLocaleString(tmp,sizeof(tmp),"%d",
2049 profile->devices[i].oclMaxComputeUnits);
2050 fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
2051 fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END), profileFile);
2052
2053 fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ), profileFile);
2054 (void) FormatLocaleString(tmp,sizeof(tmp),"%d",
2055 profile->devices[i].oclMaxClockFrequency);
2056 fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
2057 fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END), profileFile);
2058 }
2059 break;
2060 default:
2061 status = DS_UNKNOWN_DEVICE_TYPE;
2062 break;
2063 };
2064
2065 fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
2066 status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize);
2067 if (status == DS_SUCCESS && serializedScore!=NULL && serializedScoreSize > 0) {
2068 fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
2069 RelinquishMagickMemory(serializedScore);
2070 }
2071 fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
2072 fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
2073 fwrite("\n",sizeof(char),1,profileFile);
2074 }
2075 fclose(profileFile);
2076 }
2077 return status;
2078}
2079
2080
2081static ds_status readProFile(const char* fileName, char** content, size_t* contentSize) {
2082 ds_status status = DS_SUCCESS;
2083 FILE * input = NULL;
2084 size_t size = 0;
2085 size_t rsize = 0;
2086 char* binary = NULL;
2087
2088 *contentSize = 0;
2089 *content = NULL;
2090
2091 input = fopen(fileName, "rb");
2092 if(input == NULL) {
2093 return DS_FILE_ERROR;
2094 }
2095
2096 fseek(input, 0L, SEEK_END);
2097 size = ftell(input);
2098 rewind(input);
2099 binary = (char*) AcquireQuantumMemory(1,size);
2100 if(binary == NULL) {
2101 status = DS_FILE_ERROR;
2102 goto cleanup;
2103 }
2104 rsize = fread(binary, sizeof(char), size, input);
2105 if (rsize!=size
2106 || ferror(input)) {
2107 status = DS_FILE_ERROR;
2108 goto cleanup;
2109 }
2110 *contentSize = size;
2111 *content = binary;
2112
2113cleanup:
2114 if (input != NULL) fclose(input);
2115 if (status != DS_SUCCESS
2116 && binary != NULL) {
2117 RelinquishMagickMemory(binary);
2118 *content = NULL;
2119 *contentSize = 0;
2120 }
2121 return status;
2122}
2123
2124
2125static const char* findString(const char* contentStart, const char* contentEnd, const char* string) {
2126 size_t stringLength;
2127 const char* currentPosition;
2128 const char* found;
2129 found = NULL;
2130 stringLength = strlen(string);
2131 currentPosition = contentStart;
2132 for(currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) {
2133 if (*currentPosition == string[0]) {
2134 if (currentPosition+stringLength < contentEnd) {
2135 if (strncmp(currentPosition, string, stringLength) == 0) {
2136 found = currentPosition;
2137 break;
2138 }
2139 }
2140 }
2141 }
2142 return found;
2143}
2144
2145
2146typedef ds_status (*ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
2147static ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file) {
2148
2149 ds_status status = DS_SUCCESS;
2150 char* contentStart = NULL;
2151 const char* contentEnd = NULL;
2152 size_t contentSize;
2153
2154 if (profile==NULL)
2155 return DS_INVALID_PROFILE;
2156
2157 status = readProFile(file, &contentStart, &contentSize);
2158 if (status == DS_SUCCESS) {
2159 const char* currentPosition;
2160 const char* dataStart;
2161 const char* dataEnd;
2162 size_t versionStringLength;
2163
2164 contentEnd = contentStart + contentSize;
2165 currentPosition = contentStart;
2166
2167
2168 /* parse the version string */
2169 dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
2170 if (dataStart == NULL) {
2171 status = DS_PROFILE_FILE_ERROR;
2172 goto cleanup;
2173 }
2174 dataStart += strlen(DS_TAG_VERSION);
2175
2176 dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
2177 if (dataEnd==NULL) {
2178 status = DS_PROFILE_FILE_ERROR;
2179 goto cleanup;
2180 }
2181
2182 versionStringLength = strlen(profile->version);
2183 if (versionStringLength!=(size_t)(dataEnd-dataStart)
2184 || strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
2185 /* version mismatch */
2186 status = DS_PROFILE_FILE_ERROR;
2187 goto cleanup;
2188 }
2189 currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);
2190
2191 /* parse the device information */
2192DisableMSCWarning(4127)
2193 while (1) {
2194RestoreMSCWarning
2195 unsigned int i;
2196
2197 const char* deviceTypeStart;
2198 const char* deviceTypeEnd;
2199 ds_device_type deviceType;
2200
2201 const char* deviceNameStart;
2202 const char* deviceNameEnd;
2203
2204 const char* deviceScoreStart;
2205 const char* deviceScoreEnd;
2206
2207 const char* deviceDriverStart;
2208 const char* deviceDriverEnd;
2209
2210 const char* tmpStart;
2211 const char* tmpEnd;
2212 char tmp[16];
2213
2214 cl_uint maxClockFrequency;
2215 cl_uint maxComputeUnits;
2216
2217 dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
2218 if (dataStart==NULL) {
2219 /* nothing useful remain, quit...*/
2220 break;
2221 }
2222 dataStart+=strlen(DS_TAG_DEVICE);
2223 dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
2224 if (dataEnd==NULL) {
2225 status = DS_PROFILE_FILE_ERROR;
2226 goto cleanup;
2227 }
2228
2229 /* parse the device type */
2230 deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
2231 if (deviceTypeStart==NULL) {
2232 status = DS_PROFILE_FILE_ERROR;
2233 goto cleanup;
2234 }
2235 deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE);
2236 deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
2237 if (deviceTypeEnd==NULL) {
2238 status = DS_PROFILE_FILE_ERROR;
2239 goto cleanup;
2240 }
2241 memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));
2242
2243
2244 /* parse the device name */
2245 if (deviceType == DS_DEVICE_OPENCL_DEVICE) {
2246
2247 deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
2248 if (deviceNameStart==NULL) {
2249 status = DS_PROFILE_FILE_ERROR;
2250 goto cleanup;
2251 }
2252 deviceNameStart+=strlen(DS_TAG_DEVICE_NAME);
2253 deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
2254 if (deviceNameEnd==NULL) {
2255 status = DS_PROFILE_FILE_ERROR;
2256 goto cleanup;
2257 }
2258
2259
2260 deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
2261 if (deviceDriverStart==NULL) {
2262 status = DS_PROFILE_FILE_ERROR;
2263 goto cleanup;
2264 }
2265 deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION);
2266 deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
2267 if (deviceDriverEnd ==NULL) {
2268 status = DS_PROFILE_FILE_ERROR;
2269 goto cleanup;
2270 }
2271
2272
2273 tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
2274 if (tmpStart==NULL) {
2275 status = DS_PROFILE_FILE_ERROR;
2276 goto cleanup;
2277 }
2278 tmpStart+=strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
2279 tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END);
2280 if (tmpEnd ==NULL) {
2281 status = DS_PROFILE_FILE_ERROR;
2282 goto cleanup;
2283 }
2284 memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2285 tmp[tmpEnd-tmpStart] = '\0';
2286 maxComputeUnits = strtol(tmp,(char **) NULL,10);
2287
2288
2289 tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2290 if (tmpStart==NULL) {
2291 status = DS_PROFILE_FILE_ERROR;
2292 goto cleanup;
2293 }
2294 tmpStart+=strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2295 tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ_END);
2296 if (tmpEnd ==NULL) {
2297 status = DS_PROFILE_FILE_ERROR;
2298 goto cleanup;
2299 }
2300 memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2301 tmp[tmpEnd-tmpStart] = '\0';
2302 maxClockFrequency = strtol(tmp,(char **) NULL,10);
2303
2304
2305 /* check if this device is on the system */
2306 for (i = 0; i < profile->numDevices; i++) {
2307 if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) {
2308 size_t actualDeviceNameLength;
2309 size_t driverVersionLength;
2310
2311 actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
2312 driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
2313 if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
2314 && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
2315 && maxComputeUnits == profile->devices[i].oclMaxComputeUnits
2316 && maxClockFrequency == profile->devices[i].oclMaxClockFrequency
2317 && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
2318 && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==(int)0) {
2319
2320 deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2321 if (deviceNameStart==NULL) {
2322 status = DS_PROFILE_FILE_ERROR;
2323 goto cleanup;
2324 }
2325 deviceScoreStart+=strlen(DS_TAG_SCORE);
2326 deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2327 status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2328 if (status != DS_SUCCESS) {
2329 goto cleanup;
2330 }
2331 }
2332 }
2333 }
2334
2335 }
2336 else if (deviceType == DS_DEVICE_NATIVE_CPU) {
2337 for (i = 0; i < profile->numDevices; i++) {
2338 if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) {
2339 deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2340 if (deviceScoreStart==NULL) {
2341 status = DS_PROFILE_FILE_ERROR;
2342 goto cleanup;
2343 }
2344 deviceScoreStart+=strlen(DS_TAG_SCORE);
2345 deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2346 status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2347 if (status != DS_SUCCESS) {
2348 goto cleanup;
2349 }
2350 }
2351 }
2352 }
2353
2354 /* skip over the current one to find the next device */
2355 currentPosition = dataEnd+strlen(DS_TAG_DEVICE_END);
2356 }
2357 }
2358cleanup:
2359 if (contentStart!=NULL) RelinquishMagickMemory(contentStart);
2360 return status;
2361}
2362
2363
2364#if 0
2365static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
2366 unsigned int i;
2367 if (profile == NULL || num==NULL)
2368 return DS_MEMORY_ERROR;
2369 *num=0;
2370 for (i = 0; i < profile->numDevices; i++) {
2371 if (profile->devices[i].score == NULL) {
2372 (*num)++;
2373 }
2374 }
2375 return DS_SUCCESS;
2376}
2377#endif
2378
2379/*
2380 End of the OpenCL device selection infrastructure
2381*/
2382
2383
2384typedef double AccelerateScoreType;
2385
2386static ds_status AcceleratePerfEvaluator(ds_device *device,
2387 void *magick_unused(data))
2388{
2389#define ACCELERATE_PERF_DIMEN "2048x1536"
2390#define NUM_ITER 2
2391#define ReturnStatus(status) \
2392{ \
2393 if (oldClEnv != (MagickCLEnv) NULL) \
2394 defaultCLEnv=oldClEnv; \
2395 if (clEnv != (MagickCLEnv) NULL) \
2396 (void) RelinquishMagickOpenCLEnv(clEnv); \
2397 return status; \
2398}
2399
2400 AccelerateTimer
2401 timer;
2402
2403 ExceptionInfo
2404 *exception=NULL;
2405
2406 MagickBooleanType
2407 status;
2408
2409 MagickCLEnv
2410 clEnv=NULL,
2411 oldClEnv=NULL;
2412
2413 magick_unreferenced(data);
2414
2415 if (device == NULL)
2416 ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2417
2418 clEnv=AcquireMagickOpenCLEnv();
2419 exception=AcquireExceptionInfo();
2420
2421 if (device->type == DS_DEVICE_NATIVE_CPU)
2422 {
2423 /* CPU device */
2424 MagickBooleanType flag=MagickTrue;
2425 SetMagickOpenCLEnvParamInternal(clEnv,
2426 MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
2427 &flag,exception);
2428 }
2429 else if (device->type == DS_DEVICE_OPENCL_DEVICE)
2430 {
2431 /* OpenCL device */
2432 SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
2433 sizeof(cl_device_id),&device->oclDeviceID,exception);
2434 }
2435 else
2436 ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2437
2438 /* recompile the OpenCL kernels if it needs to */
2439 clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;
2440
2441 status=InitOpenCLEnvInternal(clEnv,exception);
2442 oldClEnv=defaultCLEnv;
2443 defaultCLEnv=clEnv;
2444
2445 /* microbenchmark */
2446 if (status != MagickFalse)
2447 {
2448 Image
2449 *inputImage;
2450
2451 ImageInfo
2452 *imageInfo;
2453
2454 int
2455 i;
2456
2457 imageInfo=AcquireImageInfo();
2458 CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
2459 CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
2460 inputImage=ReadImage(imageInfo,exception);
2461 if (inputImage == (Image *) NULL)
2462 ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2463
2464 initAccelerateTimer(&timer);
2465
2466 for (i=0; i<=NUM_ITER; i++)
2467 {
2468 cl_uint
2469 event_count;
2470
2471 cl_event
2472 *events;
2473
2474 Image
2475 *bluredImage,
2476 *resizedImage,
2477 *unsharpedImage;
2478
2479 if (i > 0)
2480 startAccelerateTimer(&timer);
2481
2482#ifdef MAGICKCORE_CLPERFMARKER
2483 clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
2484#endif
2485
2486 bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
2487 unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
2488 exception);
2489 resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,1.0,
2490 exception);
2491
2492 /*
2493 We need this to get a proper performance benchmark, the operations
2494 are executed asynchronous.
2495 */
2496 if (device->type != DS_DEVICE_NATIVE_CPU)
2497 {
2498 events=GetOpenCLEvents(resizedImage,&event_count);
2499 if (event_count > 0)
2500 clEnv->library->clWaitForEvents(event_count,events);
2501 events=(cl_event *) RelinquishMagickMemory(events);
2502 }
2503
2504#ifdef MAGICKCORE_CLPERFMARKER
2505 clEndPerfMarkerAMD();
2506#endif
2507
2508 if (i > 0)
2509 stopAccelerateTimer(&timer);
2510
2511 if (bluredImage)
2512 DestroyImage(bluredImage);
2513 if (unsharpedImage)
2514 DestroyImage(unsharpedImage);
2515 if (resizedImage)
2516 DestroyImage(resizedImage);
2517 }
2518 DestroyImage(inputImage);
2519 }
2520 /* end of microbenchmark */
2521
2522 if (device->score == NULL)
2523 device->score= AcquireMagickMemory(sizeof(AccelerateScoreType));
2524
2525 if (status != MagickFalse)
2526 *(AccelerateScoreType*) device->score=readAccelerateTimer(&timer);
2527 else
2528 *(AccelerateScoreType*) device->score=42;
2529
2530 ReturnStatus(DS_SUCCESS);
2531}
2532
2533ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
2534 if (device
2535 && device->score) {
2536 /* generate a string from the score */
2537 char* s = (char*) AcquireQuantumMemory(256,sizeof(char));
2538 (void) FormatLocaleString(s,256,"%.4f",*((AccelerateScoreType*)
2539 device->score));
2540 *serializedScore = (void*)s;
2541 *serializedScoreSize = (unsigned int) strlen(s);
2542 return DS_SUCCESS;
2543 }
2544 else {
2545 return DS_SCORE_SERIALIZER_ERROR;
2546 }
2547}
2548
2549ds_status AccelerateScoreDeserializer(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize) {
2550 if (device) {
2551 /* convert the string back to an int */
2552 char* s = (char*) AcquireQuantumMemory(1,serializedScoreSize+1);
2553 memcpy(s, serializedScore, serializedScoreSize);
2554 s[serializedScoreSize] = (char)'\0';
2555 device->score = AcquireMagickMemory(sizeof(AccelerateScoreType));
2556 *((AccelerateScoreType*)device->score) = (AccelerateScoreType)
2557 strtod(s, (char **) NULL);
2558 RelinquishMagickMemory(s);
2559 return DS_SUCCESS;
2560 }
2561 else {
2562 return DS_SCORE_DESERIALIZER_ERROR;
2563 }
2564}
2565
2566ds_status AccelerateScoreRelease(void* score) {
2567 if (score!=NULL) {
2568 RelinquishMagickMemory(score);
2569 }
2570 return DS_SUCCESS;
2571}
2572
2573ds_status canWriteProfileToFile(const char *path)
2574{
2575 FILE* profileFile = fopen(path, "ab");
2576
2577 if (profileFile==NULL)
2578 return DS_FILE_ERROR;
2579
2580 fclose(profileFile);
2581 return DS_SUCCESS;
2582}
2583
2584
2585#define IMAGEMAGICK_PROFILE_VERSION "ImageMagick Device Selection v0.9"
2586#define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile"
2587static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
2588
2589 MagickBooleanType mStatus = MagickFalse;
2590 ds_status status;
2591 ds_profile* profile;
2592 unsigned int numDeviceProfiled = 0;
2593 unsigned int i;
2594 unsigned int bestDeviceIndex;
2595 AccelerateScoreType bestScore;
2596 char path[MaxTextExtent];
2597 MagickBooleanType flag;
2598 ds_evaluation_type profileType;
2599
2600 LockDefaultOpenCLEnv();
2601
2602 /* Initially, just set OpenCL to off */
2603 flag = MagickTrue;
2604 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2605 , sizeof(MagickBooleanType), &flag, exception);
2606
2607 /* check and init the global lib */
2608 OpenCLLib=GetOpenCLLib();
2609 if (OpenCLLib==NULL)
2610 {
2611 mStatus=InitOpenCLEnvInternal(clEnv, exception);
2612 goto cleanup;
2613 }
2614
2615 clEnv->library=OpenCLLib;
2616
2617 status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
2618 if (status!=DS_SUCCESS) {
2619 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2620 goto cleanup;
2621 }
2622
2623 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2624 ,GetOpenCLCachedFilesDirectory()
2625 ,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
2626
2627 if (canWriteProfileToFile(path) != DS_SUCCESS) {
2628 /* We can not write out a device profile, so don't run the benchmark */
2629 /* select the first GPU device */
2630
2631 bestDeviceIndex = 0;
2632 for (i = 1; i < profile->numDevices; i++) {
2633 if ((profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) && (profile->devices[i].oclDeviceType == CL_DEVICE_TYPE_GPU)) {
2634 bestDeviceIndex = i;
2635 break;
2636 }
2637 }
2638 }
2639 else {
2640 if (clEnv->regenerateProfile != MagickFalse) {
2641 profileType = DS_EVALUATE_ALL;
2642 }
2643 else {
2644 readProfileFromFile(profile, AccelerateScoreDeserializer, path);
2645 profileType = DS_EVALUATE_NEW_ONLY;
2646 }
2647 status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
2648
2649 if (status!=DS_SUCCESS) {
2650 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2651 goto cleanup;
2652 }
2653 if (numDeviceProfiled > 0) {
2654 status = writeProfileToFile(profile, AccelerateScoreSerializer, path);
2655 if (status!=DS_SUCCESS) {
2656 (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when saving the profile into a file", "'%s'", ".");
2657 }
2658 }
2659
2660 /* pick the best device */
2661 bestDeviceIndex = 0;
2662 bestScore = *(AccelerateScoreType*)profile->devices[bestDeviceIndex].score;
2663 for (i = 1; i < profile->numDevices; i++) {
2664 AccelerateScoreType score = *(AccelerateScoreType*)profile->devices[i].score;
2665 if (score < bestScore) {
2666 bestDeviceIndex = i;
2667 bestScore = score;
2668 }
2669 }
2670 }
2671
2672 /* set up clEnv with the best device */
2673 if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
2674 /* CPU device */
2675 flag = MagickTrue;
2676 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2677 , sizeof(MagickBooleanType), &flag, exception);
2678 }
2679 else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
2680 /* OpenCL device */
2681 flag = MagickFalse;
2682 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2683 , sizeof(MagickBooleanType), &flag, exception);
2684 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2685 , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
2686 }
2687 else {
2688 status = DS_PERF_EVALUATOR_ERROR;
2689 goto cleanup;
2690 }
2691 mStatus=InitOpenCLEnvInternal(clEnv, exception);
2692
2693 status = releaseDSProfile(profile, AccelerateScoreRelease);
2694 if (status!=DS_SUCCESS) {
2695 (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
2696 }
2697
2698cleanup:
2699
2700 UnlockDefaultOpenCLEnv();
2701 return mStatus;
2702}
2703
2704
2705/*
2706%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2707% %
2708% %
2709% %
2710+ I n i t I m a g e M a g i c k O p e n C L %
2711% %
2712% %
2713% %
2714%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2715%
2716% InitImageMagickOpenCL() provides a simplified interface to initialize
2717% the OpenCL environtment in ImageMagick
2718%
2719% The format of the InitImageMagickOpenCL() method is:
2720%
2721% MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
2722% void* userSelectedDevice,
2723% void* selectedDevice)
2724%
2725% A description of each parameter follows:
2726%
2727% o mode: OpenCL mode in ImageMagick, could be off,auto,user
2728%
2729% o userSelectedDevice: when in user mode, a pointer to the selected
2730% cl_device_id
2731%
2732% o selectedDevice: a pointer to cl_device_id where the selected
2733% cl_device_id by ImageMagick could be returned
2734%
2735% o exception: exception
2736%
2737*/
2738MagickExport MagickBooleanType InitImageMagickOpenCL(
2739 ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
2740 ExceptionInfo *exception)
2741{
2742 MagickBooleanType status = MagickFalse;
2743 MagickCLEnv clEnv = NULL;
2744 MagickBooleanType flag;
2745
2746 clEnv = GetDefaultOpenCLEnv();
2747 if (clEnv!=NULL) {
2748 switch(mode) {
2749
2750 case MAGICK_OPENCL_OFF:
2751 flag = MagickTrue;
2752 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2753 , sizeof(MagickBooleanType), &flag, exception);
2754 status = InitOpenCLEnv(clEnv, exception);
2755
2756 if (selectedDevice)
2757 *(cl_device_id*)selectedDevice = NULL;
2758 break;
2759
2760 case MAGICK_OPENCL_DEVICE_SELECT_USER:
2761
2762 if (userSelectedDevice == NULL)
2763 return MagickFalse;
2764
2765 flag = MagickFalse;
2766 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2767 , sizeof(MagickBooleanType), &flag, exception);
2768
2769 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2770 , sizeof(cl_device_id), userSelectedDevice,exception);
2771
2772 status = InitOpenCLEnv(clEnv, exception);
2773 if (selectedDevice) {
2774 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2775 , sizeof(cl_device_id), selectedDevice, exception);
2776 }
2777 break;
2778
2779 case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
2780 flag = MagickTrue;
2781 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
2782 , sizeof(MagickBooleanType), &flag, exception);
2783 flag = MagickTrue;
2784 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
2785 , sizeof(MagickBooleanType), &flag, exception);
2786
2787 /* fall through here!! */
2788 case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
2789 default:
2790 {
2791 cl_device_id d = NULL;
2792 flag = MagickFalse;
2793 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2794 , sizeof(MagickBooleanType), &flag, exception);
2795 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2796 , sizeof(cl_device_id), &d,exception);
2797 status = InitOpenCLEnv(clEnv, exception);
2798 if (selectedDevice) {
2799 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2800 , sizeof(cl_device_id), selectedDevice, exception);
2801 }
2802 }
2803 break;
2804 };
2805 }
2806 return status;
2807}
2808
2809
2810MagickPrivate
2811MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2812 const char *module,const char *function,const size_t line,
2813 const ExceptionType severity,const char *tag,const char *format,...) {
2814 MagickBooleanType
2815 status;
2816
2817 MagickCLEnv clEnv;
2818
2819 status = MagickTrue;
2820
2821 clEnv = GetDefaultOpenCLEnv();
2822
2823 assert(exception != (ExceptionInfo *) NULL);
2824 assert(exception->signature == MagickCoreSignature);
2825
2826 if (severity!=0) {
2827 cl_device_type dType;
2828 clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
2829 if (dType == CL_DEVICE_TYPE_CPU) {
2830 char buffer[MaxTextExtent];
2831 clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);
2832
2833 /* Workaround for Intel OpenCL CPU runtime bug */
2834 /* Turn off OpenCL when a problem is detected! */
2835 if (strncmp(buffer, "Intel",5) == 0) {
2836
2837 InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
2838 }
2839 }
2840 }
2841
2842#ifdef OPENCLLOG_ENABLED
2843 {
2844 va_list
2845 operands;
2846 va_start(operands,format);
2847 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands);
2848 va_end(operands);
2849 }
2850#else
2851 magick_unreferenced(module);
2852 magick_unreferenced(function);
2853 magick_unreferenced(line);
2854 magick_unreferenced(tag);
2855 magick_unreferenced(format);
2856#endif
2857
2858 return(status);
2859}
2860
2861char* openclCachedFilesDirectory;
2862SemaphoreInfo* openclCachedFilesDirectoryLock;
2863
2864MagickPrivate
2865const char* GetOpenCLCachedFilesDirectory() {
2866 if (openclCachedFilesDirectory == NULL) {
2867 if (openclCachedFilesDirectoryLock == NULL)
2868 {
2869 ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
2870 }
2871 LockSemaphoreInfo(openclCachedFilesDirectoryLock);
2872 if (openclCachedFilesDirectory == NULL) {
2873 char path[MaxTextExtent];
2874 char *home = NULL;
2875 char *temp = NULL;
2876 struct stat attributes;
2877 MagickBooleanType status;
2878 int mkdirStatus = 0;
2879
2880
2881
2882 home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
2883 if (home == (char *) NULL)
2884 {
2885 home=GetEnvironmentValue("XDG_CACHE_HOME");
2886#if defined(MAGICKCORE_WINDOWS_SUPPORT) || defined(__MINGW32__)
2887 if (home == (char *) NULL)
2888 home=GetEnvironmentValue("LOCALAPPDATA");
2889 if (home == (char *) NULL)
2890 home=GetEnvironmentValue("APPDATA");
2891 if (home == (char *) NULL)
2892 home=GetEnvironmentValue("USERPROFILE");
2893#endif
2894 }
2895
2896 if (home != (char *) NULL)
2897 {
2898 /* first check if $HOME exists */
2899 (void) FormatLocaleString(path,MaxTextExtent,"%s",home);
2900 status=GetPathAttributes(path,&attributes);
2901 if (status == MagickFalse)
2902 {
2903
2904#ifdef MAGICKCORE_WINDOWS_SUPPORT
2905 mkdirStatus = mkdir(path);
2906#else
2907 mkdirStatus = mkdir(path, 0777);
2908#endif
2909 }
2910
2911 /* first check if $HOME/ImageMagick exists */
2912 if (mkdirStatus==0)
2913 {
2914 (void) FormatLocaleString(path,MaxTextExtent,
2915 "%s%sImageMagick",home,DirectorySeparator);
2916
2917 status=GetPathAttributes(path,&attributes);
2918 if (status == MagickFalse)
2919 {
2920#ifdef MAGICKCORE_WINDOWS_SUPPORT
2921 mkdirStatus = mkdir(path);
2922#else
2923 mkdirStatus = mkdir(path, 0777);
2924#endif
2925 }
2926 }
2927
2928 if (mkdirStatus==0)
2929 {
2930 temp = (char*)AcquireCriticalMemory(strlen(path)+1);
2931 CopyMagickString(temp,path,strlen(path)+1);
2932 }
2933 home=DestroyString(home);
2934 } else {
2935 home=GetEnvironmentValue("HOME");
2936 if (home != (char *) NULL)
2937 {
2938 /*
2939 */
2940
2941 /* first check if $HOME/.cache exists */
2942 (void) FormatLocaleString(path,MaxTextExtent,"%s%s.cache",
2943 home,DirectorySeparator);
2944 status=GetPathAttributes(path,&attributes);
2945 if (status == MagickFalse)
2946 {
2947
2948#ifdef MAGICKCORE_WINDOWS_SUPPORT
2949 mkdirStatus = mkdir(path);
2950#else
2951 mkdirStatus = mkdir(path, 0777);
2952#endif
2953 }
2954
2955 /* first check if $HOME/.cache/ImageMagick exists */
2956 if (mkdirStatus==0)
2957 {
2958 (void) FormatLocaleString(path,MaxTextExtent,
2959 "%s%s.cache%sImageMagick",home,DirectorySeparator,
2960 DirectorySeparator);
2961
2962 status=GetPathAttributes(path,&attributes);
2963 if (status == MagickFalse)
2964 {
2965#ifdef MAGICKCORE_WINDOWS_SUPPORT
2966 mkdirStatus = mkdir(path);
2967#else
2968 mkdirStatus = mkdir(path, 0777);
2969#endif
2970 }
2971 }
2972
2973 if (mkdirStatus==0)
2974 {
2975 temp = (char*)AcquireCriticalMemory(strlen(path)+1);
2976 CopyMagickString(temp,path,strlen(path)+1);
2977 }
2978 home=DestroyString(home);
2979 }
2980 }
2981 openclCachedFilesDirectory = temp;
2982 }
2983 UnlockSemaphoreInfo(openclCachedFilesDirectoryLock);
2984 }
2985 return openclCachedFilesDirectory;
2986}
2987
2988/* create a function for OpenCL log */
2989MagickPrivate
2990void OpenCLLog(const char* message) {
2991
2992#ifdef OPENCLLOG_ENABLED
2993#define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
2994
2995 FILE* log;
2996 if (getenv("MAGICK_OCL_LOG"))
2997 {
2998 if (message) {
2999 char path[MaxTextExtent];
3000 unsigned long allocSize;
3001
3002 MagickCLEnv clEnv;
3003
3004 clEnv = GetDefaultOpenCLEnv();
3005
3006 /* dump the source into a file */
3007 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
3008 ,GetOpenCLCachedFilesDirectory()
3009 ,DirectorySeparator,OPENCL_LOG_FILE);
3010
3011
3012 log = fopen(path, "ab");
3013 if (log == (FILE *) NULL)
3014 return;
3015 fwrite(message, sizeof(char), strlen(message), log);
3016 fwrite("\n", sizeof(char), 1, log);
3017
3018 if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
3019 {
3020 allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
3021 fprintf(log, "Devic Max Memory Alloc Size: %lu\n", allocSize);
3022 }
3023
3024 fclose(log);
3025 }
3026 }
3027#else
3028 magick_unreferenced(message);
3029#endif
3030}
3031
3032MagickPrivate void OpenCLTerminus()
3033{
3034 DumpProfileData();
3035 if (openclCachedFilesDirectory != (char *) NULL)
3036 openclCachedFilesDirectory=DestroyString(openclCachedFilesDirectory);
3037 if (openclCachedFilesDirectoryLock != (SemaphoreInfo*)NULL)
3038 DestroySemaphoreInfo(&openclCachedFilesDirectoryLock);
3039 if (defaultCLEnv != (MagickCLEnv) NULL)
3040 {
3041 (void) RelinquishMagickOpenCLEnv(defaultCLEnv);
3042 defaultCLEnv=(MagickCLEnv)NULL;
3043 }
3044 if (defaultCLEnvLock != (SemaphoreInfo*) NULL)
3045 DestroySemaphoreInfo(&defaultCLEnvLock);
3046 if (OpenCLLib != (MagickLibrary *)NULL)
3047 {
3048 if (OpenCLLib->base != (void *) NULL)
3049 (void) lt_dlclose(OpenCLLib->base);
3050 OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
3051 }
3052 if (OpenCLLibLock != (SemaphoreInfo*)NULL)
3053 DestroySemaphoreInfo(&OpenCLLibLock);
3054}
3055
3056#else
3057
3058struct _MagickCLEnv {
3059 MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */
3060};
3061
3062/*
3063* Return the OpenCL environment
3064*/
3065MagickExport MagickCLEnv GetDefaultOpenCLEnv()
3066{
3067 return (MagickCLEnv) NULL;
3068}
3069
3070MagickExport MagickBooleanType SetMagickOpenCLEnvParam(
3071 MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
3072 size_t magick_unused(dataSize),void *magick_unused(data),
3073 ExceptionInfo *magick_unused(exception))
3074{
3075 magick_unreferenced(clEnv);
3076 magick_unreferenced(param);
3077 magick_unreferenced(dataSize);
3078 magick_unreferenced(data);
3079 magick_unreferenced(exception);
3080 return(MagickFalse);
3081}
3082
3083MagickExport MagickBooleanType GetMagickOpenCLEnvParam(
3084 MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
3085 size_t magick_unused(dataSize),void *magick_unused(data),
3086 ExceptionInfo *magick_unused(exception))
3087{
3088 magick_unreferenced(clEnv);
3089 magick_unreferenced(param);
3090 magick_unreferenced(dataSize);
3091 magick_unreferenced(data);
3092 magick_unreferenced(exception);
3093 return(MagickFalse);
3094}
3095
3096MagickExport MagickBooleanType InitOpenCLEnv(MagickCLEnv magick_unused(clEnv),
3097 ExceptionInfo *magick_unused(exception))
3098{
3099 magick_unreferenced(clEnv);
3100 magick_unreferenced(exception);
3101 return(MagickFalse);
3102}
3103
3104MagickExport MagickBooleanType InitImageMagickOpenCL(
3105 ImageMagickOpenCLMode magick_unused(mode),
3106 void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
3107 ExceptionInfo *magick_unused(exception))
3108{
3109 magick_unreferenced(mode);
3110 magick_unreferenced(userSelectedDevice);
3111 magick_unreferenced(selectedDevice);
3112 magick_unreferenced(exception);
3113 return(MagickFalse);
3114}
3115
3116#endif /* MAGICKCORE_OPENCL_SUPPORT */
_ExceptionInfo
Definition exception.h:103
_ImageInfo
Definition image.h:346
_Image
Definition image.h:134
_MagickCLEnv
Definition opencl.c:3058