% MagickCore OpenCL Methods %
% %
% Software Design %
-% John Cristy %
+% Cristy %
% March 2000 %
% %
% %
#include "MagickCore/montage.h"
#include "MagickCore/morphology.h"
#include "MagickCore/nt-base.h"
+#include "MagickCore/nt-base-private.h"
#include "MagickCore/opencl.h"
#include "MagickCore/opencl-private.h"
#include "MagickCore/option.h"
#include "MagickCore/property.h"
#include "MagickCore/quantize.h"
#include "MagickCore/quantum.h"
+#include "MagickCore/random_.h"
+#include "MagickCore/random-private.h"
#include "MagickCore/resample.h"
#include "MagickCore/resource_.h"
#include "MagickCore/splay-tree.h"
#if defined(MAGICKCORE_OPENCL_SUPPORT)
-struct _MagickCLEnv {
- MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */
- MagickBooleanType OpenCLDisabled; /* whether if OpenCL has been explicitely disabled. */
+#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
+#define MAGICKCORE_OPENCL_MACOSX 1
+#endif
- /*OpenCL objects */
- cl_platform_id platform;
- cl_device_type deviceType;
- cl_device_id device;
- cl_context context;
- cl_program programs[MAGICK_OPENCL_NUM_PROGRAMS]; /* one program object maps one kernel source file */
+#define NUM_CL_RAND_GENERATORS 1024 /* number of random number generators running in parallel */
- SemaphoreInfo* lock;
-};
+/*
+ *
+ * Dynamic library loading functions
+ *
+ */
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+#else
+#include <dlfcn.h>
+#endif
+
+// dynamically load a library. returns NULL on failure
+void *OsLibraryLoad(const char *libraryName)
+{
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+ return (void *)LoadLibraryA(libraryName);
+#else
+ return (void *)dlopen(libraryName, RTLD_NOW);
+#endif
+}
+
+// get a function pointer from a loaded library. returns NULL on failure.
+void *OsLibraryGetFunctionAddress(void *library, const char *functionName)
+{
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+ if (!library || !functionName)
+ {
+ return NULL;
+ }
+ return (void *) GetProcAddress( (HMODULE)library, functionName);
+#else
+ if (!library || !functionName)
+ {
+ return NULL;
+ }
+ return (void *)dlsym(library, functionName);
+#endif
+}
+
+// unload a library.
+void OsLibraryUnload(void *library)
+{
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+ FreeLibrary( (HMODULE)library);
+#else
+ dlclose(library);
+#endif
+}
/*
if (clEnv != NULL)
{
memset(clEnv, 0, sizeof(struct _MagickCLEnv));
- AcquireSemaphoreInfo(&clEnv->lock);
+ ActivateSemaphoreInfo(&clEnv->lock);
}
return clEnv;
}
{
if (clEnv != (MagickCLEnv)NULL)
{
- RelinquishSemaphoreInfo(clEnv->lock);
+ DestroySemaphoreInfo(&clEnv->lock);
RelinquishMagickMemory(clEnv);
return MagickTrue;
}
MagickCLEnv defaultCLEnv;
SemaphoreInfo* defaultCLEnvLock;
+/*
+* OpenCL library
+*/
+MagickLibrary * OpenCLLib;
+SemaphoreInfo* OpenCLLibLock;
+
+
+static MagickBooleanType bindOpenCLFunctions(void* library)
+{
+#ifdef MAGICKCORE_OPENCL_MACOSX
+#define BIND(X) OpenCLLib->X= &X;
+#else
+#define BIND(X)\
+ if ((OpenCLLib->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(library,#X)) == NULL)\
+ return MagickFalse;
+#endif
+
+ BIND(clGetPlatformIDs);
+ BIND(clGetPlatformInfo);
+
+ BIND(clGetDeviceIDs);
+ BIND(clGetDeviceInfo);
+
+ BIND(clCreateContext);
+
+ BIND(clCreateBuffer);
+ BIND(clReleaseMemObject);
+
+ BIND(clCreateProgramWithSource);
+ BIND(clCreateProgramWithBinary);
+ BIND(clBuildProgram);
+ BIND(clGetProgramInfo);
+ BIND(clGetProgramBuildInfo);
+
+ BIND(clCreateKernel);
+ BIND(clReleaseKernel);
+ BIND(clSetKernelArg);
+
+ BIND(clFlush);
+ BIND(clFinish);
+
+ BIND(clEnqueueNDRangeKernel);
+ BIND(clEnqueueReadBuffer);
+ BIND(clEnqueueMapBuffer);
+ BIND(clEnqueueUnmapMemObject);
+
+ BIND(clCreateCommandQueue);
+ BIND(clReleaseCommandQueue);
+
+ return MagickTrue;
+}
+
+MagickLibrary * GetOpenCLLib()
+{
+ if (OpenCLLib == NULL)
+ {
+ if (OpenCLLibLock == NULL)
+ {
+ ActivateSemaphoreInfo(&OpenCLLibLock);
+ }
+
+ LockSemaphoreInfo(OpenCLLibLock);
+
+ OpenCLLib = (MagickLibrary *) AcquireMagickMemory (sizeof (MagickLibrary));
+
+ if (OpenCLLib != NULL)
+ {
+ MagickBooleanType status = MagickFalse;
+ void * library = NULL;
+
+#ifdef MAGICKCORE_OPENCL_MACOSX
+ status = bindOpenCLFunctions(library);
+#else
+
+ memset(OpenCLLib, 0, sizeof(MagickLibrary));
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+ library = OsLibraryLoad("OpenCL.dll");
+#else
+ library = OsLibraryLoad("libOpenCL.so");
+#endif
+ if (library)
+ status = bindOpenCLFunctions(library);
+
+ if (status==MagickTrue)
+ OpenCLLib->base=library;
+ else
+ OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
+#endif
+ }
+
+ UnlockSemaphoreInfo(OpenCLLibLock);
+ }
+
+
+ return OpenCLLib;
+}
+
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
{
if (defaultCLEnvLock == NULL)
{
- AcquireSemaphoreInfo(&defaultCLEnvLock);
+ ActivateSemaphoreInfo(&defaultCLEnvLock);
}
LockSemaphoreInfo(defaultCLEnvLock);
defaultCLEnv = AcquireMagickOpenCLEnv();
static void LockDefaultOpenCLEnv() {
if (defaultCLEnvLock == NULL)
{
- AcquireSemaphoreInfo(&defaultCLEnvLock);
+ ActivateSemaphoreInfo(&defaultCLEnvLock);
}
LockSemaphoreInfo(defaultCLEnvLock);
}
static void UnlockDefaultOpenCLEnv() {
if (defaultCLEnvLock == NULL)
{
- AcquireSemaphoreInfo(&defaultCLEnvLock);
+ ActivateSemaphoreInfo(&defaultCLEnvLock);
}
else
UnlockSemaphoreInfo(defaultCLEnvLock);
(void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "SetMagickOpenCLEnvParm cannot modify the OpenCL initialization state.", "'%s'", ".");
break;
+ case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
+ if (dataSize != sizeof(clEnv->disableProgramCache))
+ goto cleanup;
+ clEnv->disableProgramCache = *((MagickBooleanType*)data);
+ clEnv->OpenCLInitialized = MagickFalse;
+ status = MagickTrue;
+ break;
+
+ case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
+ if (dataSize != sizeof(clEnv->regenerateProfile))
+ goto cleanup;
+ clEnv->regenerateProfile = *((MagickBooleanType*)data);
+ clEnv->OpenCLInitialized = MagickFalse;
+ status = MagickTrue;
+ break;
+
default:
goto cleanup;
};
MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
, size_t dataSize, void* data, ExceptionInfo* exception)
{
- MagickBooleanType status;
+ MagickBooleanType
+ status;
+
+ magick_unreferenced(exception);
+
status = MagickFalse;
if (clEnv == NULL
status = MagickTrue;
break;
+ case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
+ if (dataSize != sizeof(clEnv->disableProgramCache))
+ goto cleanup;
+ *((MagickBooleanType*)data) = clEnv->disableProgramCache;
+ status = MagickTrue;
+ break;
+
+ case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
+ if (dataSize != sizeof(clEnv->regenerateProfile))
+ goto cleanup;
+ *((MagickBooleanType*)data) = clEnv->regenerateProfile;
+ status = MagickTrue;
+ break;
+
default:
goto cleanup;
};
%
*/
-MagickExport
+MagickPrivate
cl_context GetOpenCLContext(MagickCLEnv clEnv) {
if (clEnv == NULL)
return NULL;
static char* getBinaryCLProgramName(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
{
char* name;
+ char* ptr;
char path[MaxTextExtent];
char deviceName[MaxTextExtent];
const char* prefix = "magick_opencl";
- clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MaxTextExtent, deviceName, NULL);
- (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s_%s_%02d_%08x.bin"
- ,GetOpenCLCachedFilesDirectory()
- ,DirectorySeparator,prefix,deviceName, (unsigned int)prog, signature);
+ clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MaxTextExtent, deviceName, NULL);
+ ptr=deviceName;
+ /* strip out illegal characters for file names */
+ while (*ptr != '\0')
+ {
+ if ( *ptr == ' ' || *ptr == '\\' || *ptr == '/' || *ptr == ':' || *ptr == '*'
+ || *ptr == '?' || *ptr == '"' || *ptr == '<' || *ptr == '>' || *ptr == '|')
+ {
+ *ptr = '_';
+ }
+ ptr++;
+ }
+ (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s_%s_%02d_%08x_%.20g.bin",
+ GetOpenCLCachedFilesDirectory(),DirectorySeparator,prefix,deviceName,
+ (unsigned int) prog,signature,(double) sizeof(char*)*8);
name = (char*)AcquireMagickMemory(strlen(path)+1);
CopyMagickString(name,path,strlen(path)+1);
return name;
fileHandle = NULL;
saveSuccessful = MagickFalse;
- clStatus = clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binaryProgramSize, NULL);
+ clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binaryProgramSize, NULL);
if (clStatus != CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
}
binaryProgram = (unsigned char*) AcquireMagickMemory(binaryProgramSize);
- clStatus = clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARIES, sizeof(char*), &binaryProgram, NULL);
+ clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARIES, sizeof(char*), &binaryProgram, NULL);
if (clStatus != CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
return saveSuccessful;
}
-static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature, ExceptionInfo* exception)
+static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
{
MagickBooleanType loadSuccessful;
unsigned char* binaryProgram;
memset(binaryProgram, 0, length);
b_error |= fread(binaryProgram, 1, length, fileHandle) != length;
- clEnv->programs[prog] = clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
+ clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
if (clStatus != CL_SUCCESS
|| clBinaryStatus != CL_SUCCESS)
goto cleanup;
unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature;
/* try to load the binary first */
- if (!getenv("MAGICK_OCL_REC"))
- loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception);
+ if (clEnv->disableProgramCache != MagickTrue
+ && !getenv("MAGICK_OCL_REC"))
+ loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature);
if (loadSuccessful == MagickFalse)
{
/* Binary CL program unavailable, compile the program from source */
size_t programLength = strlen(MagickOpenCLProgramStrings[i]);
- clEnv->programs[i] = clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
+ clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
if (clStatus!=CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
}
}
- clStatus = clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
+ clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
if (clStatus!=CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
{
char* log;
size_t logSize;
- clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
+ clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
log = (char*)AcquireMagickMemory(logSize);
- clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);
+ clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);
(void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
,GetOpenCLCachedFilesDirectory()
if (clEnv->device != NULL)
{
- status = clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
+ status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
if (status != CL_SUCCESS) {
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
"Failed to get OpenCL platform from the selected device.", "(%d)", status);
clEnv->device = NULL;
/* Get the number of OpenCL platforms available */
- status = clGetPlatformIDs(0, NULL, &numPlatforms);
+ status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
goto cleanup;
}
- status = clGetPlatformIDs(numPlatforms, platforms, NULL);
+ status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
for (i = 0; i < numPlatforms; i++)
{
cl_uint numDevices;
- status = clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
+ status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
if (status != CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
- "clGetPlatformIDs failed.", "(%d)", status);
+ "clGetDeviceIDs failed.", "(%d)", status);
goto cleanup;
}
if (clEnv->device != NULL)
}
static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) {
- if (clEnv->OpenCLInitialized == MagickTrue
+ if (clEnv->OpenCLInitialized != MagickFalse
&& clEnv->platform != NULL
&& clEnv->device != NULL) {
clEnv->OpenCLDisabled = MagickFalse;
cl_int clStatus;
cl_context_properties cps[3];
-
+#ifdef MAGICKCORE_CLPERFMARKER
+ {
+ int status = clInitializePerfMarkerAMD();
+ if (status == AP_SUCCESS) {
+ //printf("PerfMarker successfully initialized\n");
+ }
+ }
+#endif
clEnv->OpenCLInitialized = MagickTrue;
- if (clEnv->OpenCLDisabled == MagickTrue)
+
+ /* check and init the global lib */
+ OpenCLLib=GetOpenCLLib();
+ if (OpenCLLib)
+ {
+ clEnv->library=OpenCLLib;
+ }
+ else
+ {
+ /* turn off opencl */
+ MagickBooleanType flag;
+ flag = MagickTrue;
+ SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
+ , sizeof(MagickBooleanType), &flag, exception);
+ }
+
+ if (clEnv->OpenCLDisabled != MagickFalse)
goto cleanup;
clEnv->OpenCLDisabled = MagickTrue;
cps[0] = CL_CONTEXT_PLATFORM;
cps[1] = (cl_context_properties)clEnv->platform;
cps[2] = 0;
- clEnv->context = clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
+ clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
if (clStatus != CL_SUCCESS)
{
(void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
}
status = EnableOpenCLInternal(clEnv);
+
cleanup:
return status;
}
%
*/
-MagickExport
+MagickPrivate
cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
{
if (clEnv != NULL)
- return clCreateCommandQueue(clEnv->context, clEnv->device, 0, NULL);
+ return clEnv->library->clCreateCommandQueue(clEnv->context, clEnv->device, 0, NULL);
else
return NULL;
}
%
%
*/
-MagickExport
+MagickPrivate
MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv, cl_command_queue queue)
{
if (clEnv != NULL)
{
- return ((clReleaseCommandQueue(queue) == CL_SUCCESS) ? MagickTrue:MagickFalse);
+ return ((clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ? MagickTrue:MagickFalse);
}
else
return MagickFalse;
%
*/
-MagickExport
+MagickPrivate
cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv, MagickOpenCLProgram program, const char* kernelName)
{
cl_int clStatus;
cl_kernel kernel = NULL;
if (clEnv != NULL && kernelName!=NULL)
{
- kernel = clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
+ kernel = clEnv->library->clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
}
return kernel;
}
%
*/
-MagickExport
+MagickPrivate
MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel)
{
MagickBooleanType status = MagickFalse;
if (clEnv != NULL && kernel != NULL)
{
- status = ((clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
+ status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
}
return status;
}
%
*/
-MagickExport
+MagickPrivate
unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
{
cl_ulong localMemorySize;
- clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
+ clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
return (unsigned long)localMemorySize;
}
-MagickExport
+MagickPrivate
unsigned long GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv clEnv)
{
cl_ulong maxMemAllocSize;
- clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
+ clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
return (unsigned long)maxMemAllocSize;
}
*/
-#define DS_DEVICE_NAME_LENGTH 256
-
typedef enum {
DS_SUCCESS = 0
,DS_INVALID_PROFILE = 1000
memset(profile, 0, sizeof(ds_profile));
- clGetPlatformIDs(0, NULL, &numPlatforms);
+ OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
if (numPlatforms > 0) {
platforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id));
if (platforms == NULL) {
status = DS_MEMORY_ERROR;
goto cleanup;
}
- clGetPlatformIDs(numPlatforms, platforms, NULL);
+ OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
for (i = 0; i < (unsigned int)numPlatforms; i++) {
cl_uint num;
- clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &num);
- numDevices+=num;
+ if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
+ numDevices+=num;
}
}
continue;
break;
}
- clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num);
+ if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
+ continue;
for (j = 0; j < num; j++, next++) {
- char buffer[DS_DEVICE_NAME_LENGTH];
size_t length;
profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
profile->devices[next].oclDeviceID = devices[j];
- clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
- , DS_DEVICE_NAME_LENGTH, &buffer, NULL);
- length = strlen(buffer);
- profile->devices[next].oclDeviceName = (char*)malloc(length+1);
- memcpy(profile->devices[next].oclDeviceName, buffer, length+1);
+ OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
+ , 0, NULL, &length);
+ profile->devices[next].oclDeviceName = (char*)malloc(sizeof(char)*length);
+ OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
+ , length, profile->devices[next].oclDeviceName, NULL);
- clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
- , DS_DEVICE_NAME_LENGTH, &buffer, NULL);
- length = strlen(buffer);
- profile->devices[next].oclDriverVersion = (char*)malloc(length+1);
- memcpy(profile->devices[next].oclDriverVersion, buffer, length+1);
+ OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
+ , 0, NULL, &length);
+ profile->devices[next].oclDriverVersion = (char*)malloc(sizeof(char)*length);
+ OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
+ , length, profile->devices[next].oclDriverVersion, NULL);
- clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
+ OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
, sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);
- clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
+ OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
, sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);
}
}
break;
/* else fall through */
case DS_EVALUATE_ALL:
- evaluatorStatus = evaluator(profile->devices+i, evaluatorData);
+ evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
if (evaluatorStatus != DS_SUCCESS) {
status = evaluatorStatus;
return status;
}
versionStringLength = strlen(profile->version);
- if (versionStringLength!=(dataEnd-dataStart)
+ if (versionStringLength!=(size_t)(dataEnd-dataStart)
|| strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
/* version mismatch */
status = DS_PROFILE_FILE_ERROR;
currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);
/* parse the device information */
+DisableMSCWarning(4127)
while (1) {
+RestoreMSCWarning
unsigned int i;
const char* deviceTypeStart;
actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
- if (actualDeviceNameLength == (deviceNameEnd - deviceNameStart)
- && driverVersionLength == (deviceDriverEnd - deviceDriverStart)
+ if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
+ && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
&& maxComputeUnits == profile->devices[i].oclMaxComputeUnits
&& maxClockFrequency == profile->devices[i].oclMaxClockFrequency
&& strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
return status;
}
+
+#if 0
static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
unsigned int i;
if (profile == NULL || num==NULL)
*num=0;
for (i = 0; i < profile->numDevices; i++) {
if (profile->devices[i].score == NULL) {
- *num++;
+ (*num)++;
}
}
return DS_SUCCESS;
}
+#endif
/*
End of the OpenCL device selection infrastructure
*/
+typedef double AccelerateScoreType;
-typedef struct _AccelerateTimer {
- long long _freq;
- long long _clocks;
- long long _start;
-} AccelerateTimer;
-
-static void startAccelerateTimer(AccelerateTimer* timer) {
-#ifdef _WIN32
- QueryPerformanceCounter((LARGE_INTEGER*)&timer->_start);
-
-
-#else
- struct timeval s;
- gettimeofday(&s, 0);
- timer->_start = (long long)s.tv_sec * (long long)1.0E3 + (long long)s.tv_usec / (long long)1.0E3;
-#endif
-}
-
-static void stopAccelerateTimer(AccelerateTimer* timer) {
- long long n=0;
-#ifdef _WIN32
- QueryPerformanceCounter((LARGE_INTEGER*)&(n));
-#else
- struct timeval s;
- gettimeofday(&s, 0);
- n = (long long)s.tv_sec * (long long)1.0E3+ (long long)s.tv_usec / (long long)1.0E3;
-#endif
- n -= timer->_start;
- timer->_start = 0;
- timer->_clocks += n;
-}
-
-static void resetAccelerateTimer(AccelerateTimer* timer) {
- timer->_clocks = 0;
- timer->_start = 0;
+static ds_status AcceleratePerfEvaluator(ds_device *device,
+ void *magick_unused(data))
+{
+#define ACCELERATE_PERF_DIMEN "2048x1536"
+#define NUM_ITER 2
+#define ReturnStatus(status) \
+{ \
+ if (clEnv!=NULL) \
+ RelinquishMagickOpenCLEnv(clEnv); \
+ if (oldClEnv!=NULL) \
+ defaultCLEnv = oldClEnv; \
+ return status; \
}
+ AccelerateTimer
+ timer;
-static void initAccelerateTimer(AccelerateTimer* timer) {
-#ifdef _WIN32
- QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq);
-#else
- timer->_freq = (long long)1.0E3;
-#endif
- resetAccelerateTimer(timer);
-}
+ ExceptionInfo
+ *exception=NULL;
-double readAccelerateTimer(AccelerateTimer* timer) { return (double)timer->_clocks/(double)timer->_freq; };
+ MagickCLEnv
+ clEnv=NULL,
+ oldClEnv=NULL;
+ magick_unreferenced(data);
-typedef double AccelerateScoreType;
+ if (device == NULL)
+ ReturnStatus(DS_PERF_EVALUATOR_ERROR);
-static ds_status AcceleratePerfEvaluator(ds_device* device, void* data) {
+ clEnv=AcquireMagickOpenCLEnv();
+ exception=AcquireExceptionInfo();
- ds_status status = DS_SUCCESS;
- MagickCLEnv clEnv = NULL;
- MagickCLEnv oldClEnv = NULL;
- ExceptionInfo* exception = NULL;
- AccelerateTimer timer;
+ if (device->type == DS_DEVICE_NATIVE_CPU)
+ {
+ /* CPU device */
+ MagickBooleanType flag=MagickTrue;
+ SetMagickOpenCLEnvParamInternal(clEnv,
+ MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
+ &flag,exception);
+ }
+ else if (device->type == DS_DEVICE_OPENCL_DEVICE)
+ {
+ /* OpenCL device */
+ SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
+ sizeof(cl_device_id),&device->oclDeviceID,exception);
+ }
+ else
+ ReturnStatus(DS_PERF_EVALUATOR_ERROR);
- if (device == NULL) {
- status = DS_PERF_EVALUATOR_ERROR;
- goto cleanup;
- }
+ /* recompile the OpenCL kernels if it needs to */
+ clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;
- clEnv = AcquireMagickOpenCLEnv();
- exception = AcquireExceptionInfo();
-
- if (device->type == DS_DEVICE_NATIVE_CPU) {
- /* CPU device */
- MagickBooleanType flag = MagickTrue;
- SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
- , sizeof(MagickBooleanType), &flag, exception);
- }
- else if (device->type == DS_DEVICE_OPENCL_DEVICE) {
- /* OpenCL device */
- SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
- , sizeof(cl_device_id), &device->oclDeviceID,exception);
- }
- else {
- status = DS_PERF_EVALUATOR_ERROR;
- goto cleanup;
- }
- InitOpenCLEnvInternal(clEnv, exception);
- oldClEnv = defaultCLEnv;
- defaultCLEnv = clEnv;
+ InitOpenCLEnvInternal(clEnv,exception);
+ oldClEnv=defaultCLEnv;
+ defaultCLEnv=clEnv;
/* microbenchmark */
{
-#define ACCELERATE_PERF_DIMEN "2048x1536"
-#define NUM_ITER 2
+ Image
+ *inputImage;
+
+ ImageInfo
+ *imageInfo;
- Image* inputImage;
- ImageInfo* imageInfo;
- int i;
+ int
+ i;
- imageInfo = AcquireImageInfo();
+ imageInfo=AcquireImageInfo();
CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
- inputImage = ReadImage(imageInfo,exception);
+ inputImage=ReadImage(imageInfo,exception);
initAccelerateTimer(&timer);
- for (i = 0; i <=NUM_ITER; i++) {
-
- Image* bluredImage;
- Image* unsharpedImage;
- Image* resizedImage;
+ for (i=0; i<=NUM_ITER; i++)
+ {
+ Image
+ *bluredImage,
+ *resizedImage,
+ *unsharpedImage;
if (i > 0)
startAccelerateTimer(&timer);
#ifdef MAGICKCORE_CLPERFMARKER
- clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
+ clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
#endif
- bluredImage = BlurImage(inputImage, 10.0f, 3.5f, exception);
- unsharpedImage = UnsharpMaskImage(bluredImage, 2.0f,2.0f,50.0f,10.0f,exception);
- resizedImage = ResizeImage(unsharpedImage,640,480,LanczosFilter,1.0,exception);
+ bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
+ unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
+ exception);
+ resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,1.0,
+ exception);
#ifdef MAGICKCORE_CLPERFMARKER
- clEndPerfMarkerAMD();
+ clEndPerfMarkerAMD();
#endif
if (i > 0)
stopAccelerateTimer(&timer);
- if (bluredImage) DestroyImage(bluredImage);
- if (unsharpedImage) DestroyImage(unsharpedImage);
- if (resizedImage) DestroyImage(resizedImage);
+ if (bluredImage)
+ DestroyImage(bluredImage);
+ if (unsharpedImage)
+ DestroyImage(unsharpedImage);
+ if (resizedImage)
+ DestroyImage(resizedImage);
}
DestroyImage(inputImage);
}
/* end of microbenchmark */
- if (device->score == NULL) {
- device->score = malloc(sizeof(AccelerateScoreType));
- }
- *(AccelerateScoreType*)device->score = readAccelerateTimer(&timer);
+ if (device->score == NULL)
+ device->score=malloc(sizeof(AccelerateScoreType));
+ *(AccelerateScoreType*)device->score=readAccelerateTimer(&timer);
-cleanup:
- if (clEnv!=NULL)
- RelinquishMagickOpenCLEnv(clEnv);
- if (oldClEnv!=NULL)
- defaultCLEnv = oldClEnv;
- return status;
+ ReturnStatus(DS_SUCCESS);
}
-
-
ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
if (device
&& device->score) {
unsigned int bestDeviceIndex;
AccelerateScoreType bestScore;
char path[MaxTextExtent];
-
+ MagickBooleanType flag;
+ ds_evaluation_type profileType;
LockDefaultOpenCLEnv();
+ /* Initially, just set OpenCL to off */
+ flag = MagickTrue;
+ SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
+ , sizeof(MagickBooleanType), &flag, exception);
+
+ /* check and init the global lib */
+ OpenCLLib=GetOpenCLLib();
+ if (OpenCLLib==NULL)
+ {
+ mStatus=InitOpenCLEnvInternal(clEnv, exception);
+ goto cleanup;
+ }
+
status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
if (status!=DS_SUCCESS) {
(void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
,GetOpenCLCachedFilesDirectory()
,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
- readProfileFromFile(profile, AccelerateScoreDeserializer, path);
- status = profileDevices(profile, DS_EVALUATE_NEW_ONLY, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
+ if (clEnv->regenerateProfile != MagickFalse) {
+ profileType = DS_EVALUATE_ALL;
+ }
+ else {
+ readProfileFromFile(profile, AccelerateScoreDeserializer, path);
+ profileType = DS_EVALUATE_NEW_ONLY;
+ }
+ status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
+
if (status!=DS_SUCCESS) {
(void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
goto cleanup;
/* set up clEnv with the best device */
if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
/* CPU device */
- MagickBooleanType flag = MagickTrue;
+ flag = MagickTrue;
SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
, sizeof(MagickBooleanType), &flag, exception);
}
else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
/* OpenCL device */
+ flag = MagickFalse;
+ SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
+ , sizeof(MagickBooleanType), &flag, exception);
SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
, sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
}
status = DS_PERF_EVALUATOR_ERROR;
goto cleanup;
}
- InitOpenCLEnvInternal(clEnv, exception);
+ mStatus=InitOpenCLEnvInternal(clEnv, exception);
status = releaseDSProfile(profile, AccelerateScoreRelease);
if (status!=DS_SUCCESS) {
(void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
}
- mStatus = MagickTrue;
cleanup:
% o exception: exception
%
*/
-MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
- void* userSelectedDevice,
- void* selectedDevice,
- ExceptionInfo* exception) {
-
- MagickBooleanType status = MagickTrue;
+MagickExport MagickBooleanType InitImageMagickOpenCL(
+ ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
+ ExceptionInfo *exception)
+{
+ MagickBooleanType status = MagickFalse;
MagickCLEnv clEnv = NULL;
MagickBooleanType flag;
- exception = AcquireExceptionInfo();
clEnv = GetDefaultOpenCLEnv();
if (clEnv!=NULL) {
switch(mode) {
}
break;
+ case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
+ flag = MagickTrue;
+ SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
+ , sizeof(MagickBooleanType), &flag, exception);
+ flag = MagickTrue;
+ SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
+ , sizeof(MagickBooleanType), &flag, exception);
+
+ /* fall through here!! */
case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
default:
{
}
+MagickPrivate
+MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
+ const char *module,const char *function,const size_t line,
+ const ExceptionType severity,const char *tag,const char *format,...) {
+ MagickBooleanType
+ status;
+
+ MagickCLEnv clEnv;
+
+ status = MagickTrue;
+
+ clEnv = GetDefaultOpenCLEnv();
+
+ assert(exception != (ExceptionInfo *) NULL);
+ assert(exception->signature == MagickSignature);
+
+ if (severity!=0) {
+ cl_device_type dType;
+ clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
+ if (dType == CL_DEVICE_TYPE_CPU) {
+ char buffer[MaxTextExtent];
+ clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);
+
+ /* Workaround for Intel OpenCL CPU runtime bug */
+ /* Turn off OpenCL when a problem is detected! */
+ if (strncmp(buffer, "Intel",5) == 0) {
+
+ InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
+ }
+ }
+ }
+
+#ifdef OPENCLLOG_ENABLED
+ {
+ va_list
+ operands;
+ va_start(operands,format);
+ status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands);
+ va_end(operands);
+ }
+#else
+ magick_unreferenced(module);
+ magick_unreferenced(function);
+ magick_unreferenced(line);
+ magick_unreferenced(tag);
+ magick_unreferenced(format);
+#endif
+
+ return(status);
+}
+
+MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
+{
+ LockSemaphoreInfo(clEnv->lock);
+ if (clEnv->seedsLock == NULL)
+ {
+ ActivateSemaphoreInfo(&clEnv->seedsLock);
+ }
+ LockSemaphoreInfo(clEnv->seedsLock);
+
+ if (clEnv->seeds == NULL)
+ {
+ cl_int clStatus;
+ clEnv->numGenerators = NUM_CL_RAND_GENERATORS;
+ clEnv->seeds = clEnv->library->clCreateBuffer(clEnv->context, CL_MEM_READ_WRITE,
+ clEnv->numGenerators*4*sizeof(unsigned int),
+ NULL, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ clEnv->seeds = NULL;
+ }
+ else
+ {
+ unsigned int i;
+ cl_command_queue queue = NULL;
+ unsigned int *seeds;
+
+ queue = AcquireOpenCLCommandQueue(clEnv);
+ seeds = (unsigned int*) clEnv->library->clEnqueueMapBuffer(queue, clEnv->seeds, CL_TRUE,
+ CL_MAP_WRITE, 0,
+ clEnv->numGenerators*4
+ *sizeof(unsigned int),
+ 0, NULL, NULL, &clStatus);
+ if (clStatus!=CL_SUCCESS)
+ {
+ clEnv->library->clReleaseMemObject(clEnv->seeds);
+ goto cleanup;
+ }
+
+ for (i = 0; i < clEnv->numGenerators; i++) {
+ RandomInfo* randomInfo = AcquireRandomInfo();
+ const unsigned long* s = GetRandomInfoSeed(randomInfo);
+ if (i == 0)
+ clEnv->randNormalize = GetRandomInfoNormalize(randomInfo);
+
+ seeds[i*4] = (unsigned int) s[0];
+ seeds[i*4+1] = (unsigned int) 0x50a7f451;
+ seeds[i*4+2] = (unsigned int) 0x5365417e;
+ seeds[i*4+3] = (unsigned int) 0xc3a4171a;
+
+ randomInfo = DestroyRandomInfo(randomInfo);
+ }
+ clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, clEnv->seeds, seeds, 0,
+ NULL, NULL);
+ clEnv->library->clFinish(queue);
+cleanup:
+ if (queue != NULL)
+ RelinquishOpenCLCommandQueue(clEnv, queue);
+ }
+ }
+ UnlockSemaphoreInfo(clEnv->lock);
+ return clEnv->seeds;
+}
+
+MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv) {
+ if (clEnv->seedsLock == NULL)
+ {
+ ActivateSemaphoreInfo(&clEnv->seedsLock);
+ }
+ else
+ UnlockSemaphoreInfo(clEnv->seedsLock);
+}
+
+MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
+{
+ return clEnv->numGenerators;
+}
+
+
+MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
+{
+ return clEnv->randNormalize;
+}
+
#else
struct _MagickCLEnv {
MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */
};
-extern MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
+MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
{
return NULL;
}
-extern MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(
+MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(
MagickCLEnv magick_unused(clEnv))
{
magick_unreferenced(clEnv);
return MagickFalse;
}
-MagickExport cl_command_queue AcquireOpenCLCommandQueue(
+MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(
MagickCLEnv magick_unused(clEnv))
{
magick_unreferenced(clEnv);
return (cl_command_queue) NULL;
}
-MagickExport MagickBooleanType RelinquishCommandQueue(
+MagickPrivate MagickBooleanType RelinquishCommandQueue(
MagickCLEnv magick_unused(clEnv),cl_command_queue magick_unused(queue))
{
magick_unreferenced(clEnv);
return MagickFalse;
}
-MagickExport cl_kernel AcquireOpenCLKernel(
+MagickPrivate cl_kernel AcquireOpenCLKernel(
MagickCLEnv magick_unused(clEnv),MagickOpenCLProgram magick_unused(program),
const char *magick_unused(kernelName))
{
return (cl_kernel)NULL;
}
-MagickExport MagickBooleanType RelinquishOpenCLKernel(
+MagickPrivate MagickBooleanType RelinquishOpenCLKernel(
MagickCLEnv magick_unused(clEnv),cl_kernel magick_unused(kernel))
{
magick_unreferenced(clEnv);
return MagickFalse;
}
-MagickExport unsigned long GetOpenCLDeviceLocalMemorySize(
+MagickPrivate unsigned long GetOpenCLDeviceLocalMemorySize(
MagickCLEnv magick_unused(clEnv))
{
magick_unreferenced(clEnv);
return 0;
}
-MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
- void* userSelectedDevice,
- void* selectedDevice,
- ExceptionInfo* exception)
+MagickExport MagickBooleanType InitImageMagickOpenCL(
+ ImageMagickOpenCLMode magick_unused(mode),
+ void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
+ ExceptionInfo *magick_unused(exception))
{
magick_unreferenced(mode);
magick_unreferenced(userSelectedDevice);
return MagickFalse;
}
+
+MagickPrivate
+MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
+ const char *module,const char *function,const size_t line,
+ const ExceptionType severity,const char *tag,const char *format,...)
+{
+ magick_unreferenced(exception);
+ magick_unreferenced(module);
+ magick_unreferenced(function);
+ magick_unreferenced(line);
+ magick_unreferenced(severity);
+ magick_unreferenced(tag);
+ magick_unreferenced(format);
+ return(MagickFalse);
+}
+
+
+MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
+{
+ magick_unreferenced(clEnv);
+ return NULL;
+}
+
+
+MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv)
+{
+ magick_unreferenced(clEnv);
+}
+
+MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
+{
+ magick_unreferenced(clEnv);
+ return 0;
+}
+
+MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
+{
+ magick_unreferenced(clEnv);
+ return 0.0f;
+}
+
#endif /* MAGICKCORE_OPENCL_SUPPORT */
char* openclCachedFilesDirectory;
SemaphoreInfo* openclCachedFilesDirectoryLock;
-MagickExport
+MagickPrivate
const char* GetOpenCLCachedFilesDirectory() {
if (openclCachedFilesDirectory == NULL) {
if (openclCachedFilesDirectoryLock == NULL)
{
- AcquireSemaphoreInfo(&openclCachedFilesDirectoryLock);
+ ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
}
LockSemaphoreInfo(openclCachedFilesDirectoryLock);
if (openclCachedFilesDirectory == NULL) {
struct stat attributes;
MagickBooleanType status;
-#ifdef MAGICKCORE_WINDOWS_SUPPORT
- home=GetEnvironmentValue("LOCALAPPDATA");
- if (home == (char *) NULL)
- home=GetEnvironmentValue("APPDATA");
+
+
+ home=GetEnvironmentValue("IMAGEMAGICK_OPENCL_CACHE_DIR");
if (home == (char *) NULL)
- home=GetEnvironmentValue("USERPROFILE");
+ {
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+ home=GetEnvironmentValue("LOCALAPPDATA");
+ if (home == (char *) NULL)
+ home=GetEnvironmentValue("APPDATA");
+ if (home == (char *) NULL)
+ home=GetEnvironmentValue("USERPROFILE");
#else
- home=GetEnvironmentValue("HOME");
+ home=GetEnvironmentValue("HOME");
#endif
+ }
+
if (home != (char *) NULL)
{
+ int mkdirStatus = 0;
/*
- Search $HOME/.magick.
*/
- (void) FormatLocaleString(path,MaxTextExtent,"%s%s.magick",home,
- DirectorySeparator);
- home=DestroyString(home);
- temp = (char*)AcquireMagickMemory(strlen(path)+1);
- CopyMagickString(temp,path,strlen(path)+1);
+
+ /* first check if $HOME/.config exists */
+ (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config",
+ home,DirectorySeparator);
status=GetPathAttributes(path,&attributes);
- if (status == MagickFalse) {
+ if (status == MagickFalse)
+ {
+
#ifdef MAGICKCORE_WINDOWS_SUPPORT
- mkdir(path);
+ mkdirStatus = mkdir(path);
#else
- mkdir(path, 0777);
+ mkdirStatus = mkdir(path, 0777);
#endif
}
+
+ /* first check if $HOME/.config/ImageMagick exists */
+ if (mkdirStatus==0)
+ {
+ (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config%sImageMagick",
+ home,DirectorySeparator,DirectorySeparator);
+
+ status=GetPathAttributes(path,&attributes);
+ if (status == MagickFalse)
+ {
+#ifdef MAGICKCORE_WINDOWS_SUPPORT
+ mkdirStatus = mkdir(path);
+#else
+ mkdirStatus = mkdir(path, 0777);
+#endif
+ }
+ }
+
+ if (mkdirStatus==0)
+ {
+ temp = (char*)AcquireMagickMemory(strlen(path)+1);
+ CopyMagickString(temp,path,strlen(path)+1);
+ }
+ home=DestroyString(home);
}
openclCachedFilesDirectory = temp;
}
return openclCachedFilesDirectory;
}
+void startAccelerateTimer(AccelerateTimer* timer) {
+#ifdef _WIN32
+ QueryPerformanceCounter((LARGE_INTEGER*)&timer->_start);
-/* create a loggin function */
-MagickExport
+
+#else
+ struct timeval s;
+ gettimeofday(&s, 0);
+ timer->_start = (long long)s.tv_sec * (long long)1.0E3 + (long long)s.tv_usec / (long long)1.0E3;
+#endif
+}
+
+void stopAccelerateTimer(AccelerateTimer* timer) {
+ long long n=0;
+#ifdef _WIN32
+ QueryPerformanceCounter((LARGE_INTEGER*)&(n));
+#else
+ struct timeval s;
+ gettimeofday(&s, 0);
+ n = (long long)s.tv_sec * (long long)1.0E3+ (long long)s.tv_usec / (long long)1.0E3;
+#endif
+ n -= timer->_start;
+ timer->_start = 0;
+ timer->_clocks += n;
+}
+
+void resetAccelerateTimer(AccelerateTimer* timer) {
+ timer->_clocks = 0;
+ timer->_start = 0;
+}
+
+
+void initAccelerateTimer(AccelerateTimer* timer) {
+#ifdef _WIN32
+ QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq);
+#else
+ timer->_freq = (long long)1.0E3;
+#endif
+ resetAccelerateTimer(timer);
+}
+
+double readAccelerateTimer(AccelerateTimer* timer) {
+ return (double)timer->_clocks/(double)timer->_freq;
+};
+
+
+/* create a function for OpenCL log */
+MagickPrivate
void OpenCLLog(const char* message) {
+#ifdef OPENCLLOG_ENABLED
#define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
FILE* log;
- if (message) {
- char path[MaxTextExtent];
+ if (getenv("MAGICK_OCL_LOG"))
+ {
+ if (message) {
+ char path[MaxTextExtent];
+ unsigned long allocSize;
+
+ MagickCLEnv clEnv;
+
+ clEnv = GetDefaultOpenCLEnv();
- /* dump the source into a file */
- (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
- ,GetOpenCLCachedFilesDirectory()
- ,DirectorySeparator,OPENCL_LOG_FILE);
+ /* dump the source into a file */
+ (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
+ ,GetOpenCLCachedFilesDirectory()
+ ,DirectorySeparator,OPENCL_LOG_FILE);
- log = fopen(path, "ab");
- fwrite(message, sizeof(char), strlen(message), log);
- fwrite("\n", sizeof(char), 1, log);
- fclose(log);
+ log = fopen(path, "ab");
+ fwrite(message, sizeof(char), strlen(message), log);
+ fwrite("\n", sizeof(char), 1, log);
+
+ if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
+ {
+ allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
+ fprintf(log, "Devic Max Memory Alloc Size: %ld\n", allocSize);
+ }
+
+ fclose(log);
+ }
}
+#else
+ magick_unreferenced(message);
+#endif
}
+
projects / imagemagick / blobdiff