2 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
6 % OOO PPPP EEEEE N N CCCC L %
8 % O O PPPP EEE N N N C L %
10 % OOO P EEEEE N N CCCC LLLLL %
13 % MagickCore OpenCL Methods %
20 % Copyright 1999-2014 ImageMagick Studio LLC, a non-profit organization %
21 % dedicated to making software imaging solutions freely available. %
23 % You may not use this file except in compliance with the License. You may %
24 % obtain a copy of the License at %
26 % http://www.imagemagick.org/script/license.php %
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. %
34 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
43 #include "MagickCore/studio.h"
44 #include "MagickCore/artifact.h"
45 #include "MagickCore/cache.h"
46 #include "MagickCore/color.h"
47 #include "MagickCore/compare.h"
48 #include "MagickCore/constitute.h"
49 #include "MagickCore/distort.h"
50 #include "MagickCore/draw.h"
51 #include "MagickCore/effect.h"
52 #include "MagickCore/exception.h"
53 #include "MagickCore/exception-private.h"
54 #include "MagickCore/fx.h"
55 #include "MagickCore/gem.h"
56 #include "MagickCore/geometry.h"
57 #include "MagickCore/image.h"
58 #include "MagickCore/image-private.h"
59 #include "MagickCore/layer.h"
60 #include "MagickCore/mime-private.h"
61 #include "MagickCore/memory_.h"
62 #include "MagickCore/monitor.h"
63 #include "MagickCore/montage.h"
64 #include "MagickCore/morphology.h"
65 #include "MagickCore/nt-base.h"
66 #include "MagickCore/nt-base-private.h"
67 #include "MagickCore/opencl.h"
68 #include "MagickCore/opencl-private.h"
69 #include "MagickCore/option.h"
70 #include "MagickCore/policy.h"
71 #include "MagickCore/property.h"
72 #include "MagickCore/quantize.h"
73 #include "MagickCore/quantum.h"
74 #include "MagickCore/random_.h"
75 #include "MagickCore/random-private.h"
76 #include "MagickCore/resample.h"
77 #include "MagickCore/resource_.h"
78 #include "MagickCore/splay-tree.h"
79 #include "MagickCore/semaphore.h"
80 #include "MagickCore/statistic.h"
81 #include "MagickCore/string_.h"
82 #include "MagickCore/token.h"
83 #include "MagickCore/utility.h"
85 #ifdef MAGICKCORE_CLPERFMARKER
86 #include "CLPerfMarker.h"
90 #if defined(MAGICKCORE_OPENCL_SUPPORT)
92 #ifdef MAGICKCORE_HAVE_OPENCL_CL_H
93 #define MAGICKCORE_OPENCL_MACOSX 1
97 #define NUM_CL_RAND_GENERATORS 1024 /* number of random number generators running in parallel */
101 * Dynamic library loading functions
104 #ifdef MAGICKCORE_WINDOWS_SUPPORT
109 // dynamically load a library. returns NULL on failure
110 void *OsLibraryLoad(const char *libraryName)
112 #ifdef MAGICKCORE_WINDOWS_SUPPORT
113 return (void *)LoadLibraryA(libraryName);
115 return (void *)dlopen(libraryName, RTLD_NOW);
119 // get a function pointer from a loaded library. returns NULL on failure.
120 void *OsLibraryGetFunctionAddress(void *library, const char *functionName)
122 #ifdef MAGICKCORE_WINDOWS_SUPPORT
123 if (!library || !functionName)
127 return (void *) GetProcAddress( (HMODULE)library, functionName);
129 if (!library || !functionName)
133 return (void *)dlsym(library, functionName);
138 void OsLibraryUnload(void *library)
140 #ifdef MAGICKCORE_WINDOWS_SUPPORT
141 FreeLibrary( (HMODULE)library);
149 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
153 + A c q u i r e M a g i c k O p e n C L E n v %
157 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
159 % AcquireMagickOpenCLEnv() allocates the MagickCLEnv structure
163 MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
166 clEnv = (MagickCLEnv) AcquireMagickMemory(sizeof(struct _MagickCLEnv));
169 memset(clEnv, 0, sizeof(struct _MagickCLEnv));
170 ActivateSemaphoreInfo(&clEnv->lock);
177 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
181 + 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 %
185 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
187 % RelinquishMagickOpenCLEnv() destroy the MagickCLEnv structure
189 % The format of the RelinquishMagickOpenCLEnv method is:
191 % MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv)
193 % A description of each parameter follows:
195 % o clEnv: MagickCLEnv structure to destroy
199 MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv)
201 if (clEnv != (MagickCLEnv)NULL)
203 DestroySemaphoreInfo(&clEnv->lock);
204 RelinquishMagickMemory(clEnv);
212 * Default OpenCL environment
214 MagickCLEnv defaultCLEnv;
215 SemaphoreInfo* defaultCLEnvLock;
220 MagickLibrary * OpenCLLib;
221 SemaphoreInfo* OpenCLLibLock;
224 static MagickBooleanType bindOpenCLFunctions(void* library)
226 #ifdef MAGICKCORE_OPENCL_MACOSX
227 #define BIND(X) OpenCLLib->X= &X;
230 if ((OpenCLLib->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(library,#X)) == NULL)\
234 BIND(clGetPlatformIDs);
235 BIND(clGetPlatformInfo);
237 BIND(clGetDeviceIDs);
238 BIND(clGetDeviceInfo);
240 BIND(clCreateContext);
242 BIND(clCreateBuffer);
243 BIND(clReleaseMemObject);
245 BIND(clCreateProgramWithSource);
246 BIND(clCreateProgramWithBinary);
247 BIND(clBuildProgram);
248 BIND(clGetProgramInfo);
249 BIND(clGetProgramBuildInfo);
251 BIND(clCreateKernel);
252 BIND(clReleaseKernel);
253 BIND(clSetKernelArg);
258 BIND(clEnqueueNDRangeKernel);
259 BIND(clEnqueueReadBuffer);
260 BIND(clEnqueueMapBuffer);
261 BIND(clEnqueueUnmapMemObject);
263 BIND(clCreateCommandQueue);
264 BIND(clReleaseCommandQueue);
269 MagickLibrary * GetOpenCLLib()
271 if (OpenCLLib == NULL)
273 if (OpenCLLibLock == NULL)
275 ActivateSemaphoreInfo(&OpenCLLibLock);
278 LockSemaphoreInfo(OpenCLLibLock);
280 OpenCLLib = (MagickLibrary *) AcquireMagickMemory (sizeof (MagickLibrary));
282 if (OpenCLLib != NULL)
284 MagickBooleanType status = MagickFalse;
285 void * library = NULL;
287 #ifdef MAGICKCORE_OPENCL_MACOSX
288 status = bindOpenCLFunctions(library);
291 memset(OpenCLLib, 0, sizeof(MagickLibrary));
292 #ifdef MAGICKCORE_WINDOWS_SUPPORT
293 library = OsLibraryLoad("OpenCL.dll");
295 library = OsLibraryLoad("libOpenCL.so");
298 status = bindOpenCLFunctions(library);
300 if (status==MagickTrue)
301 OpenCLLib->base=library;
303 OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
307 UnlockSemaphoreInfo(OpenCLLibLock);
316 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
320 + G e t D e f a u l t O p e n C L E n v %
324 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
326 % GetDefaultOpenCLEnv() returns the default OpenCL env
328 % The format of the GetDefaultOpenCLEnv method is:
330 % MagickCLEnv GetDefaultOpenCLEnv()
332 % A description of each parameter follows:
334 % o exception: return any errors or warnings.
338 MagickExport MagickCLEnv GetDefaultOpenCLEnv()
340 if (defaultCLEnv == NULL)
342 if (defaultCLEnvLock == NULL)
344 ActivateSemaphoreInfo(&defaultCLEnvLock);
346 LockSemaphoreInfo(defaultCLEnvLock);
347 defaultCLEnv = AcquireMagickOpenCLEnv();
348 UnlockSemaphoreInfo(defaultCLEnvLock);
353 static void LockDefaultOpenCLEnv() {
354 if (defaultCLEnvLock == NULL)
356 ActivateSemaphoreInfo(&defaultCLEnvLock);
358 LockSemaphoreInfo(defaultCLEnvLock);
361 static void UnlockDefaultOpenCLEnv() {
362 if (defaultCLEnvLock == NULL)
364 ActivateSemaphoreInfo(&defaultCLEnvLock);
367 UnlockSemaphoreInfo(defaultCLEnvLock);
372 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
376 + S e t D e f a u l t O p e n C L E n v %
380 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
382 % SetDefaultOpenCLEnv() sets the new OpenCL environment as default
383 % and returns the old OpenCL environment
385 % The format of the SetDefaultOpenCLEnv() method is:
387 % MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv)
389 % A description of each parameter follows:
391 % o clEnv: the new default OpenCL environment.
394 MagickExport MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv)
397 LockDefaultOpenCLEnv();
398 oldEnv = defaultCLEnv;
399 defaultCLEnv = clEnv;
400 UnlockDefaultOpenCLEnv();
407 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
411 + S e t M a g i c k O p e n C L E n v P a r a m %
415 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
417 % SetMagickOpenCLEnvParam() sets the parameters in the OpenCL environment
419 % The format of the SetMagickOpenCLEnvParam() method is:
421 % MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv,
422 % MagickOpenCLEnvParam param, size_t dataSize, void* data,
423 % ExceptionInfo* exception)
425 % A description of each parameter follows:
427 % o clEnv: the OpenCL environment.
429 % o param: the parameter to be set.
431 % o dataSize: the data size of the parameter value.
433 % o data: the pointer to the new parameter value
435 % o exception: return any errors or warnings
439 static MagickBooleanType SetMagickOpenCLEnvParamInternal(MagickCLEnv clEnv, MagickOpenCLEnvParam param
440 , size_t dataSize, void* data, ExceptionInfo* exception)
442 MagickBooleanType status = MagickFalse;
450 case MAGICK_OPENCL_ENV_PARAM_DEVICE:
451 if (dataSize != sizeof(clEnv->device))
453 clEnv->device = *((cl_device_id*)data);
454 clEnv->OpenCLInitialized = MagickFalse;
458 case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED:
459 if (dataSize != sizeof(clEnv->OpenCLDisabled))
461 clEnv->OpenCLDisabled = *((MagickBooleanType*)data);
462 clEnv->OpenCLInitialized = MagickFalse;
466 case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED:
467 (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "SetMagickOpenCLEnvParm cannot modify the OpenCL initialization state.", "'%s'", ".");
470 case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
471 if (dataSize != sizeof(clEnv->disableProgramCache))
473 clEnv->disableProgramCache = *((MagickBooleanType*)data);
474 clEnv->OpenCLInitialized = MagickFalse;
478 case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
479 if (dataSize != sizeof(clEnv->regenerateProfile))
481 clEnv->regenerateProfile = *((MagickBooleanType*)data);
482 clEnv->OpenCLInitialized = MagickFalse;
495 MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
496 , size_t dataSize, void* data, ExceptionInfo* exception) {
497 MagickBooleanType status = MagickFalse;
499 LockSemaphoreInfo(clEnv->lock);
500 status = SetMagickOpenCLEnvParamInternal(clEnv,param,dataSize,data,exception);
501 UnlockSemaphoreInfo(clEnv->lock);
507 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
511 + G e t M a g i c k O p e n C L E n v P a r a m %
515 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
517 % GetMagickOpenCLEnvParam() gets the parameters in the OpenCL environment
519 % The format of the GetMagickOpenCLEnvParam() method is:
521 % MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv,
522 % MagickOpenCLEnvParam param, size_t dataSize, void* data,
523 % ExceptionInfo* exception)
525 % A description of each parameter follows:
527 % o clEnv: the OpenCL environment.
529 % o param: the parameter to be returned.
531 % o dataSize: the data size of the parameter value.
533 % o data: the location where the returned parameter value will be stored
535 % o exception: return any errors or warnings
540 MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
541 , size_t dataSize, void* data, ExceptionInfo* exception)
546 magick_unreferenced(exception);
548 status = MagickFalse;
556 case MAGICK_OPENCL_ENV_PARAM_DEVICE:
557 if (dataSize != sizeof(cl_device_id))
559 *((cl_device_id*)data) = clEnv->device;
563 case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED:
564 if (dataSize != sizeof(clEnv->OpenCLDisabled))
566 *((MagickBooleanType*)data) = clEnv->OpenCLDisabled;
570 case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED:
571 if (dataSize != sizeof(clEnv->OpenCLDisabled))
573 *((MagickBooleanType*)data) = clEnv->OpenCLInitialized;
577 case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
578 if (dataSize != sizeof(clEnv->disableProgramCache))
580 *((MagickBooleanType*)data) = clEnv->disableProgramCache;
584 case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
585 if (dataSize != sizeof(clEnv->regenerateProfile))
587 *((MagickBooleanType*)data) = clEnv->regenerateProfile;
601 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
605 + G e t O p e n C L C o n t e x t %
609 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
611 % GetOpenCLContext() returns the OpenCL context
613 % The format of the GetOpenCLContext() method is:
615 % cl_context GetOpenCLContext(MagickCLEnv clEnv)
617 % A description of each parameter follows:
619 % o clEnv: OpenCL environment
624 cl_context GetOpenCLContext(MagickCLEnv clEnv) {
628 return clEnv->context;
631 static char* getBinaryCLProgramName(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
635 char path[MaxTextExtent];
636 char deviceName[MaxTextExtent];
637 const char* prefix = "magick_opencl";
638 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MaxTextExtent, deviceName, NULL);
640 /* strip out illegal characters for file names */
643 if ( *ptr == ' ' || *ptr == '\\' || *ptr == '/' || *ptr == ':' || *ptr == '*'
644 || *ptr == '?' || *ptr == '"' || *ptr == '<' || *ptr == '>' || *ptr == '|')
650 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s_%s_%02d_%08x_%.20g.bin",
651 GetOpenCLCachedFilesDirectory(),DirectorySeparator,prefix,deviceName,
652 (unsigned int) prog,signature,(double) sizeof(char*)*8);
653 name = (char*)AcquireMagickMemory(strlen(path)+1);
654 CopyMagickString(name,path,strlen(path)+1);
658 static MagickBooleanType saveBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature, ExceptionInfo* exception)
660 MagickBooleanType saveSuccessful;
662 size_t binaryProgramSize;
663 unsigned char* binaryProgram;
664 char* binaryFileName;
667 #ifdef MAGICKCORE_CLPERFMARKER
668 clBeginPerfMarkerAMD(__FUNCTION__,"");
671 binaryProgram = NULL;
672 binaryFileName = NULL;
674 saveSuccessful = MagickFalse;
676 clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binaryProgramSize, NULL);
677 if (clStatus != CL_SUCCESS)
679 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
683 binaryProgram = (unsigned char*) AcquireMagickMemory(binaryProgramSize);
684 clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARIES, sizeof(char*), &binaryProgram, NULL);
685 if (clStatus != CL_SUCCESS)
687 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
691 binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
692 fileHandle = fopen(binaryFileName, "wb");
693 if (fileHandle != NULL)
695 fwrite(binaryProgram, sizeof(char), binaryProgramSize, fileHandle);
696 saveSuccessful = MagickTrue;
700 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
701 "Saving binary kernel failed.", "'%s'", ".");
705 if (fileHandle != NULL)
707 if (binaryProgram != NULL)
708 RelinquishMagickMemory(binaryProgram);
709 if (binaryFileName != NULL)
710 free(binaryFileName);
712 #ifdef MAGICKCORE_CLPERFMARKER
713 clEndPerfMarkerAMD();
716 return saveSuccessful;
719 static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
721 MagickBooleanType loadSuccessful;
722 unsigned char* binaryProgram;
723 char* binaryFileName;
726 #ifdef MAGICKCORE_CLPERFMARKER
727 clBeginPerfMarkerAMD(__FUNCTION__,"");
730 binaryProgram = NULL;
731 binaryFileName = NULL;
733 loadSuccessful = MagickFalse;
735 binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
736 fileHandle = fopen(binaryFileName, "rb");
737 if (fileHandle != NULL)
742 cl_int clBinaryStatus;
746 b_error |= fseek( fileHandle, 0, SEEK_END ) < 0;
747 b_error |= ( length = ftell( fileHandle ) ) <= 0;
748 b_error |= fseek( fileHandle, 0, SEEK_SET ) < 0;
752 binaryProgram = (unsigned char*)AcquireMagickMemory(length);
753 if (binaryProgram == NULL)
756 memset(binaryProgram, 0, length);
757 b_error |= fread(binaryProgram, 1, length, fileHandle) != length;
759 clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
760 if (clStatus != CL_SUCCESS
761 || clBinaryStatus != CL_SUCCESS)
764 loadSuccessful = MagickTrue;
768 if (fileHandle != NULL)
770 if (binaryFileName != NULL)
771 free(binaryFileName);
772 if (binaryProgram != NULL)
773 RelinquishMagickMemory(binaryProgram);
775 #ifdef MAGICKCORE_CLPERFMARKER
776 clEndPerfMarkerAMD();
779 return loadSuccessful;
782 static unsigned int stringSignature(const char* string)
784 unsigned int stringLength;
786 unsigned int signature;
790 const unsigned int* u;
793 #ifdef MAGICKCORE_CLPERFMARKER
794 clBeginPerfMarkerAMD(__FUNCTION__,"");
797 stringLength = strlen(string);
798 signature = stringLength;
799 n = stringLength/sizeof(unsigned int);
801 for (i = 0; i < n; i++)
805 if (n * sizeof(unsigned int) != stringLength)
808 j = n * sizeof(unsigned int);
809 for (i = 0; i < 4; i++,j++)
811 if (j < stringLength)
820 #ifdef MAGICKCORE_CLPERFMARKER
821 clEndPerfMarkerAMD();
827 /* OpenCL kernels for accelerate.c */
828 extern const char *accelerateKernels, *accelerateKernels2;
830 static MagickBooleanType CompileOpenCLKernels(MagickCLEnv clEnv, ExceptionInfo* exception)
832 MagickBooleanType status = MagickFalse;
835 char* accelerateKernelsBuffer = NULL;
837 /* The index of the program strings in this array has to match the value of the enum MagickOpenCLProgram */
838 const char* MagickOpenCLProgramStrings[MAGICK_OPENCL_NUM_PROGRAMS];
840 char options[MaxTextExtent];
841 unsigned int optionsSignature;
843 #ifdef MAGICKCORE_CLPERFMARKER
844 clBeginPerfMarkerAMD(__FUNCTION__,"");
847 /* Get additional options */
848 (void) FormatLocaleString(options, MaxTextExtent, CLOptions, (float)QuantumRange,
849 (float)QuantumScale, (float)CLCharQuantumScale, (float)MagickEpsilon, (float)MagickPI, (unsigned int)MaxMap, (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
852 if (getenv("MAGICK_OCL_DEF"))
855 strcat(options,getenv("MAGICK_OCL_DEF"));
860 if (getenv("MAGICK_OCL_BUILD"))
861 printf("options: %s\n", options);
864 optionsSignature = stringSignature(options);
866 /* get all the OpenCL program strings here */
867 accelerateKernelsBuffer = (char*) AcquireMagickMemory(strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
868 sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
869 MagickOpenCLProgramStrings[MAGICK_OPENCL_ACCELERATE] = accelerateKernelsBuffer;
871 for (i = 0; i < MAGICK_OPENCL_NUM_PROGRAMS; i++)
873 MagickBooleanType loadSuccessful = MagickFalse;
874 unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature;
876 /* try to load the binary first */
877 if (clEnv->disableProgramCache != MagickTrue
878 && !getenv("MAGICK_OCL_REC"))
879 loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature);
881 if (loadSuccessful == MagickFalse)
883 /* Binary CL program unavailable, compile the program from source */
884 size_t programLength = strlen(MagickOpenCLProgramStrings[i]);
885 clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
886 if (clStatus!=CL_SUCCESS)
888 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
889 "clCreateProgramWithSource failed.", "(%d)", (int)clStatus);
895 clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
896 if (clStatus!=CL_SUCCESS)
898 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
899 "clBuildProgram failed.", "(%d)", (int)clStatus);
901 if (loadSuccessful == MagickFalse)
903 char path[MaxTextExtent];
906 /* dump the source into a file */
907 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
908 ,GetOpenCLCachedFilesDirectory()
909 ,DirectorySeparator,"magick_badcl.cl");
910 fileHandle = fopen(path, "wb");
911 if (fileHandle != NULL)
913 fwrite(MagickOpenCLProgramStrings[i], sizeof(char), strlen(MagickOpenCLProgramStrings[i]), fileHandle);
917 /* dump the build log */
921 clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
922 log = (char*)AcquireMagickMemory(logSize);
923 clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);
925 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
926 ,GetOpenCLCachedFilesDirectory()
927 ,DirectorySeparator,"magick_badcl_build.log");
928 fileHandle = fopen(path, "wb");
929 if (fileHandle != NULL)
931 const char* buildOptionsTitle = "build options: ";
932 fwrite(buildOptionsTitle, sizeof(char), strlen(buildOptionsTitle), fileHandle);
933 fwrite(options, sizeof(char), strlen(options), fileHandle);
934 fwrite("\n",sizeof(char), 1, fileHandle);
935 fwrite(log, sizeof(char), logSize, fileHandle);
938 RelinquishMagickMemory(log);
944 if (loadSuccessful == MagickFalse)
946 /* Save the binary to a file to avoid re-compilation of the kernels in the future */
947 saveBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception);
955 if (accelerateKernelsBuffer!=NULL) RelinquishMagickMemory(accelerateKernelsBuffer);
957 #ifdef MAGICKCORE_CLPERFMARKER
958 clEndPerfMarkerAMD();
964 static MagickBooleanType InitOpenCLPlatformDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
967 cl_uint numPlatforms = 0;
968 cl_platform_id *platforms = NULL;
969 char* MAGICK_OCL_DEVICE = NULL;
970 MagickBooleanType OpenCLAvailable = MagickFalse;
972 #ifdef MAGICKCORE_CLPERFMARKER
973 clBeginPerfMarkerAMD(__FUNCTION__,"");
976 /* check if there's an environment variable overriding the device selection */
977 MAGICK_OCL_DEVICE = getenv("MAGICK_OCL_DEVICE");
978 if (MAGICK_OCL_DEVICE != NULL)
980 if (strcmp(MAGICK_OCL_DEVICE, "CPU") == 0)
982 clEnv->deviceType = CL_DEVICE_TYPE_CPU;
984 else if (strcmp(MAGICK_OCL_DEVICE, "GPU") == 0)
986 clEnv->deviceType = CL_DEVICE_TYPE_GPU;
988 else if (strcmp(MAGICK_OCL_DEVICE, "OFF") == 0)
990 /* OpenCL disabled */
994 else if (clEnv->deviceType == 0) {
995 clEnv->deviceType = CL_DEVICE_TYPE_ALL;
998 if (clEnv->device != NULL)
1000 status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
1001 if (status != CL_SUCCESS) {
1002 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1003 "Failed to get OpenCL platform from the selected device.", "(%d)", status);
1007 else if (clEnv->platform != NULL)
1010 platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id));
1011 if (platforms == (cl_platform_id *) NULL)
1013 (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1014 "AcquireMagickMemory failed.",".");
1017 platforms[0] = clEnv->platform;
1021 clEnv->device = NULL;
1023 /* Get the number of OpenCL platforms available */
1024 status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms);
1025 if (status != CL_SUCCESS)
1027 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1028 "clGetplatformIDs failed.", "(%d)", status);
1032 /* No OpenCL available, just leave */
1033 if (numPlatforms == 0) {
1037 platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id));
1038 if (platforms == (cl_platform_id *) NULL)
1040 (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1041 "AcquireMagickMemory failed.",".");
1045 status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL);
1046 if (status != CL_SUCCESS)
1048 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1049 "clGetPlatformIDs failed.", "(%d)", status);
1054 /* Device selection */
1055 clEnv->device = NULL;
1056 for (j = 0; j < 2; j++)
1059 cl_device_type deviceType;
1060 if (clEnv->deviceType == CL_DEVICE_TYPE_ALL)
1063 deviceType = CL_DEVICE_TYPE_GPU;
1065 deviceType = CL_DEVICE_TYPE_CPU;
1072 deviceType = clEnv->deviceType;
1074 for (i = 0; i < numPlatforms; i++)
1077 status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
1078 if (status != CL_SUCCESS)
1080 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1081 "clGetDeviceIDs failed.", "(%d)", status);
1084 if (clEnv->device != NULL)
1086 clEnv->platform = platforms[i];
1093 if (platforms!=NULL)
1094 RelinquishMagickMemory(platforms);
1096 OpenCLAvailable = (clEnv->platform!=NULL
1097 && clEnv->device!=NULL)?MagickTrue:MagickFalse;
1099 #ifdef MAGICKCORE_CLPERFMARKER
1100 clEndPerfMarkerAMD();
1103 return OpenCLAvailable;
1106 static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) {
1107 if (clEnv->OpenCLInitialized != MagickFalse
1108 && clEnv->platform != NULL
1109 && clEnv->device != NULL) {
1110 clEnv->OpenCLDisabled = MagickFalse;
1113 clEnv->OpenCLDisabled = MagickTrue;
1118 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception);
1120 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1124 + I n i t O p e n C L E n v %
1128 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1130 % InitOpenCLEnv() initialize the OpenCL environment
1132 % The format of the RelinquishMagickOpenCLEnv method is:
1134 % MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception)
1136 % A description of each parameter follows:
1138 % o clEnv: OpenCL environment structure
1140 % o exception: return any errors or warnings.
1145 MagickBooleanType InitOpenCLEnvInternal(MagickCLEnv clEnv, ExceptionInfo* exception) {
1146 MagickBooleanType status = MagickTrue;
1148 cl_context_properties cps[3];
1150 #ifdef MAGICKCORE_CLPERFMARKER
1152 int status = clInitializePerfMarkerAMD();
1153 if (status == AP_SUCCESS) {
1154 //printf("PerfMarker successfully initialized\n");
1158 clEnv->OpenCLInitialized = MagickTrue;
1160 /* check and init the global lib */
1161 OpenCLLib=GetOpenCLLib();
1164 clEnv->library=OpenCLLib;
1168 /* turn off opencl */
1169 MagickBooleanType flag;
1171 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
1172 , sizeof(MagickBooleanType), &flag, exception);
1175 if (clEnv->OpenCLDisabled != MagickFalse)
1178 clEnv->OpenCLDisabled = MagickTrue;
1179 /* setup the OpenCL platform and device */
1180 status = InitOpenCLPlatformDevice(clEnv, exception);
1181 if (status == MagickFalse) {
1182 /* No OpenCL device available */
1186 /* create an OpenCL context */
1187 cps[0] = CL_CONTEXT_PLATFORM;
1188 cps[1] = (cl_context_properties)clEnv->platform;
1190 clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
1191 if (clStatus != CL_SUCCESS)
1193 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1194 "clCreateContext failed.", "(%d)", clStatus);
1195 status = MagickFalse;
1199 status = CompileOpenCLKernels(clEnv, exception);
1200 if (status == MagickFalse) {
1201 (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1202 "clCreateCommandQueue failed.", "(%d)", status);
1204 status = MagickFalse;
1208 status = EnableOpenCLInternal(clEnv);
1216 MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) {
1217 MagickBooleanType status = MagickFalse;
1222 #ifdef MAGICKCORE_CLPERFMARKER
1223 clBeginPerfMarkerAMD(__FUNCTION__,"");
1226 LockSemaphoreInfo(clEnv->lock);
1227 if (clEnv->OpenCLInitialized == MagickFalse) {
1228 if (clEnv->device==NULL
1229 && clEnv->OpenCLDisabled == MagickFalse)
1230 status = autoSelectDevice(clEnv, exception);
1232 status = InitOpenCLEnvInternal(clEnv, exception);
1234 UnlockSemaphoreInfo(clEnv->lock);
1236 #ifdef MAGICKCORE_CLPERFMARKER
1237 clEndPerfMarkerAMD();
1244 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1248 + 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 %
1252 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1254 % AcquireOpenCLCommandQueue() acquires an OpenCL command queue
1256 % The format of the AcquireOpenCLCommandQueue method is:
1258 % cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1260 % A description of each parameter follows:
1262 % o clEnv: the OpenCL environment.
1267 cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1270 return clEnv->library->clCreateCommandQueue(clEnv->context, clEnv->device, 0, NULL);
1277 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1281 + 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 %
1285 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1287 % RelinquishOpenCLCommandQueue() releases the OpenCL command queue
1289 % The format of the RelinquishOpenCLCommandQueue method is:
1291 % MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
1292 % cl_command_queue queue)
1294 % A description of each parameter follows:
1296 % o clEnv: the OpenCL environment.
1298 % o queue: the OpenCL queue to be released.
1303 MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv, cl_command_queue queue)
1307 return ((clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ? MagickTrue:MagickFalse);
1316 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1320 + A c q u i r e O p e n C L K e r n e l %
1324 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1326 % AcquireOpenCLKernel() acquires an OpenCL kernel
1328 % The format of the AcquireOpenCLKernel method is:
1330 % cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
1331 % MagickOpenCLProgram program, const char* kernelName)
1333 % A description of each parameter follows:
1335 % o clEnv: the OpenCL environment.
1337 % o program: the OpenCL program module that the kernel belongs to.
1339 % o kernelName: the name of the kernel
1344 cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv, MagickOpenCLProgram program, const char* kernelName)
1347 cl_kernel kernel = NULL;
1348 if (clEnv != NULL && kernelName!=NULL)
1350 kernel = clEnv->library->clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
1357 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1361 + R e l i n q u i s h O p e n C L K e r n e l %
1365 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1367 % RelinquishOpenCLKernel() releases an OpenCL kernel
1369 % The format of the RelinquishOpenCLKernel method is:
1371 % MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv,
1374 % A description of each parameter follows:
1376 % o clEnv: the OpenCL environment.
1378 % o kernel: the OpenCL kernel object to be released.
1384 MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel)
1386 MagickBooleanType status = MagickFalse;
1387 if (clEnv != NULL && kernel != NULL)
1389 status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
1395 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1399 + 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 %
1403 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1405 % GetOpenCLDeviceLocalMemorySize() returns local memory size of the device
1407 % The format of the GetOpenCLDeviceLocalMemorySize method is:
1409 % unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1411 % A description of each parameter follows:
1413 % o clEnv: the OpenCL environment.
1419 unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1421 cl_ulong localMemorySize;
1422 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
1423 return (unsigned long)localMemorySize;
1427 unsigned long GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv clEnv)
1429 cl_ulong maxMemAllocSize;
1430 clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
1431 return (unsigned long)maxMemAllocSize;
1436 Beginning of the OpenCL device selection infrastructure
1442 ,DS_INVALID_PROFILE = 1000
1444 ,DS_INVALID_PERF_EVALUATOR_TYPE
1445 ,DS_INVALID_PERF_EVALUATOR
1446 ,DS_PERF_EVALUATOR_ERROR
1448 ,DS_UNKNOWN_DEVICE_TYPE
1449 ,DS_PROFILE_FILE_ERROR
1450 ,DS_SCORE_SERIALIZER_ERROR
1451 ,DS_SCORE_DESERIALIZER_ERROR
1456 DS_DEVICE_NATIVE_CPU = 0
1457 ,DS_DEVICE_OPENCL_DEVICE
1462 ds_device_type type;
1463 cl_device_id oclDeviceID;
1464 char* oclDeviceName;
1465 char* oclDriverVersion;
1466 cl_uint oclMaxClockFrequency;
1467 cl_uint oclMaxComputeUnits;
1468 void* score; /* a pointer to the score data, the content/format is application defined */
1472 unsigned int numDevices;
1474 const char* version;
1477 /* deallocate memory used by score */
1478 typedef ds_status (*ds_score_release)(void* score);
1480 static ds_status releaseDeviceResource(ds_device* device, ds_score_release sr) {
1481 ds_status status = DS_SUCCESS;
1483 if (device->oclDeviceName) free(device->oclDeviceName);
1484 if (device->oclDriverVersion) free(device->oclDriverVersion);
1485 if (device->score) status = sr(device->score);
1490 static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) {
1491 ds_status status = DS_SUCCESS;
1492 if (profile!=NULL) {
1493 if (profile->devices!=NULL && sr!=NULL) {
1495 for (i = 0; i < profile->numDevices; i++) {
1496 status = releaseDeviceResource(profile->devices+i,sr);
1497 if (status != DS_SUCCESS)
1500 free(profile->devices);
1508 static ds_status initDSProfile(ds_profile** p, const char* version) {
1510 cl_uint numPlatforms = 0;
1511 cl_platform_id* platforms = NULL;
1512 cl_device_id* devices = NULL;
1513 ds_status status = DS_SUCCESS;
1514 ds_profile* profile = NULL;
1515 unsigned int next = 0;
1519 return DS_INVALID_PROFILE;
1521 profile = (ds_profile*)malloc(sizeof(ds_profile));
1522 if (profile == NULL)
1523 return DS_MEMORY_ERROR;
1525 memset(profile, 0, sizeof(ds_profile));
1527 OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
1528 if (numPlatforms > 0) {
1529 platforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id));
1530 if (platforms == NULL) {
1531 status = DS_MEMORY_ERROR;
1534 OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
1535 for (i = 0; i < (unsigned int)numPlatforms; i++) {
1537 if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
1542 profile->numDevices = numDevices+1; /* +1 to numDevices to include the native CPU */
1544 profile->devices = (ds_device*)malloc(profile->numDevices*sizeof(ds_device));
1545 if (profile->devices == NULL) {
1546 profile->numDevices = 0;
1547 status = DS_MEMORY_ERROR;
1550 memset(profile->devices, 0, profile->numDevices*sizeof(ds_device));
1552 if (numDevices > 0) {
1553 devices = (cl_device_id*)malloc(numDevices*sizeof(cl_device_id));
1554 if (devices == NULL) {
1555 status = DS_MEMORY_ERROR;
1558 for (i = 0; i < (unsigned int)numPlatforms; i++) {
1562 for (d = 0; d < 2; d++) {
1564 cl_device_type deviceType;
1567 deviceType = CL_DEVICE_TYPE_GPU;
1570 deviceType = CL_DEVICE_TYPE_CPU;
1576 if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
1578 for (j = 0; j < num; j++, next++) {
1581 profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
1582 profile->devices[next].oclDeviceID = devices[j];
1584 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1585 , 0, NULL, &length);
1586 profile->devices[next].oclDeviceName = (char*)malloc(sizeof(char)*length);
1587 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1588 , length, profile->devices[next].oclDeviceName, NULL);
1590 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1591 , 0, NULL, &length);
1592 profile->devices[next].oclDriverVersion = (char*)malloc(sizeof(char)*length);
1593 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1594 , length, profile->devices[next].oclDriverVersion, NULL);
1596 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
1597 , sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);
1599 OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
1600 , sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);
1606 profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
1607 profile->version = version;
1610 if (platforms) free(platforms);
1611 if (devices) free(devices);
1612 if (status == DS_SUCCESS) {
1617 if (profile->devices)
1618 free(profile->devices);
1625 /* Pointer to a function that calculates the score of a device (ex: device->score)
1626 update the data size of score. The encoding and the format of the score data
1627 is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
1629 typedef ds_status (*ds_perf_evaluator)(ds_device* device, void* data);
1633 ,DS_EVALUATE_NEW_ONLY
1634 } ds_evaluation_type;
1636 static ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type
1637 ,ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates) {
1638 ds_status status = DS_SUCCESS;
1640 unsigned int updates = 0;
1642 if (profile == NULL) {
1643 return DS_INVALID_PROFILE;
1645 if (evaluator == NULL) {
1646 return DS_INVALID_PERF_EVALUATOR;
1649 for (i = 0; i < profile->numDevices; i++) {
1650 ds_status evaluatorStatus;
1653 case DS_EVALUATE_NEW_ONLY:
1654 if (profile->devices[i].score != NULL)
1656 /* else fall through */
1657 case DS_EVALUATE_ALL:
1658 evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
1659 if (evaluatorStatus != DS_SUCCESS) {
1660 status = evaluatorStatus;
1666 return DS_INVALID_PERF_EVALUATOR_TYPE;
1671 *numUpdates = updates;
1676 #define DS_TAG_VERSION "<version>"
1677 #define DS_TAG_VERSION_END "</version>"
1678 #define DS_TAG_DEVICE "<device>"
1679 #define DS_TAG_DEVICE_END "</device>"
1680 #define DS_TAG_SCORE "<score>"
1681 #define DS_TAG_SCORE_END "</score>"
1682 #define DS_TAG_DEVICE_TYPE "<type>"
1683 #define DS_TAG_DEVICE_TYPE_END "</type>"
1684 #define DS_TAG_DEVICE_NAME "<name>"
1685 #define DS_TAG_DEVICE_NAME_END "</name>"
1686 #define DS_TAG_DEVICE_DRIVER_VERSION "<driver>"
1687 #define DS_TAG_DEVICE_DRIVER_VERSION_END "</driver>"
1688 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS "<max cu>"
1689 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END "</max cu>"
1690 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ "<max clock>"
1691 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ_END "</max clock>"
1693 #define DS_DEVICE_NATIVE_CPU_STRING "native_cpu"
1697 typedef ds_status (*ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
1698 static ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file) {
1699 ds_status status = DS_SUCCESS;
1700 FILE* profileFile = NULL;
1703 if (profile == NULL)
1704 return DS_INVALID_PROFILE;
1706 profileFile = fopen(file, "wb");
1707 if (profileFile==NULL) {
1708 status = DS_FILE_ERROR;
1713 /* write version string */
1714 fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
1715 fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
1716 fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
1717 fwrite("\n", sizeof(char), 1, profileFile);
1719 for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) {
1720 void* serializedScore;
1721 unsigned int serializedScoreSize;
1723 fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);
1725 fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
1726 fwrite(&profile->devices[i].type,sizeof(ds_device_type),1, profileFile);
1727 fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);
1729 switch(profile->devices[i].type) {
1730 case DS_DEVICE_NATIVE_CPU:
1732 /* There's no need to emit a device name for the native CPU device. */
1734 fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
1735 fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
1736 fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
1740 case DS_DEVICE_OPENCL_DEVICE:
1744 fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
1745 fwrite(profile->devices[i].oclDeviceName,sizeof(char),strlen(profile->devices[i].oclDeviceName), profileFile);
1746 fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
1748 fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
1749 fwrite(profile->devices[i].oclDriverVersion,sizeof(char),strlen(profile->devices[i].oclDriverVersion), profileFile);
1750 fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);
1752 fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS), profileFile);
1753 sprintf(tmp,"%d",profile->devices[i].oclMaxComputeUnits);
1754 fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
1755 fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END), profileFile);
1757 fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ), profileFile);
1758 sprintf(tmp,"%d",profile->devices[i].oclMaxClockFrequency);
1759 fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
1760 fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END), profileFile);
1764 status = DS_UNKNOWN_DEVICE_TYPE;
1768 fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
1769 status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize);
1770 if (status == DS_SUCCESS && serializedScore!=NULL && serializedScoreSize > 0) {
1771 fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
1772 free(serializedScore);
1774 fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
1775 fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
1776 fwrite("\n",sizeof(char),1,profileFile);
1778 fclose(profileFile);
1784 static ds_status readProFile(const char* fileName, char** content, size_t* contentSize) {
1785 ds_status status = DS_SUCCESS;
1786 FILE * input = NULL;
1789 char* binary = NULL;
1794 input = fopen(fileName, "rb");
1796 return DS_FILE_ERROR;
1799 fseek(input, 0L, SEEK_END);
1800 size = ftell(input);
1802 binary = (char*)malloc(size);
1803 if(binary == NULL) {
1804 status = DS_FILE_ERROR;
1807 rsize = fread(binary, sizeof(char), size, input);
1810 status = DS_FILE_ERROR;
1813 *contentSize = size;
1817 if (input != NULL) fclose(input);
1818 if (status != DS_SUCCESS
1819 && binary != NULL) {
1828 static const char* findString(const char* contentStart, const char* contentEnd, const char* string) {
1829 size_t stringLength;
1830 const char* currentPosition;
1833 stringLength = strlen(string);
1834 currentPosition = contentStart;
1835 for(currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) {
1836 if (*currentPosition == string[0]) {
1837 if (currentPosition+stringLength < contentEnd) {
1838 if (strncmp(currentPosition, string, stringLength) == 0) {
1839 found = currentPosition;
1849 typedef ds_status (*ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
1850 static ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file) {
1852 ds_status status = DS_SUCCESS;
1853 char* contentStart = NULL;
1854 const char* contentEnd = NULL;
1858 return DS_INVALID_PROFILE;
1860 status = readProFile(file, &contentStart, &contentSize);
1861 if (status == DS_SUCCESS) {
1862 const char* currentPosition;
1863 const char* dataStart;
1864 const char* dataEnd;
1865 size_t versionStringLength;
1867 contentEnd = contentStart + contentSize;
1868 currentPosition = contentStart;
1871 /* parse the version string */
1872 dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
1873 if (dataStart == NULL) {
1874 status = DS_PROFILE_FILE_ERROR;
1877 dataStart += strlen(DS_TAG_VERSION);
1879 dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
1880 if (dataEnd==NULL) {
1881 status = DS_PROFILE_FILE_ERROR;
1885 versionStringLength = strlen(profile->version);
1886 if (versionStringLength!=(size_t)(dataEnd-dataStart)
1887 || strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
1888 /* version mismatch */
1889 status = DS_PROFILE_FILE_ERROR;
1892 currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);
1894 /* parse the device information */
1895 DisableMSCWarning(4127)
1900 const char* deviceTypeStart;
1901 const char* deviceTypeEnd;
1902 ds_device_type deviceType;
1904 const char* deviceNameStart;
1905 const char* deviceNameEnd;
1907 const char* deviceScoreStart;
1908 const char* deviceScoreEnd;
1910 const char* deviceDriverStart;
1911 const char* deviceDriverEnd;
1913 const char* tmpStart;
1917 cl_uint maxClockFrequency;
1918 cl_uint maxComputeUnits;
1920 dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
1921 if (dataStart==NULL) {
1922 /* nothing useful remain, quit...*/
1925 dataStart+=strlen(DS_TAG_DEVICE);
1926 dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
1927 if (dataEnd==NULL) {
1928 status = DS_PROFILE_FILE_ERROR;
1932 /* parse the device type */
1933 deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
1934 if (deviceTypeStart==NULL) {
1935 status = DS_PROFILE_FILE_ERROR;
1938 deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE);
1939 deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
1940 if (deviceTypeEnd==NULL) {
1941 status = DS_PROFILE_FILE_ERROR;
1944 memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));
1947 /* parse the device name */
1948 if (deviceType == DS_DEVICE_OPENCL_DEVICE) {
1950 deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
1951 if (deviceNameStart==NULL) {
1952 status = DS_PROFILE_FILE_ERROR;
1955 deviceNameStart+=strlen(DS_TAG_DEVICE_NAME);
1956 deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
1957 if (deviceNameEnd==NULL) {
1958 status = DS_PROFILE_FILE_ERROR;
1963 deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
1964 if (deviceDriverStart==NULL) {
1965 status = DS_PROFILE_FILE_ERROR;
1968 deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION);
1969 deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
1970 if (deviceDriverEnd ==NULL) {
1971 status = DS_PROFILE_FILE_ERROR;
1976 tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
1977 if (tmpStart==NULL) {
1978 status = DS_PROFILE_FILE_ERROR;
1981 tmpStart+=strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
1982 tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END);
1983 if (tmpEnd ==NULL) {
1984 status = DS_PROFILE_FILE_ERROR;
1987 memcpy(tmp,tmpStart,tmpEnd-tmpStart);
1988 tmp[tmpEnd-tmpStart] = '\0';
1989 maxComputeUnits = atoi(tmp);
1992 tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ);
1993 if (tmpStart==NULL) {
1994 status = DS_PROFILE_FILE_ERROR;
1997 tmpStart+=strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ);
1998 tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ_END);
1999 if (tmpEnd ==NULL) {
2000 status = DS_PROFILE_FILE_ERROR;
2003 memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2004 tmp[tmpEnd-tmpStart] = '\0';
2005 maxClockFrequency = atoi(tmp);
2008 /* check if this device is on the system */
2009 for (i = 0; i < profile->numDevices; i++) {
2010 if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) {
2011 size_t actualDeviceNameLength;
2012 size_t driverVersionLength;
2014 actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
2015 driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
2016 if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
2017 && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
2018 && maxComputeUnits == profile->devices[i].oclMaxComputeUnits
2019 && maxClockFrequency == profile->devices[i].oclMaxClockFrequency
2020 && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
2021 && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==(int)0) {
2023 deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2024 if (deviceNameStart==NULL) {
2025 status = DS_PROFILE_FILE_ERROR;
2028 deviceScoreStart+=strlen(DS_TAG_SCORE);
2029 deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2030 status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2031 if (status != DS_SUCCESS) {
2039 else if (deviceType == DS_DEVICE_NATIVE_CPU) {
2040 for (i = 0; i < profile->numDevices; i++) {
2041 if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) {
2042 deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2043 if (deviceScoreStart==NULL) {
2044 status = DS_PROFILE_FILE_ERROR;
2047 deviceScoreStart+=strlen(DS_TAG_SCORE);
2048 deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2049 status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2050 if (status != DS_SUCCESS) {
2057 /* skip over the current one to find the next device */
2058 currentPosition = dataEnd+strlen(DS_TAG_DEVICE_END);
2062 if (contentStart!=NULL) free(contentStart);
2068 static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
2070 if (profile == NULL || num==NULL)
2071 return DS_MEMORY_ERROR;
2073 for (i = 0; i < profile->numDevices; i++) {
2074 if (profile->devices[i].score == NULL) {
2083 End of the OpenCL device selection infrastructure
2087 typedef double AccelerateScoreType;
2089 static ds_status AcceleratePerfEvaluator(ds_device *device,
2090 void *magick_unused(data))
2092 #define ACCELERATE_PERF_DIMEN "2048x1536"
2094 #define ReturnStatus(status) \
2097 RelinquishMagickOpenCLEnv(clEnv); \
2098 if (oldClEnv!=NULL) \
2099 defaultCLEnv = oldClEnv; \
2113 magick_unreferenced(data);
2116 ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2118 clEnv=AcquireMagickOpenCLEnv();
2119 exception=AcquireExceptionInfo();
2121 if (device->type == DS_DEVICE_NATIVE_CPU)
2124 MagickBooleanType flag=MagickTrue;
2125 SetMagickOpenCLEnvParamInternal(clEnv,
2126 MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
2129 else if (device->type == DS_DEVICE_OPENCL_DEVICE)
2132 SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
2133 sizeof(cl_device_id),&device->oclDeviceID,exception);
2136 ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2138 /* recompile the OpenCL kernels if it needs to */
2139 clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;
2141 InitOpenCLEnvInternal(clEnv,exception);
2142 oldClEnv=defaultCLEnv;
2145 /* microbenchmark */
2156 imageInfo=AcquireImageInfo();
2157 CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
2158 CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
2159 inputImage=ReadImage(imageInfo,exception);
2161 initAccelerateTimer(&timer);
2163 for (i=0; i<=NUM_ITER; i++)
2171 startAccelerateTimer(&timer);
2173 #ifdef MAGICKCORE_CLPERFMARKER
2174 clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
2177 bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
2178 unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
2180 resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,1.0,
2183 #ifdef MAGICKCORE_CLPERFMARKER
2184 clEndPerfMarkerAMD();
2188 stopAccelerateTimer(&timer);
2191 DestroyImage(bluredImage);
2193 DestroyImage(unsharpedImage);
2195 DestroyImage(resizedImage);
2197 DestroyImage(inputImage);
2199 /* end of microbenchmark */
2201 if (device->score == NULL)
2202 device->score=malloc(sizeof(AccelerateScoreType));
2203 *(AccelerateScoreType*)device->score=readAccelerateTimer(&timer);
2205 ReturnStatus(DS_SUCCESS);
2208 ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
2211 /* generate a string from the score */
2212 char* s = (char*)malloc(sizeof(char)*256);
2213 sprintf(s,"%.4f",*((AccelerateScoreType*)device->score));
2214 *serializedScore = (void*)s;
2215 *serializedScoreSize = strlen(s);
2219 return DS_SCORE_SERIALIZER_ERROR;
2223 ds_status AccelerateScoreDeserializer(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize) {
2225 /* convert the string back to an int */
2226 char* s = (char*)malloc(serializedScoreSize+1);
2227 memcpy(s, serializedScore, serializedScoreSize);
2228 s[serializedScoreSize] = (char)'\0';
2229 device->score = malloc(sizeof(AccelerateScoreType));
2230 *((AccelerateScoreType*)device->score) = (AccelerateScoreType)atof(s);
2235 return DS_SCORE_DESERIALIZER_ERROR;
2239 ds_status AccelerateScoreRelease(void* score) {
2247 #define IMAGEMAGICK_PROFILE_VERSION "ImageMagick Device Selection v0.9"
2248 #define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile"
2249 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
2251 MagickBooleanType mStatus = MagickFalse;
2253 ds_profile* profile;
2254 unsigned int numDeviceProfiled = 0;
2256 unsigned int bestDeviceIndex;
2257 AccelerateScoreType bestScore;
2258 char path[MaxTextExtent];
2259 MagickBooleanType flag;
2260 ds_evaluation_type profileType;
2262 LockDefaultOpenCLEnv();
2264 /* Initially, just set OpenCL to off */
2266 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2267 , sizeof(MagickBooleanType), &flag, exception);
2269 /* check and init the global lib */
2270 OpenCLLib=GetOpenCLLib();
2271 if (OpenCLLib==NULL)
2273 mStatus=InitOpenCLEnvInternal(clEnv, exception);
2277 status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
2278 if (status!=DS_SUCCESS) {
2279 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2283 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2284 ,GetOpenCLCachedFilesDirectory()
2285 ,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
2287 if (clEnv->regenerateProfile != MagickFalse) {
2288 profileType = DS_EVALUATE_ALL;
2291 readProfileFromFile(profile, AccelerateScoreDeserializer, path);
2292 profileType = DS_EVALUATE_NEW_ONLY;
2294 status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
2296 if (status!=DS_SUCCESS) {
2297 (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2300 if (numDeviceProfiled > 0) {
2301 status = writeProfileToFile(profile, AccelerateScoreSerializer, path);
2302 if (status!=DS_SUCCESS) {
2303 (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when saving the profile into a file", "'%s'", ".");
2307 /* pick the best device */
2308 bestDeviceIndex = 0;
2309 bestScore = *(AccelerateScoreType*)profile->devices[bestDeviceIndex].score;
2310 for (i = 1; i < profile->numDevices; i++) {
2311 AccelerateScoreType score = *(AccelerateScoreType*)profile->devices[i].score;
2312 if (score < bestScore) {
2313 bestDeviceIndex = i;
2318 /* set up clEnv with the best device */
2319 if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
2322 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2323 , sizeof(MagickBooleanType), &flag, exception);
2325 else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
2328 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2329 , sizeof(MagickBooleanType), &flag, exception);
2330 SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2331 , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
2334 status = DS_PERF_EVALUATOR_ERROR;
2337 mStatus=InitOpenCLEnvInternal(clEnv, exception);
2339 status = releaseDSProfile(profile, AccelerateScoreRelease);
2340 if (status!=DS_SUCCESS) {
2341 (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
2346 UnlockDefaultOpenCLEnv();
2352 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2356 + I n i t I m a g e M a g i c k O p e n C L %
2360 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2362 % InitImageMagickOpenCL() provides a simplified interface to initialize
2363 % the OpenCL environtment in ImageMagick
2365 % The format of the InitImageMagickOpenCL() method is:
2367 % MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
2368 % void* userSelectedDevice,
2369 % void* selectedDevice)
2371 % A description of each parameter follows:
2373 % o mode: OpenCL mode in ImageMagick, could be off,auto,user
2375 % o userSelectedDevice: when in user mode, a pointer to the selected
2378 % o selectedDevice: a pointer to cl_device_id where the selected
2379 % cl_device_id by ImageMagick could be returned
2381 % o exception: exception
2384 MagickExport MagickBooleanType InitImageMagickOpenCL(
2385 ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
2386 ExceptionInfo *exception)
2388 MagickBooleanType status = MagickTrue;
2389 MagickCLEnv clEnv = NULL;
2390 MagickBooleanType flag;
2392 clEnv = GetDefaultOpenCLEnv();
2396 case MAGICK_OPENCL_OFF:
2398 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2399 , sizeof(MagickBooleanType), &flag, exception);
2400 status = InitOpenCLEnv(clEnv, exception);
2403 *(cl_device_id*)selectedDevice = NULL;
2406 case MAGICK_OPENCL_DEVICE_SELECT_USER:
2408 if (userSelectedDevice == NULL)
2412 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2413 , sizeof(MagickBooleanType), &flag, exception);
2415 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2416 , sizeof(cl_device_id), userSelectedDevice,exception);
2418 status = InitOpenCLEnv(clEnv, exception);
2419 if (selectedDevice) {
2420 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2421 , sizeof(cl_device_id), selectedDevice, exception);
2425 case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
2427 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
2428 , sizeof(MagickBooleanType), &flag, exception);
2430 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
2431 , sizeof(MagickBooleanType), &flag, exception);
2433 /* fall through here!! */
2434 case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
2437 cl_device_id d = NULL;
2439 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2440 , sizeof(MagickBooleanType), &flag, exception);
2441 SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2442 , sizeof(cl_device_id), &d,exception);
2443 status = InitOpenCLEnv(clEnv, exception);
2444 if (selectedDevice) {
2445 GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2446 , sizeof(cl_device_id), selectedDevice, exception);
2457 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2458 const char *module,const char *function,const size_t line,
2459 const ExceptionType severity,const char *tag,const char *format,...) {
2465 status = MagickTrue;
2467 clEnv = GetDefaultOpenCLEnv();
2469 assert(exception != (ExceptionInfo *) NULL);
2470 assert(exception->signature == MagickSignature);
2473 cl_device_type dType;
2474 clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
2475 if (dType == CL_DEVICE_TYPE_CPU) {
2476 char buffer[MaxTextExtent];
2477 clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);
2479 /* Workaround for Intel OpenCL CPU runtime bug */
2480 /* Turn off OpenCL when a problem is detected! */
2481 if (strncmp(buffer, "Intel",5) == 0) {
2483 InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
2488 #ifdef OPENCLLOG_ENABLED
2492 va_start(operands,format);
2493 status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands);
2497 magick_unreferenced(module);
2498 magick_unreferenced(function);
2499 magick_unreferenced(line);
2500 magick_unreferenced(tag);
2501 magick_unreferenced(format);
2507 MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
2509 LockSemaphoreInfo(clEnv->lock);
2510 if (clEnv->seedsLock == NULL)
2512 ActivateSemaphoreInfo(&clEnv->seedsLock);
2514 LockSemaphoreInfo(clEnv->seedsLock);
2516 if (clEnv->seeds == NULL)
2519 clEnv->numGenerators = NUM_CL_RAND_GENERATORS;
2520 clEnv->seeds = clEnv->library->clCreateBuffer(clEnv->context, CL_MEM_READ_WRITE,
2521 clEnv->numGenerators*4*sizeof(unsigned int),
2523 if (clStatus != CL_SUCCESS)
2525 clEnv->seeds = NULL;
2530 cl_command_queue queue = NULL;
2531 unsigned int *seeds;
2533 queue = AcquireOpenCLCommandQueue(clEnv);
2534 seeds = (unsigned int*) clEnv->library->clEnqueueMapBuffer(queue, clEnv->seeds, CL_TRUE,
2536 clEnv->numGenerators*4
2537 *sizeof(unsigned int),
2538 0, NULL, NULL, &clStatus);
2539 if (clStatus!=CL_SUCCESS)
2541 clEnv->library->clReleaseMemObject(clEnv->seeds);
2545 for (i = 0; i < clEnv->numGenerators; i++) {
2546 RandomInfo* randomInfo = AcquireRandomInfo();
2547 const unsigned long* s = GetRandomInfoSeed(randomInfo);
2549 clEnv->randNormalize = GetRandomInfoNormalize(randomInfo);
2551 seeds[i*4] = (unsigned int) s[0];
2552 seeds[i*4+1] = (unsigned int) 0x50a7f451;
2553 seeds[i*4+2] = (unsigned int) 0x5365417e;
2554 seeds[i*4+3] = (unsigned int) 0xc3a4171a;
2556 randomInfo = DestroyRandomInfo(randomInfo);
2558 clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, clEnv->seeds, seeds, 0,
2560 clEnv->library->clFinish(queue);
2563 RelinquishOpenCLCommandQueue(clEnv, queue);
2566 UnlockSemaphoreInfo(clEnv->lock);
2567 return clEnv->seeds;
2570 MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv) {
2571 if (clEnv->seedsLock == NULL)
2573 ActivateSemaphoreInfo(&clEnv->seedsLock);
2576 UnlockSemaphoreInfo(clEnv->seedsLock);
2579 MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
2581 return clEnv->numGenerators;
2585 MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
2587 return clEnv->randNormalize;
2592 struct _MagickCLEnv {
2593 MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */
2596 MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
2601 MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(
2602 MagickCLEnv magick_unused(clEnv))
2604 magick_unreferenced(clEnv);
2610 * Return the OpenCL environment
2612 MagickExport MagickCLEnv GetDefaultOpenCLEnv(
2613 ExceptionInfo *magick_unused(exception))
2615 magick_unreferenced(exception);
2617 return (MagickCLEnv) NULL;
2620 MagickExport MagickCLEnv SetDefaultOpenCLEnv(
2621 MagickCLEnv magick_unused(clEnv))
2623 magick_unreferenced(clEnv);
2625 return (MagickCLEnv) NULL;
2628 MagickExport MagickBooleanType SetMagickOpenCLEnvParam(
2629 MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
2630 size_t magick_unused(dataSize),void *magick_unused(data),
2631 ExceptionInfo *magick_unused(exception))
2633 magick_unreferenced(clEnv);
2634 magick_unreferenced(param);
2635 magick_unreferenced(dataSize);
2636 magick_unreferenced(data);
2637 magick_unreferenced(exception);
2642 MagickExport MagickBooleanType GetMagickOpenCLEnvParam(
2643 MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
2644 size_t magick_unused(dataSize),void *magick_unused(data),
2645 ExceptionInfo *magick_unused(exception))
2647 magick_unreferenced(clEnv);
2648 magick_unreferenced(param);
2649 magick_unreferenced(dataSize);
2650 magick_unreferenced(data);
2651 magick_unreferenced(exception);
2656 MagickExport MagickBooleanType InitOpenCLEnv(MagickCLEnv magick_unused(clEnv),
2657 ExceptionInfo *magick_unused(exception))
2659 magick_unreferenced(clEnv);
2660 magick_unreferenced(exception);
2665 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(
2666 MagickCLEnv magick_unused(clEnv))
2668 magick_unreferenced(clEnv);
2670 return (cl_command_queue) NULL;
2673 MagickPrivate MagickBooleanType RelinquishCommandQueue(
2674 MagickCLEnv magick_unused(clEnv),cl_command_queue magick_unused(queue))
2676 magick_unreferenced(clEnv);
2677 magick_unreferenced(queue);
2682 MagickPrivate cl_kernel AcquireOpenCLKernel(
2683 MagickCLEnv magick_unused(clEnv),MagickOpenCLProgram magick_unused(program),
2684 const char *magick_unused(kernelName))
2686 magick_unreferenced(clEnv);
2687 magick_unreferenced(program);
2688 magick_unreferenced(kernelName);
2690 return (cl_kernel)NULL;
2693 MagickPrivate MagickBooleanType RelinquishOpenCLKernel(
2694 MagickCLEnv magick_unused(clEnv),cl_kernel magick_unused(kernel))
2696 magick_unreferenced(clEnv);
2697 magick_unreferenced(kernel);
2702 MagickPrivate unsigned long GetOpenCLDeviceLocalMemorySize(
2703 MagickCLEnv magick_unused(clEnv))
2705 magick_unreferenced(clEnv);
2710 MagickExport MagickBooleanType InitImageMagickOpenCL(
2711 ImageMagickOpenCLMode magick_unused(mode),
2712 void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
2713 ExceptionInfo *magick_unused(exception))
2715 magick_unreferenced(mode);
2716 magick_unreferenced(userSelectedDevice);
2717 magick_unreferenced(selectedDevice);
2718 magick_unreferenced(exception);
2724 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2725 const char *module,const char *function,const size_t line,
2726 const ExceptionType severity,const char *tag,const char *format,...)
2728 magick_unreferenced(exception);
2729 magick_unreferenced(module);
2730 magick_unreferenced(function);
2731 magick_unreferenced(line);
2732 magick_unreferenced(severity);
2733 magick_unreferenced(tag);
2734 magick_unreferenced(format);
2735 return(MagickFalse);
2739 MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
2741 magick_unreferenced(clEnv);
2746 MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv)
2748 magick_unreferenced(clEnv);
2751 MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
2753 magick_unreferenced(clEnv);
2757 MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
2759 magick_unreferenced(clEnv);
2763 #endif /* MAGICKCORE_OPENCL_SUPPORT */
2765 char* openclCachedFilesDirectory;
2766 SemaphoreInfo* openclCachedFilesDirectoryLock;
2769 const char* GetOpenCLCachedFilesDirectory() {
2770 if (openclCachedFilesDirectory == NULL) {
2771 if (openclCachedFilesDirectoryLock == NULL)
2773 ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
2775 LockSemaphoreInfo(openclCachedFilesDirectoryLock);
2776 if (openclCachedFilesDirectory == NULL) {
2777 char path[MaxTextExtent];
2780 struct stat attributes;
2781 MagickBooleanType status;
2785 home=GetEnvironmentValue("IMAGEMAGICK_OPENCL_CACHE_DIR");
2786 if (home == (char *) NULL)
2788 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2789 home=GetEnvironmentValue("LOCALAPPDATA");
2790 if (home == (char *) NULL)
2791 home=GetEnvironmentValue("APPDATA");
2792 if (home == (char *) NULL)
2793 home=GetEnvironmentValue("USERPROFILE");
2795 home=GetEnvironmentValue("HOME");
2799 if (home != (char *) NULL)
2801 int mkdirStatus = 0;
2805 /* first check if $HOME/.config exists */
2806 (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config",
2807 home,DirectorySeparator);
2808 status=GetPathAttributes(path,&attributes);
2809 if (status == MagickFalse)
2812 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2813 mkdirStatus = mkdir(path);
2815 mkdirStatus = mkdir(path, 0777);
2819 /* first check if $HOME/.config/ImageMagick exists */
2822 (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config%sImageMagick",
2823 home,DirectorySeparator,DirectorySeparator);
2825 status=GetPathAttributes(path,&attributes);
2826 if (status == MagickFalse)
2828 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2829 mkdirStatus = mkdir(path);
2831 mkdirStatus = mkdir(path, 0777);
2838 temp = (char*)AcquireMagickMemory(strlen(path)+1);
2839 CopyMagickString(temp,path,strlen(path)+1);
2841 home=DestroyString(home);
2843 openclCachedFilesDirectory = temp;
2845 UnlockSemaphoreInfo(openclCachedFilesDirectoryLock);
2847 return openclCachedFilesDirectory;
2850 void startAccelerateTimer(AccelerateTimer* timer) {
2852 QueryPerformanceCounter((LARGE_INTEGER*)&timer->_start);
2857 gettimeofday(&s, 0);
2858 timer->_start = (long long)s.tv_sec * (long long)1.0E3 + (long long)s.tv_usec / (long long)1.0E3;
2862 void stopAccelerateTimer(AccelerateTimer* timer) {
2865 QueryPerformanceCounter((LARGE_INTEGER*)&(n));
2868 gettimeofday(&s, 0);
2869 n = (long long)s.tv_sec * (long long)1.0E3+ (long long)s.tv_usec / (long long)1.0E3;
2873 timer->_clocks += n;
2876 void resetAccelerateTimer(AccelerateTimer* timer) {
2882 void initAccelerateTimer(AccelerateTimer* timer) {
2884 QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq);
2886 timer->_freq = (long long)1.0E3;
2888 resetAccelerateTimer(timer);
2891 double readAccelerateTimer(AccelerateTimer* timer) {
2892 return (double)timer->_clocks/(double)timer->_freq;
2896 /* create a function for OpenCL log */
2898 void OpenCLLog(const char* message) {
2900 #ifdef OPENCLLOG_ENABLED
2901 #define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
2904 if (getenv("MAGICK_OCL_LOG"))
2907 char path[MaxTextExtent];
2908 unsigned long allocSize;
2912 clEnv = GetDefaultOpenCLEnv();
2914 /* dump the source into a file */
2915 (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2916 ,GetOpenCLCachedFilesDirectory()
2917 ,DirectorySeparator,OPENCL_LOG_FILE);
2920 log = fopen(path, "ab");
2921 fwrite(message, sizeof(char), strlen(message), log);
2922 fwrite("\n", sizeof(char), 1, log);
2924 if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
2926 allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
2927 fprintf(log, "Devic Max Memory Alloc Size: %ld\n", allocSize);
2934 magick_unreferenced(message);