king
/
v811_spc009
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
You can not select more than 25 topics Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
Branches Tags
${ item.name }
Create tag ${ searchTerm }
Create branch ${ searchTerm }
from 'main'
${ noResults }
v811_spc009/external/ImageMagick/MagickCore/opencl.c

3312 lines
104 KiB
Raw Permalink Blame History Escape

This file contains invisible Unicode characters!

This file contains invisible Unicode characters that may be processed differently from what appears below. If your use case is intentional and legitimate, you can safely ignore this warning. Use the Escape button to reveal hidden characters.

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% OOO PPPP EEEEE N N CCCC L %
% O O P P E NN N C L %
% O O PPPP EEE N N N C L %
% O O P E N NN C L %
% OOO P EEEEE N N CCCC LLLLL %
% %
% %
% MagickCore OpenCL Methods %
% %
% Software Design %
% Cristy %
% March 2000 %
% %
% %
% Copyright 1999-2021 ImageMagick Studio LLC, a non-profit organization %
% dedicated to making software imaging solutions freely available. %
% %
% You may not use this file except in compliance with the License. You may %
% obtain a copy of the License at %
% %
% https://imagemagick.org/script/license.php %
% %
% Unless required by applicable law or agreed to in writing, software %
% distributed under the License is distributed on an "AS IS" BASIS, %
% WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. %
% See the License for the specific language governing permissions and %
% limitations under the License. %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%
%
*/
/*
Include declarations.
*/
#include "MagickCore/studio.h"
#include "MagickCore/artifact.h"
#include "MagickCore/cache.h"
#include "MagickCore/cache-private.h"
#include "MagickCore/color.h"
#include "MagickCore/compare.h"
#include "MagickCore/constitute.h"
#include "MagickCore/configure.h"
#include "MagickCore/distort.h"
#include "MagickCore/draw.h"
#include "MagickCore/effect.h"
#include "MagickCore/exception.h"
#include "MagickCore/exception-private.h"
#include "MagickCore/fx.h"
#include "MagickCore/gem.h"
#include "MagickCore/geometry.h"
#include "MagickCore/image.h"
#include "MagickCore/image-private.h"
#include "MagickCore/layer.h"
#include "MagickCore/mime-private.h"
#include "MagickCore/memory_.h"
#include "MagickCore/memory-private.h"
#include "MagickCore/monitor.h"
#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/policy.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"
#include "MagickCore/semaphore.h"
#include "MagickCore/statistic.h"
#include "MagickCore/string_.h"
#include "MagickCore/string-private.h"
#include "MagickCore/token.h"
#include "MagickCore/utility.h"
#include "MagickCore/utility-private.h"
#if defined(MAGICKCORE_OPENCL_SUPPORT)
#if defined(MAGICKCORE_LTDL_DELEGATE)
#include "ltdl.h"
#endif
#ifndef MAGICKCORE_WINDOWS_SUPPORT
#include <dlfcn.h>
#endif
#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
#define MAGICKCORE_OPENCL_MACOSX 1
#endif
/*
Define declarations.
*/
#define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"
/*
Typedef declarations.
*/
typedef struct
{
long long freq;
long long clocks;
long long start;
} AccelerateTimer;
typedef struct
{
char
*name,
*platform_name,
*vendor_name,
*version;
cl_uint
max_clock_frequency,
max_compute_units;
double
score;
} MagickCLDeviceBenchmark;
/*
Forward declarations.
*/
static MagickBooleanType
HasOpenCLDevices(MagickCLEnv,ExceptionInfo *),
LoadOpenCLLibrary(void);
static MagickCLDevice
RelinquishMagickCLDevice(MagickCLDevice);
static MagickCLEnv
RelinquishMagickCLEnv(MagickCLEnv);
static void
BenchmarkOpenCLDevices(MagickCLEnv);
extern const char
*accelerateKernels, *accelerateKernels2;
/* OpenCL library */
MagickLibrary
*openCL_library;
/* Default OpenCL environment */
MagickCLEnv
default_CLEnv;
MagickThreadType
test_thread_id=0;
SemaphoreInfo
*openCL_lock;
/* Cached location of the OpenCL cache files */
char
*cache_directory;
SemaphoreInfo
*cache_directory_lock;
static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
MagickCLDevice b)
{
if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
(LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
(LocaleCompare(a->name,b->name) == 0) &&
(LocaleCompare(a->version,b->version) == 0) &&
(a->max_clock_frequency == b->max_clock_frequency) &&
(a->max_compute_units == b->max_compute_units))
return(MagickTrue);
return(MagickFalse);
}
static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
MagickCLDeviceBenchmark *b)
{
if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
(LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
(LocaleCompare(a->name,b->name) == 0) &&
(LocaleCompare(a->version,b->version) == 0) &&
(a->max_clock_frequency == b->max_clock_frequency) &&
(a->max_compute_units == b->max_compute_units))
return(MagickTrue);
return(MagickFalse);
}
static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
{
size_t
i;
if (clEnv->devices != (MagickCLDevice *) NULL)
{
for (i = 0; i < clEnv->number_devices; i++)
clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
}
clEnv->number_devices=0;
}
static inline MagickBooleanType MagickCreateDirectory(const char *path)
{
int
status;
#ifdef MAGICKCORE_WINDOWS_SUPPORT
status=mkdir(path);
#else
status=mkdir(path,0777);
#endif
return(status == 0 ? MagickTrue : MagickFalse);
}
static inline void InitAccelerateTimer(AccelerateTimer *timer)
{
#ifdef _WIN32
QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
#else
timer->freq=(long long)1.0E3;
#endif
timer->clocks=0;
timer->start=0;
}
static inline double ReadAccelerateTimer(AccelerateTimer *timer)
{
return (double)timer->clocks/(double)timer->freq;
}
static inline 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 inline void StopAccelerateTimer(AccelerateTimer *timer)
{
long long
n;
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 const char *GetOpenCLCacheDirectory()
{
if (cache_directory == (char *) NULL)
{
if (cache_directory_lock == (SemaphoreInfo *) NULL)
ActivateSemaphoreInfo(&cache_directory_lock);
LockSemaphoreInfo(cache_directory_lock);
if (cache_directory == (char *) NULL)
{
char
*home,
path[MagickPathExtent],
*temp;
MagickBooleanType
status;
struct stat
attributes;
temp=(char *) NULL;
home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
if (home == (char *) NULL)
{
home=GetEnvironmentValue("XDG_CACHE_HOME");
#if defined(MAGICKCORE_WINDOWS_SUPPORT) || defined(__MINGW32__)
if (home == (char *) NULL)
home=GetEnvironmentValue("LOCALAPPDATA");
if (home == (char *) NULL)
home=GetEnvironmentValue("APPDATA");
if (home == (char *) NULL)
home=GetEnvironmentValue("USERPROFILE");
#endif
}
if (home != (char *) NULL)
{
/* first check if $HOME exists */
(void) FormatLocaleString(path,MagickPathExtent,"%s",home);
status=GetPathAttributes(path,&attributes);
if (status == MagickFalse)
status=MagickCreateDirectory(path);
/* first check if $HOME/ImageMagick exists */
if (status != MagickFalse)
{
(void) FormatLocaleString(path,MagickPathExtent,
"%s%sImageMagick",home,DirectorySeparator);
status=GetPathAttributes(path,&attributes);
if (status == MagickFalse)
status=MagickCreateDirectory(path);
}
if (status != MagickFalse)
{
temp=(char*) AcquireCriticalMemory(strlen(path)+1);
CopyMagickString(temp,path,strlen(path)+1);
}
home=DestroyString(home);
}
else
{
home=GetEnvironmentValue("HOME");
if (home != (char *) NULL)
{
/* first check if $HOME/.cache exists */
(void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache",
home,DirectorySeparator);
status=GetPathAttributes(path,&attributes);
if (status == MagickFalse)
status=MagickCreateDirectory(path);
/* first check if $HOME/.cache/ImageMagick exists */
if (status != MagickFalse)
{
(void) FormatLocaleString(path,MagickPathExtent,
"%s%s.cache%sImageMagick",home,DirectorySeparator,
DirectorySeparator);
status=GetPathAttributes(path,&attributes);
if (status == MagickFalse)
status=MagickCreateDirectory(path);
}
if (status != MagickFalse)
{
temp=(char*) AcquireCriticalMemory(strlen(path)+1);
CopyMagickString(temp,path,strlen(path)+1);
}
home=DestroyString(home);
}
}
if (temp == (char *) NULL)
{
temp=AcquireString("?");
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Cannot use cache directory: \"%s\"",path);
}
else
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Using cache directory: \"%s\"",temp);
cache_directory=temp;
}
UnlockSemaphoreInfo(cache_directory_lock);
}
if (*cache_directory == '?')
return((const char *) NULL);
return(cache_directory);
}
static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
{
MagickCLDevice
device;
size_t
i,
j;
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Selecting device for type: %d",(int) type);
for (i = 0; i < clEnv->number_devices; i++)
clEnv->devices[i]->enabled=MagickFalse;
for (i = 0; i < clEnv->number_devices; i++)
{
device=clEnv->devices[i];
if (device->type != type)
continue;
device->enabled=MagickTrue;
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Selected device: %s",device->name);
for (j = i+1; j < clEnv->number_devices; j++)
{
MagickCLDevice
other_device;
other_device=clEnv->devices[j];
if (IsSameOpenCLDevice(device,other_device))
other_device->enabled=MagickTrue;
}
}
}
static size_t StringSignature(const char* string)
{
size_t
n,
i,
j,
signature,
stringLength;
union
{
const char* s;
const size_t* u;
} p;
stringLength=(size_t) strlen(string);
signature=stringLength;
n=stringLength/sizeof(size_t);
p.s=string;
for (i = 0; i < n; i++)
signature^=p.u[i];
if (n * sizeof(size_t) != stringLength)
{
char
padded[4];
j=n*sizeof(size_t);
for (i = 0; i < 4; i++, j++)
{
if (j < stringLength)
padded[i]=p.s[j];
else
padded[i]=0;
}
p.s=padded;
signature^=p.u[0];
}
return(signature);
}
static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
{
ssize_t
i;
for (i=0; i < (ssize_t) info->event_count; i++)
openCL_library->clReleaseEvent(info->events[i]);
info->events=(cl_event *) RelinquishMagickMemory(info->events);
if (info->buffer != (cl_mem) NULL)
openCL_library->clReleaseMemObject(info->buffer);
RelinquishSemaphoreInfo(&info->events_semaphore);
ReleaseOpenCLDevice(info->device);
RelinquishMagickMemory(info);
}
/*
Provide call to OpenCL library methods
*/
MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
cl_mem_flags flags,size_t size,void *host_ptr)
{
return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr,
(cl_int *) NULL));
}
MagickPrivate void ReleaseOpenCLKernel(cl_kernel kernel)
{
(void) openCL_library->clReleaseKernel(kernel);
}
MagickPrivate void ReleaseOpenCLMemObject(cl_mem memobj)
{
(void) openCL_library->clReleaseMemObject(memobj);
}
MagickPrivate void RetainOpenCLMemObject(cl_mem memobj)
{
(void) openCL_library->clRetainMemObject(memobj);
}
MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,size_t arg_index,
size_t arg_size,const void *arg_value)
{
return(openCL_library->clSetKernelArg(kernel,(cl_uint) arg_index,arg_size,
arg_value));
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ A c q u i r e M a g i c k C L C a c h e I n f o %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AcquireMagickCLCacheInfo() acquires an OpenCL cache info structure.
%
% The format of the AcquireMagickCLCacheInfo method is:
%
% MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
% Quantum *pixels,const MagickSizeType length)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
%
% o pixels: the pixel buffer of the image.
%
% o length: the length of the pixel buffer.
%
*/
MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
Quantum *pixels,const MagickSizeType length)
{
cl_int
status;
MagickCLCacheInfo
info;
info=(MagickCLCacheInfo) AcquireCriticalMemory(sizeof(*info));
(void) memset(info,0,sizeof(*info));
LockSemaphoreInfo(openCL_lock);
device->requested++;
UnlockSemaphoreInfo(openCL_lock);
info->device=device;
info->length=length;
info->pixels=pixels;
info->events_semaphore=AcquireSemaphoreInfo();
info->buffer=openCL_library->clCreateBuffer(device->context,
CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,
&status);
if (status == CL_SUCCESS)
return(info);
DestroyMagickCLCacheInfo(info);
return((MagickCLCacheInfo) NULL);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% A c q u i r e M a g i c k C L D e v i c e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AcquireMagickCLDevice() acquires an OpenCL device
%
% The format of the AcquireMagickCLDevice method is:
%
% MagickCLDevice AcquireMagickCLDevice()
%
*/
static MagickCLDevice AcquireMagickCLDevice()
{
MagickCLDevice
device;
device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device));
if (device != NULL)
{
(void) memset(device,0,sizeof(*device));
ActivateSemaphoreInfo(&device->lock);
device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
device->command_queues_index=-1;
device->enabled=MagickTrue;
}
return(device);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% A c q u i r e M a g i c k C L E n v %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AcquireMagickCLEnv() allocates the MagickCLEnv structure
%
*/
static MagickCLEnv AcquireMagickCLEnv(void)
{
const char
*option;
MagickCLEnv
clEnv;
clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv));
if (clEnv != (MagickCLEnv) NULL)
{
(void) memset(clEnv,0,sizeof(*clEnv));
ActivateSemaphoreInfo(&clEnv->lock);
clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
clEnv->enabled=MagickFalse;
option=getenv("MAGICK_OCL_DEVICE");
if (option != (const char *) NULL)
{
if ((IsStringTrue(option) != MagickFalse) ||
(strcmp(option,"GPU") == 0) ||
(strcmp(option,"CPU") == 0))
clEnv->enabled=MagickTrue;
}
}
return clEnv;
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ 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 %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AcquireOpenCLCommandQueue() acquires an OpenCL command queue
%
% The format of the AcquireOpenCLCommandQueue method is:
%
% cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
%
*/
MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
{
cl_command_queue
queue;
cl_command_queue_properties
properties;
assert(device != (MagickCLDevice) NULL);
LockSemaphoreInfo(device->lock);
if ((device->profile_kernels == MagickFalse) &&
(device->command_queues_index >= 0))
{
queue=device->command_queues[device->command_queues_index--];
UnlockSemaphoreInfo(device->lock);
}
else
{
UnlockSemaphoreInfo(device->lock);
properties=0;
if (device->profile_kernels != MagickFalse)
properties=CL_QUEUE_PROFILING_ENABLE;
queue=openCL_library->clCreateCommandQueue(device->context,
device->deviceID,properties,(cl_int *) NULL);
}
return(queue);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ A c q u i r e O p e n C L K e r n e l %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AcquireOpenCLKernel() acquires an OpenCL kernel
%
% The format of the AcquireOpenCLKernel method is:
%
% cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
% MagickOpenCLProgram program, const char* kernelName)
%
% A description of each parameter follows:
%
% o clEnv: the OpenCL environment.
%
% o program: the OpenCL program module that the kernel belongs to.
%
% o kernelName: the name of the kernel
%
*/
MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
const char *kernel_name)
{
cl_kernel
kernel;
assert(device != (MagickCLDevice) NULL);
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Using kernel: %s",
kernel_name);
kernel=openCL_library->clCreateKernel(device->program,kernel_name,
(cl_int *) NULL);
return(kernel);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% A u t o S e l e c t O p e n C L D e v i c e s %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AutoSelectOpenCLDevices() determines the best device based on the
% information from the micro-benchmark.
%
% The format of the AutoSelectOpenCLDevices method is:
%
% void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception)
%
% A description of each parameter follows:
%
% o clEnv: the OpenCL environment.
%
% o exception: return any errors or warnings in this structure.
%
*/
static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
{
char
keyword[MagickPathExtent],
*token;
const char
*q;
MagickCLDeviceBenchmark
*device_benchmark;
size_t
i,
extent;
if (xml == (char *) NULL)
return;
device_benchmark=(MagickCLDeviceBenchmark *) NULL;
token=AcquireString(xml);
extent=strlen(token)+MagickPathExtent;
for (q=(char *) xml; *q != '\0'; )
{
/*
Interpret XML.
*/
(void) GetNextToken(q,&q,extent,token);
if (*token == '\0')
break;
(void) CopyMagickString(keyword,token,MagickPathExtent);
if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
{
/*
Doctype element.
*/
while ((LocaleNCompare(q,"]>",2) != 0) && (*q != '\0'))
(void) GetNextToken(q,&q,extent,token);
continue;
}
if (LocaleNCompare(keyword,"<!--",4) == 0)
{
/*
Comment element.
*/
while ((LocaleNCompare(q,"->",2) != 0) && (*q != '\0'))
(void) GetNextToken(q,&q,extent,token);
continue;
}
if (LocaleCompare(keyword,"<device") == 0)
{
/*
Device element.
*/
device_benchmark=(MagickCLDeviceBenchmark *) AcquireQuantumMemory(1,
sizeof(*device_benchmark));
if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
break;
(void) memset(device_benchmark,0,sizeof(*device_benchmark));
device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
continue;
}
if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
continue;
if (LocaleCompare(keyword,"/>") == 0)
{
if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
{
if (LocaleCompare(device_benchmark->name,"CPU") == 0)
clEnv->cpu_score=device_benchmark->score;
else
{
MagickCLDevice
device;
/*
Set the score for all devices that match this device.
*/
for (i = 0; i < clEnv->number_devices; i++)
{
device=clEnv->devices[i];
if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
device->score=device_benchmark->score;
}
}
}
device_benchmark->platform_name=RelinquishMagickMemory(
device_benchmark->platform_name);
device_benchmark->vendor_name=RelinquishMagickMemory(
device_benchmark->vendor_name);
device_benchmark->name=RelinquishMagickMemory(device_benchmark->name);
device_benchmark->version=RelinquishMagickMemory(
device_benchmark->version);
device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
device_benchmark);
continue;
}
(void) GetNextToken(q,(const char **) NULL,extent,token);
if (*token != '=')
continue;
(void) GetNextToken(q,&q,extent,token);
(void) GetNextToken(q,&q,extent,token);
switch (*keyword)
{
case 'M':
case 'm':
{
if (LocaleCompare((char *) keyword,"maxClockFrequency") == 0)
{
device_benchmark->max_clock_frequency=StringToInteger(token);
break;
}
if (LocaleCompare((char *) keyword,"maxComputeUnits") == 0)
{
device_benchmark->max_compute_units=StringToInteger(token);
break;
}
break;
}
case 'N':
case 'n':
{
if (LocaleCompare((char *) keyword,"name") == 0)
device_benchmark->name=ConstantString(token);
break;
}
case 'P':
case 'p':
{
if (LocaleCompare((char *) keyword,"platform") == 0)
device_benchmark->platform_name=ConstantString(token);
break;
}
case 'S':
case 's':
{
if (LocaleCompare((char *) keyword,"score") == 0)
device_benchmark->score=StringToDouble(token,(char **) NULL);
break;
}
case 'V':
case 'v':
{
if (LocaleCompare((char *) keyword,"vendor") == 0)
device_benchmark->vendor_name=ConstantString(token);
if (LocaleCompare((char *) keyword,"version") == 0)
device_benchmark->version=ConstantString(token);
break;
}
default:
break;
}
}
token=(char *) RelinquishMagickMemory(token);
device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
device_benchmark);
}
static MagickBooleanType CanWriteProfileToFile(const char *filename)
{
FILE
*profileFile;
profileFile=fopen(filename,"ab");
if (profileFile == (FILE *) NULL)
{
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Unable to save profile to: \"%s\"",filename);
return(MagickFalse);
}
fclose(profileFile);
return(MagickTrue);
}
static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
{
char
filename[MagickPathExtent];
StringInfo
*option;
size_t
i;
(void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
/*
We don't run the benchmark when we can not write out a device profile. The
first GPU device will be used.
*/
#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
if (CanWriteProfileToFile(filename) == MagickFalse)
#endif
{
for (i = 0; i < clEnv->number_devices; i++)
clEnv->devices[i]->score=1.0;
SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
return(MagickFalse);
}
option=ConfigureFileToStringInfo(filename);
LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(option));
option=DestroyStringInfo(option);
return(MagickTrue);
}
static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
{
const char
*option;
double
best_score;
MagickBooleanType
benchmark;
size_t
i;
option=getenv("MAGICK_OCL_DEVICE");
if (option != (const char *) NULL)
{
if (strcmp(option,"GPU") == 0)
SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
else if (strcmp(option,"CPU") == 0)
SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
}
if (LoadOpenCLBenchmarks(clEnv) == MagickFalse)
return;
benchmark=MagickFalse;
if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
benchmark=MagickTrue;
else
{
for (i = 0; i < clEnv->number_devices; i++)
{
if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
{
benchmark=MagickTrue;
break;
}
}
}
if (benchmark != MagickFalse)
BenchmarkOpenCLDevices(clEnv);
best_score=clEnv->cpu_score;
for (i = 0; i < clEnv->number_devices; i++)
best_score=MagickMin(clEnv->devices[i]->score,best_score);
for (i = 0; i < clEnv->number_devices; i++)
{
if (clEnv->devices[i]->score != best_score)
clEnv->devices[i]->enabled=MagickFalse;
}
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% B e n c h m a r k O p e n C L D e v i c e s %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% BenchmarkOpenCLDevices() benchmarks the OpenCL devices and the CPU to help
% the automatic selection of the best device.
%
% The format of the BenchmarkOpenCLDevices method is:
%
% void BenchmarkOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
%
% A description of each parameter follows:
%
% o clEnv: the OpenCL environment.
%
% o exception: return any errors or warnings
*/
static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
{
AccelerateTimer
timer;
ExceptionInfo
*exception;
Image
*inputImage;
ImageInfo
*imageInfo;
size_t
i;
exception=AcquireExceptionInfo();
imageInfo=AcquireImageInfo();
CloneString(&imageInfo->size,"2048x1536");
CopyMagickString(imageInfo->filename,"xc:none",MagickPathExtent);
inputImage=ReadImage(imageInfo,exception);
InitAccelerateTimer(&timer);
for (i=0; i<=2; i++)
{
Image
*bluredImage,
*resizedImage,
*unsharpedImage;
if (i > 0)
StartAccelerateTimer(&timer);
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,
exception);
/*
We need this to get a proper performance benchmark, the operations
are executed asynchronous.
*/
if (is_cpu == MagickFalse)
{
CacheInfo
*cache_info;
cache_info=(CacheInfo *) resizedImage->cache;
if (cache_info->opencl != (MagickCLCacheInfo) NULL)
openCL_library->clWaitForEvents(cache_info->opencl->event_count,
cache_info->opencl->events);
}
if (i > 0)
StopAccelerateTimer(&timer);
if (bluredImage != (Image *) NULL)
DestroyImage(bluredImage);
if (unsharpedImage != (Image *) NULL)
DestroyImage(unsharpedImage);
if (resizedImage != (Image *) NULL)
DestroyImage(resizedImage);
}
DestroyImage(inputImage);
return(ReadAccelerateTimer(&timer));
}
static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
MagickCLDevice device)
{
testEnv->devices[0]=device;
default_CLEnv=testEnv;
device->score=RunOpenCLBenchmark(MagickFalse);
default_CLEnv=clEnv;
testEnv->devices[0]=(MagickCLDevice) NULL;
}
static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
{
char
filename[MagickPathExtent];
FILE
*cache_file;
MagickCLDevice
device;
size_t
i,
j;
(void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
GetOpenCLCacheDirectory(),DirectorySeparator,
IMAGEMAGICK_PROFILE_FILE);
cache_file=fopen_utf8(filename,"wb");
if (cache_file == (FILE *) NULL)
return;
fwrite("<devices>\n",sizeof(char),10,cache_file);
fprintf(cache_file," <device name=\"CPU\" score=\"%.4g\"/>\n",
clEnv->cpu_score);
for (i = 0; i < clEnv->number_devices; i++)
{
MagickBooleanType
duplicate;
device=clEnv->devices[i];
duplicate=MagickFalse;
for (j = 0; j < i; j++)
{
if (IsSameOpenCLDevice(clEnv->devices[j],device))
{
duplicate=MagickTrue;
break;
}
}
if (duplicate)
continue;
if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
fprintf(cache_file," <device platform=\"%s\" vendor=\"%s\" name=\"%s\"\
version=\"%s\" maxClockFrequency=\"%d\" maxComputeUnits=\"%d\"\
score=\"%.4g\"/>\n",
device->platform_name,device->vendor_name,device->name,device->version,
(int)device->max_clock_frequency,(int)device->max_compute_units,
device->score);
}
fwrite("</devices>",sizeof(char),10,cache_file);
fclose(cache_file);
}
static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
{
MagickCLDevice
device;
MagickCLEnv
testEnv;
size_t
i,
j;
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Starting benchmark");
testEnv=AcquireMagickCLEnv();
testEnv->library=openCL_library;
testEnv->devices=(MagickCLDevice *) AcquireCriticalMemory(
sizeof(MagickCLDevice));
testEnv->number_devices=1;
testEnv->benchmark_thread_id=GetMagickThreadId();
testEnv->initialized=MagickTrue;
for (i = 0; i < clEnv->number_devices; i++)
clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
for (i = 0; i < clEnv->number_devices; i++)
{
device=clEnv->devices[i];
if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
RunDeviceBenckmark(clEnv,testEnv,device);
/* Set the score on all the other devices that are the same */
for (j = i+1; j < clEnv->number_devices; j++)
{
MagickCLDevice
other_device;
other_device=clEnv->devices[j];
if (IsSameOpenCLDevice(device,other_device))
other_device->score=device->score;
}
}
testEnv->enabled=MagickFalse;
default_CLEnv=testEnv;
clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
default_CLEnv=clEnv;
testEnv=RelinquishMagickCLEnv(testEnv);
CacheOpenCLBenchmarks(clEnv);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% C o m p i l e O p e n C L K e r n e l %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% CompileOpenCLKernel() compiles the kernel for the specified device. The
% kernel will be cached on disk to reduce the compilation time.
%
% The format of the CompileOpenCLKernel method is:
%
% MagickBooleanType AcquireOpenCLKernel(MagickCLDevice clEnv,
% unsigned int signature,const char *kernel,const char *options,
% ExceptionInfo *exception)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
%
% o kernel: the source code of the kernel.
%
% o options: options for the compiler.
%
% o signature: a number to uniquely identify the kernel
%
% o exception: return any errors or warnings in this structure.
%
*/
static void CacheOpenCLKernel(MagickCLDevice device,char *filename,
ExceptionInfo *exception)
{
cl_uint
status;
size_t
binaryProgramSize;
unsigned char
*binaryProgram;
status=openCL_library->clGetProgramInfo(device->program,
CL_PROGRAM_BINARY_SIZES,sizeof(size_t),&binaryProgramSize,NULL);
if (status != CL_SUCCESS)
return;
binaryProgram=(unsigned char*) AcquireQuantumMemory(1,binaryProgramSize);
if (binaryProgram == (unsigned char *) NULL)
{
(void) ThrowMagickException(exception,GetMagickModule(),
ResourceLimitError,"MemoryAllocationFailed","`%s'",filename);
return;
}
status=openCL_library->clGetProgramInfo(device->program,
CL_PROGRAM_BINARIES,sizeof(unsigned char*),&binaryProgram,NULL);
if (status == CL_SUCCESS)
{
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Creating cache file: \"%s\"",filename);
(void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
}
binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
}
static MagickBooleanType LoadCachedOpenCLKernels(MagickCLDevice device,
const char *filename)
{
cl_int
binaryStatus,
status;
ExceptionInfo
*sans_exception;
size_t
length;
unsigned char
*binaryProgram;
sans_exception=AcquireExceptionInfo();
binaryProgram=(unsigned char *) FileToBlob(filename,~0UL,&length,
sans_exception);
sans_exception=DestroyExceptionInfo(sans_exception);
if (binaryProgram == (unsigned char *) NULL)
return(MagickFalse);
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Loaded cached kernels: \"%s\"",filename);
device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
&device->deviceID,&length,(const unsigned char**)&binaryProgram,
&binaryStatus,&status);
binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
MagickTrue);
}
static void LogOpenCLBuildFailure(MagickCLDevice device,const char *kernel,
ExceptionInfo *exception)
{
char
filename[MagickPathExtent],
*log;
size_t
log_size;
(void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.cl");
(void) remove_utf8(filename);
(void) BlobToFile(filename,kernel,strlen(kernel),exception);
openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
CL_PROGRAM_BUILD_LOG,0,NULL,&log_size);
log=(char*)AcquireCriticalMemory(log_size);
openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
CL_PROGRAM_BUILD_LOG,log_size,log,&log_size);
(void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.log");
(void) remove_utf8(filename);
(void) BlobToFile(filename,log,log_size,exception);
log=(char*)RelinquishMagickMemory(log);
}
static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
const char *kernel,const char *options,size_t signature,
ExceptionInfo *exception)
{
char
deviceName[MagickPathExtent],
filename[MagickPathExtent],
*ptr;
cl_int
status;
MagickBooleanType
loaded;
size_t
length;
(void) CopyMagickString(deviceName,device->name,MagickPathExtent);
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(filename,MagickPathExtent,
"%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
DirectorySeparator,"magick_opencl",deviceName,(unsigned int) signature,
(double) sizeof(char*)*8);
loaded=LoadCachedOpenCLKernels(device,filename);
if (loaded == MagickFalse)
{
/* Binary CL program unavailable, compile the program from source */
length=strlen(kernel);
device->program=openCL_library->clCreateProgramWithSource(
device->context,1,&kernel,&length,&status);
if (status != CL_SUCCESS)
return(MagickFalse);
}
status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
options,NULL,NULL);
if (status != CL_SUCCESS)
{
(void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
"clBuildProgram failed.","(%d)",(int)status);
LogOpenCLBuildFailure(device,kernel,exception);
return(MagickFalse);
}
/* Save the binary to a file to avoid re-compilation of the kernels */
if (loaded == MagickFalse)
CacheOpenCLKernel(device,filename,exception);
return(MagickTrue);
}
static cl_event* CopyOpenCLEvents(MagickCLCacheInfo first,
MagickCLCacheInfo second,cl_uint *event_count)
{
cl_event
*events;
size_t
i;
size_t
j;
assert(first != (MagickCLCacheInfo) NULL);
assert(event_count != (cl_uint *) NULL);
events=(cl_event *) NULL;
LockSemaphoreInfo(first->events_semaphore);
if (second != (MagickCLCacheInfo) NULL)
LockSemaphoreInfo(second->events_semaphore);
*event_count=first->event_count;
if (second != (MagickCLCacheInfo) NULL)
*event_count+=second->event_count;
if (*event_count > 0)
{
events=AcquireQuantumMemory(*event_count,sizeof(*events));
if (events == (cl_event *) NULL)
*event_count=0;
else
{
j=0;
for (i=0; i < first->event_count; i++, j++)
events[j]=first->events[i];
if (second != (MagickCLCacheInfo) NULL)
{
for (i=0; i < second->event_count; i++, j++)
events[j]=second->events[i];
}
}
}
UnlockSemaphoreInfo(first->events_semaphore);
if (second != (MagickCLCacheInfo) NULL)
UnlockSemaphoreInfo(second->events_semaphore);
return(events);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ C o p y M a g i c k C L C a c h e I n f o %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% CopyMagickCLCacheInfo() copies the memory from the device into host memory.
%
% The format of the CopyMagickCLCacheInfo method is:
%
% void CopyMagickCLCacheInfo(MagickCLCacheInfo info)
%
% A description of each parameter follows:
%
% o info: the OpenCL cache info.
%
*/
MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
{
cl_command_queue
queue;
cl_event
*events;
cl_uint
event_count;
Quantum
*pixels;
if (info == (MagickCLCacheInfo) NULL)
return((MagickCLCacheInfo) NULL);
events=CopyOpenCLEvents(info,(MagickCLCacheInfo) NULL,&event_count);
if (events != (cl_event *) NULL)
{
queue=AcquireOpenCLCommandQueue(info->device);
pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
CL_MAP_READ | CL_MAP_WRITE,0,info->length,event_count,events,
(cl_event *) NULL,(cl_int *) NULL);
assert(pixels == info->pixels);
ReleaseOpenCLCommandQueue(info->device,queue);
events=(cl_event *) RelinquishMagickMemory(events);
}
return(RelinquishMagickCLCacheInfo(info,MagickFalse));
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ D u m p O p e n C L P r o f i l e D a t a %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% DumpOpenCLProfileData() dumps the kernel profile data.
%
% The format of the DumpProfileData method is:
%
% void DumpProfileData()
%
*/
MagickPrivate void DumpOpenCLProfileData()
{
#define OpenCLLog(message) \
fwrite(message,sizeof(char),strlen(message),log); \
fwrite("\n",sizeof(char),1,log);
char
buf[4096],
filename[MagickPathExtent],
indent[160];
FILE
*log;
size_t
i,
j;
if (default_CLEnv == (MagickCLEnv) NULL)
return;
for (i = 0; i < default_CLEnv->number_devices; i++)
if (default_CLEnv->devices[i]->profile_kernels != MagickFalse)
break;
if (i == default_CLEnv->number_devices)
return;
(void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
log=fopen_utf8(filename,"wb");
if (log == (FILE *) NULL)
return;
for (i = 0; i < default_CLEnv->number_devices; i++)
{
MagickCLDevice
device;
device=default_CLEnv->devices[i];
if ((device->profile_kernels == MagickFalse) ||
(device->profile_records == (KernelProfileRecord *) NULL))
continue;
OpenCLLog("====================================================");
fprintf(log,"Device: %s\n",device->name);
fprintf(log,"Version: %s\n",device->version);
OpenCLLog("====================================================");
OpenCLLog(" average calls min max");
OpenCLLog(" ------- ----- --- ---");
j=0;
while (device->profile_records[j] != (KernelProfileRecord) NULL)
{
KernelProfileRecord
profile;
profile=device->profile_records[j];
strcpy(indent," ");
CopyMagickString(indent,profile->kernel_name,MagickMin(strlen(
profile->kernel_name),strlen(indent)));
sprintf(buf,"%s %7d %7d %7d %7d",indent,(int) (profile->total/
profile->count),(int) profile->count,(int) profile->min,
(int) profile->max);
OpenCLLog(buf);
j++;
}
OpenCLLog("====================================================");
fwrite("\n\n",sizeof(char),2,log);
}
fclose(log);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ E n q u e u e O p e n C L K e r n e l %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
% events with the images.
%
% The format of the EnqueueOpenCLKernel method is:
%
% MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
% const size_t *global_work_offset,const size_t *global_work_size,
% const size_t *local_work_size,const Image *input_image,
% const Image *output_image,ExceptionInfo *exception)
%
% A description of each parameter follows:
%
% o kernel: the OpenCL kernel.
%
% o work_dim: the number of dimensions used to specify the global work-items
% and work-items in the work-group.
%
% o offset: can be used to specify an array of work_dim unsigned values
% that describe the offset used to calculate the global ID of a
% work-item.
%
% o gsize: points to an array of work_dim unsigned values that describe the
% number of global work-items in work_dim dimensions that will
% execute the kernel function.
%
% o lsize: points to an array of work_dim unsigned values that describe the
% number of work-items that make up a work-group that will execute
% the kernel specified by kernel.
%
% o input_image: the input image of the operation.
%
% o output_image: the output or secondairy image of the operation.
%
% o exception: return any errors or warnings in this structure.
%
*/
static MagickBooleanType RegisterCacheEvent(MagickCLCacheInfo info,
cl_event event)
{
assert(info != (MagickCLCacheInfo) NULL);
assert(event != (cl_event) NULL);
if (openCL_library->clRetainEvent(event) != CL_SUCCESS)
{
openCL_library->clWaitForEvents(1,&event);
return(MagickFalse);
}
LockSemaphoreInfo(info->events_semaphore);
if (info->events == (cl_event *) NULL)
{
info->events=AcquireMagickMemory(sizeof(*info->events));
info->event_count=1;
}
else
info->events=ResizeQuantumMemory(info->events,++info->event_count,
sizeof(*info->events));
if (info->events == (cl_event *) NULL)
ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
info->events[info->event_count-1]=event;
UnlockSemaphoreInfo(info->events_semaphore);
return(MagickTrue);
}
MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
const size_t *lsize,const Image *input_image,const Image *output_image,
MagickBooleanType flush,ExceptionInfo *exception)
{
CacheInfo
*output_info,
*input_info;
cl_event
event,
*events;
cl_int
status;
cl_uint
event_count;
assert(input_image != (const Image *) NULL);
input_info=(CacheInfo *) input_image->cache;
assert(input_info != (CacheInfo *) NULL);
assert(input_info->opencl != (MagickCLCacheInfo) NULL);
output_info=(CacheInfo *) NULL;
if (output_image == (const Image *) NULL)
events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
&event_count);
else
{
output_info=(CacheInfo *) output_image->cache;
assert(output_info != (CacheInfo *) NULL);
assert(output_info->opencl != (MagickCLCacheInfo) NULL);
events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
&event_count);
}
status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
gsize,lsize,event_count,events,&event);
/* This can fail due to memory issues and calling clFinish might help. */
if ((status != CL_SUCCESS) && (event_count > 0))
{
openCL_library->clFinish(queue);
status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
offset,gsize,lsize,event_count,events,&event);
}
events=(cl_event *) RelinquishMagickMemory(events);
if (status != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(input_info->opencl->device,exception,
GetMagickModule(),ResourceLimitWarning,
"clEnqueueNDRangeKernel failed.","'%s'",".");
return(MagickFalse);
}
if (flush != MagickFalse)
openCL_library->clFlush(queue);
if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
{
if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
{
if (output_info != (CacheInfo *) NULL)
(void) RegisterCacheEvent(output_info->opencl,event);
}
}
openCL_library->clReleaseEvent(event);
return(MagickTrue);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ G e t C u r r e n t O p e n C L E n v %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetCurrentOpenCLEnv() returns the current OpenCL env
%
% The format of the GetCurrentOpenCLEnv method is:
%
% MagickCLEnv GetCurrentOpenCLEnv()
%
*/
MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
{
if (default_CLEnv != (MagickCLEnv) NULL)
{
if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
(default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
return((MagickCLEnv) NULL);
else
return(default_CLEnv);
}
if (GetOpenCLCacheDirectory() == (char *) NULL)
return((MagickCLEnv) NULL);
if (openCL_lock == (SemaphoreInfo *) NULL)
ActivateSemaphoreInfo(&openCL_lock);
LockSemaphoreInfo(openCL_lock);
if (default_CLEnv == (MagickCLEnv) NULL)
default_CLEnv=AcquireMagickCLEnv();
UnlockSemaphoreInfo(openCL_lock);
return(default_CLEnv);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e B e n c h m a r k D u r a t i o n %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
% device. The score is determined by the duration of the micro benchmark so
% that means a lower score is better than a higher score.
%
% The format of the GetOpenCLDeviceBenchmarkScore method is:
%
% double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
*/
MagickExport double GetOpenCLDeviceBenchmarkScore(
const MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
return(device->score);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e E n a b l e d %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDeviceEnabled() returns true if the device is enabled.
%
% The format of the GetOpenCLDeviceEnabled method is:
%
% MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
*/
MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
const MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return(MagickFalse);
return(device->enabled);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e N a m e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDeviceName() returns the name of the device.
%
% The format of the GetOpenCLDeviceName method is:
%
% const char *GetOpenCLDeviceName(const MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
*/
MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return((const char *) NULL);
return(device->name);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e V e n d o r N a m e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDeviceVendorName() returns the vendor name of the device.
%
% The format of the GetOpenCLDeviceVendorName method is:
%
% const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
*/
MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return((const char *) NULL);
return(device->vendor_name);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e s %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
% value of length to the number of devices that are available.
%
% The format of the GetOpenCLDevices method is:
%
% const MagickCLDevice *GetOpenCLDevices(size_t *length,
% ExceptionInfo *exception)
%
% A description of each parameter follows:
%
% o length: the number of device.
%
% o exception: return any errors or warnings in this structure.
%
*/
MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
ExceptionInfo *exception)
{
MagickCLEnv
clEnv;
clEnv=GetCurrentOpenCLEnv();
if (clEnv == (MagickCLEnv) NULL)
{
if (length != (size_t *) NULL)
*length=0;
return((MagickCLDevice *) NULL);
}
InitializeOpenCL(clEnv,exception);
if (length != (size_t *) NULL)
*length=clEnv->number_devices;
return(clEnv->devices);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e T y p e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDeviceType() returns the type of the device.
%
% The format of the GetOpenCLDeviceType method is:
%
% MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
*/
MagickExport MagickCLDeviceType GetOpenCLDeviceType(
const MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return(UndefinedCLDeviceType);
if (device->type == CL_DEVICE_TYPE_GPU)
return(GpuCLDeviceType);
if (device->type == CL_DEVICE_TYPE_CPU)
return(CpuCLDeviceType);
return(UndefinedCLDeviceType);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L D e v i c e V e r s i o n %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLDeviceVersion() returns the version of the device.
%
% The format of the GetOpenCLDeviceName method is:
%
% const char *GetOpenCLDeviceVersion(MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
*/
MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return((const char *) NULL);
return(device->version);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L E n a b l e d %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
%
% The format of the GetOpenCLEnabled method is:
%
% MagickBooleanType GetOpenCLEnabled()
%
*/
MagickExport MagickBooleanType GetOpenCLEnabled(void)
{
MagickCLEnv
clEnv;
clEnv=GetCurrentOpenCLEnv();
if (clEnv == (MagickCLEnv) NULL)
return(MagickFalse);
return(clEnv->enabled);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% G e t O p e n C L K e r n e l P r o f i l e R e c o r d s %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% GetOpenCLKernelProfileRecords() returns the profile records for the
% specified device and sets length to the number of profile records.
%
% The format of the GetOpenCLKernelProfileRecords method is:
%
% const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
%
% A description of each parameter follows:
%
% o length: the number of profiles records.
*/
MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
const MagickCLDevice device,size_t *length)
{
if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
(KernelProfileRecord *) NULL))
{
if (length != (size_t *) NULL)
*length=0;
return((const KernelProfileRecord *) NULL);
}
if (length != (size_t *) NULL)
{
*length=0;
LockSemaphoreInfo(device->lock);
while (device->profile_records[*length] != (KernelProfileRecord) NULL)
*length=*length+1;
UnlockSemaphoreInfo(device->lock);
}
return(device->profile_records);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% H a s O p e n C L D e v i c e s %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% HasOpenCLDevices() checks if the OpenCL environment has devices that are
% enabled and compiles the kernel for the device when necessary. False will be
% returned if no enabled devices could be found
%
% The format of the HasOpenCLDevices method is:
%
% MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
% ExceptionInfo exception)
%
% A description of each parameter follows:
%
% o clEnv: the OpenCL environment.
%
% o exception: return any errors or warnings in this structure.
%
*/
static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
ExceptionInfo *exception)
{
char
*accelerateKernelsBuffer,
options[MagickPathExtent];
MagickStatusType
status;
size_t
i;
size_t
signature;
/* Check if there are enabled devices */
for (i = 0; i < clEnv->number_devices; i++)
{
if ((clEnv->devices[i]->enabled != MagickFalse))
break;
}
if (i == clEnv->number_devices)
return(MagickFalse);
/* Check if we need to compile a kernel for one of the devices */
status=MagickTrue;
for (i = 0; i < clEnv->number_devices; i++)
{
if ((clEnv->devices[i]->enabled != MagickFalse) &&
(clEnv->devices[i]->program == (cl_program) NULL))
{
status=MagickFalse;
break;
}
}
if (status != MagickFalse)
return(MagickTrue);
/* Get additional options */
(void) FormatLocaleString(options,MagickPathExtent,CLOptions,
(float)QuantumRange,(float)QuantumScale,(float)CLCharQuantumScale,
(float)MagickEpsilon,(float)MagickPI,(unsigned int)MaxMap,
(unsigned int)MAGICKCORE_QUANTUM_DEPTH);
signature=StringSignature(options);
accelerateKernelsBuffer=(char*) AcquireQuantumMemory(1,
strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
if (accelerateKernelsBuffer == (char*) NULL)
return(MagickFalse);
sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
signature^=StringSignature(accelerateKernelsBuffer);
status=MagickTrue;
for (i = 0; i < clEnv->number_devices; i++)
{
MagickCLDevice
device;
size_t
device_signature;
device=clEnv->devices[i];
if ((device->enabled == MagickFalse) ||
(device->program != (cl_program) NULL))
continue;
LockSemaphoreInfo(device->lock);
if (device->program != (cl_program) NULL)
{
UnlockSemaphoreInfo(device->lock);
continue;
}
device_signature=signature;
device_signature^=StringSignature(device->platform_name);
status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
device_signature,exception);
UnlockSemaphoreInfo(device->lock);
if (status == MagickFalse)
break;
}
accelerateKernelsBuffer=RelinquishMagickMemory(accelerateKernelsBuffer);
return(status);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ I n i t i a l i z e O p e n C L %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% InitializeOpenCL() is used to initialize the OpenCL environment. This method
% makes sure the devices are propertly initialized and benchmarked.
%
% The format of the InitializeOpenCL method is:
%
% MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
%
% A description of each parameter follows:
%
% o exception: return any errors or warnings in this structure.
%
*/
static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
{
char
version[MagickPathExtent];
cl_uint
num;
if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
MagickPathExtent,version,NULL) != CL_SUCCESS)
return(0);
if (strncmp(version,"OpenCL 1.0 ",11) == 0)
return(0);
if (clEnv->library->clGetDeviceIDs(platform,
CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
return(0);
return(num);
}
static inline char *GetOpenCLPlatformString(cl_platform_id platform,
cl_platform_info param_name)
{
char
*value;
size_t
length;
openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
value=AcquireCriticalMemory(length*sizeof(*value));
openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
return(value);
}
static inline char *GetOpenCLDeviceString(cl_device_id device,
cl_device_info param_name)
{
char
*value;
size_t
length;
openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
value=AcquireCriticalMemory(length*sizeof(*value));
openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
return(value);
}
static void LoadOpenCLDevices(MagickCLEnv clEnv)
{
cl_context_properties
properties[3];
cl_device_id
*devices;
cl_int
status;
cl_platform_id
*platforms;
cl_uint
i,
j,
next,
number_devices,
number_platforms;
size_t
length;
number_platforms=0;
if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
return;
if (number_platforms == 0)
return;
platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
sizeof(cl_platform_id));
if (platforms == (cl_platform_id *) NULL)
return;
if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
{
platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
return;
}
for (i = 0; i < number_platforms; i++)
{
char
*platform_name;
number_devices=0;
platform_name=GetOpenCLPlatformString(platforms[i],CL_PLATFORM_NAME);
/* NVIDIA is disabled by default due to reported access violation */
if (strncmp(platform_name,"NVIDIA",6) != 0)
{
number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
clEnv->number_devices+=number_devices;
}
platform_name=(char *) RelinquishMagickMemory(platform_name);
if (number_devices == 0)
platforms[i]=(cl_platform_id) NULL;
}
if (clEnv->number_devices == 0)
{
platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
return;
}
clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
sizeof(MagickCLDevice));
if (clEnv->devices == (MagickCLDevice *) NULL)
{
RelinquishMagickCLDevices(clEnv);
platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
return;
}
(void) memset(clEnv->devices,0,clEnv->number_devices*sizeof(MagickCLDevice));
devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
sizeof(cl_device_id));
if (devices == (cl_device_id *) NULL)
{
platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
RelinquishMagickCLDevices(clEnv);
return;
}
(void) memset(devices,0,clEnv->number_devices*sizeof(cl_device_id));
clEnv->number_contexts=(size_t) number_platforms;
clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
sizeof(cl_context));
if (clEnv->contexts == (cl_context *) NULL)
{
devices=(cl_device_id *) RelinquishMagickMemory(devices);
platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
RelinquishMagickCLDevices(clEnv);
return;
}
(void) memset(clEnv->contexts,0,clEnv->number_contexts*sizeof(cl_context));
next=0;
for (i = 0; i < number_platforms; i++)
{
if (platforms[i] == (cl_platform_id) NULL)
continue;
status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
if (status != CL_SUCCESS)
continue;
properties[0]=CL_CONTEXT_PLATFORM;
properties[1]=(cl_context_properties) platforms[i];
properties[2]=0;
clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
devices,NULL,NULL,&status);
if (status != CL_SUCCESS)
continue;
for (j = 0; j < number_devices; j++,next++)
{
MagickCLDevice
device;
device=AcquireMagickCLDevice();
if (device == (MagickCLDevice) NULL)
break;
device->context=clEnv->contexts[i];
device->deviceID=devices[j];
device->platform_name=GetOpenCLPlatformString(platforms[i],
CL_PLATFORM_NAME);
device->vendor_name=GetOpenCLPlatformString(platforms[i],
CL_PLATFORM_VENDOR);
device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);
device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);
openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
sizeof(cl_uint),&device->max_clock_frequency,NULL);
openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(cl_uint),&device->max_compute_units,NULL);
openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
sizeof(cl_device_type),&device->type,NULL);
openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
sizeof(cl_ulong),&device->local_memory_size,NULL);
clEnv->devices[next]=device;
(void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Found device: %s (%s)",device->name,device->platform_name);
}
}
if (next != clEnv->number_devices)
RelinquishMagickCLDevices(clEnv);
platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
devices=(cl_device_id *) RelinquishMagickMemory(devices);
}
MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
ExceptionInfo *exception)
{
register
size_t i;
LockSemaphoreInfo(clEnv->lock);
if (clEnv->initialized != MagickFalse)
{
UnlockSemaphoreInfo(clEnv->lock);
return(HasOpenCLDevices(clEnv,exception));
}
if (LoadOpenCLLibrary() != MagickFalse)
{
clEnv->library=openCL_library;
LoadOpenCLDevices(clEnv);
if (clEnv->number_devices > 0)
AutoSelectOpenCLDevices(clEnv);
}
clEnv->initialized=MagickTrue;
/* NVIDIA is disabled by default due to reported access violation */
for (i=0; i < (ssize_t) clEnv->number_devices; i++)
{
if (strncmp(clEnv->devices[i]->platform_name,"NVIDIA",6) == 0)
clEnv->devices[i]->enabled=MagickFalse;
}
UnlockSemaphoreInfo(clEnv->lock);
return(HasOpenCLDevices(clEnv,exception));
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% L o a d O p e n C L L i b r a r y %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% LoadOpenCLLibrary() load and binds the OpenCL library.
%
% The format of the LoadOpenCLLibrary method is:
%
% MagickBooleanType LoadOpenCLLibrary(void)
%
*/
void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
{
if ((library == (void *) NULL) || (functionName == (const char *) NULL))
return (void *) NULL;
#ifdef MAGICKCORE_WINDOWS_SUPPORT
return (void *) GetProcAddress((HMODULE)library,functionName);
#else
return (void *) dlsym(library,functionName);
#endif
}
static MagickBooleanType BindOpenCLFunctions()
{
#ifdef MAGICKCORE_OPENCL_MACOSX
#define BIND(X) openCL_library->X= &X;
#else
(void) memset(openCL_library,0,sizeof(MagickLibrary));
#ifdef MAGICKCORE_WINDOWS_SUPPORT
openCL_library->library=(void *)LoadLibraryA("OpenCL.dll");
#else
openCL_library->library=(void *)dlopen("libOpenCL.so",RTLD_NOW);
#endif
#define BIND(X) \
if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
return(MagickFalse);
#endif
if (openCL_library->library == (void*) NULL)
return(MagickFalse);
BIND(clGetPlatformIDs);
BIND(clGetPlatformInfo);
BIND(clGetDeviceIDs);
BIND(clGetDeviceInfo);
BIND(clCreateBuffer);
BIND(clReleaseMemObject);
BIND(clRetainMemObject);
BIND(clCreateContext);
BIND(clReleaseContext);
BIND(clCreateCommandQueue);
BIND(clReleaseCommandQueue);
BIND(clFlush);
BIND(clFinish);
BIND(clCreateProgramWithSource);
BIND(clCreateProgramWithBinary);
BIND(clReleaseProgram);
BIND(clBuildProgram);
BIND(clGetProgramBuildInfo);
BIND(clGetProgramInfo);
BIND(clCreateKernel);
BIND(clReleaseKernel);
BIND(clSetKernelArg);
BIND(clGetKernelInfo);
BIND(clEnqueueReadBuffer);
BIND(clEnqueueMapBuffer);
BIND(clEnqueueUnmapMemObject);
BIND(clEnqueueNDRangeKernel);
BIND(clGetEventInfo);
BIND(clWaitForEvents);
BIND(clReleaseEvent);
BIND(clRetainEvent);
BIND(clSetEventCallback);
BIND(clGetEventProfilingInfo);
return(MagickTrue);
}
static MagickBooleanType LoadOpenCLLibrary(void)
{
openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
if (openCL_library == (MagickLibrary *) NULL)
return(MagickFalse);
if (BindOpenCLFunctions() == MagickFalse)
{
openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
return(MagickFalse);
}
return(MagickTrue);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ O p e n C L T e r m i n u s %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% OpenCLTerminus() destroys the OpenCL component.
%
% The format of the OpenCLTerminus method is:
%
% OpenCLTerminus(void)
%
*/
MagickPrivate void OpenCLTerminus()
{
DumpOpenCLProfileData();
if (cache_directory != (char *) NULL)
cache_directory=DestroyString(cache_directory);
if (cache_directory_lock != (SemaphoreInfo *) NULL)
RelinquishSemaphoreInfo(&cache_directory_lock);
if (default_CLEnv != (MagickCLEnv) NULL)
default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
if (openCL_lock != (SemaphoreInfo *) NULL)
RelinquishSemaphoreInfo(&openCL_lock);
if (openCL_library != (MagickLibrary *) NULL)
{
if (openCL_library->library != (void *) NULL)
(void) lt_dlclose(openCL_library->library);
openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
}
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ O p e n C L T h r o w M a g i c k E x c e p t i o n %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% OpenCLThrowMagickException logs an OpenCL exception as determined by the log
% configuration file. If an error occurs, MagickFalse is returned
% otherwise MagickTrue.
%
% The format of the OpenCLThrowMagickException method is:
%
% MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
% const char *module,const char *function,const size_t line,
% const ExceptionType severity,const char *tag,const char *format,...)
%
% A description of each parameter follows:
%
% o exception: the exception info.
%
% o filename: the source module filename.
%
% o function: the function name.
%
% o line: the line number of the source module.
%
% o severity: Specifies the numeric error category.
%
% o tag: the locale tag.
%
% o format: the output format.
%
*/
MagickPrivate MagickBooleanType OpenCLThrowMagickException(
MagickCLDevice device,ExceptionInfo *exception,const char *module,
const char *function,const size_t line,const ExceptionType severity,
const char *tag,const char *format,...)
{
MagickBooleanType
status;
assert(device != (MagickCLDevice) NULL);
assert(exception != (ExceptionInfo *) NULL);
assert(exception->signature == MagickCoreSignature);
(void) exception;
status=MagickTrue;
if (severity != 0)
{
if (device->type == CL_DEVICE_TYPE_CPU)
{
/* Workaround for Intel OpenCL CPU runtime bug */
/* Turn off OpenCL when a problem is detected! */
if (strncmp(device->platform_name,"Intel",5) == 0)
default_CLEnv->enabled=MagickFalse;
}
}
#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);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ R e c o r d P r o f i l e D a t a %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% RecordProfileData() records profile data.
%
% The format of the RecordProfileData method is:
%
% void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
% cl_event event)
%
% A description of each parameter follows:
%
% o device: the OpenCL device that did the operation.
%
% o event: the event that contains the profiling data.
%
*/
MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
cl_kernel kernel,cl_event event)
{
char
*name;
cl_int
status;
cl_ulong
elapsed,
end,
start;
KernelProfileRecord
profile_record;
size_t
i,
length;
if (device->profile_kernels == MagickFalse)
return(MagickFalse);
status=openCL_library->clWaitForEvents(1,&event);
if (status != CL_SUCCESS)
return(MagickFalse);
status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
&length);
if (status != CL_SUCCESS)
return(MagickTrue);
name=AcquireQuantumMemory(length,sizeof(*name));
if (name == (char *) NULL)
return(MagickTrue);
start=end=elapsed=0;
status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
name,(size_t *) NULL);
status|=openCL_library->clGetEventProfilingInfo(event,
CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
status|=openCL_library->clGetEventProfilingInfo(event,
CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
if (status != CL_SUCCESS)
{
name=DestroyString(name);
return(MagickTrue);
}
start/=1000; /* usecs */
end/=1000;
elapsed=end-start;
LockSemaphoreInfo(device->lock);
i=0;
profile_record=(KernelProfileRecord) NULL;
if (device->profile_records != (KernelProfileRecord *) NULL)
{
while (device->profile_records[i] != (KernelProfileRecord) NULL)
{
if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
{
profile_record=device->profile_records[i];
break;
}
i++;
}
}
if (profile_record != (KernelProfileRecord) NULL)
name=DestroyString(name);
else
{
profile_record=AcquireCriticalMemory(sizeof(*profile_record));
(void) memset(profile_record,0,sizeof(*profile_record));
profile_record->kernel_name=name;
device->profile_records=ResizeQuantumMemory(device->profile_records,(i+2),
sizeof(*device->profile_records));
if (device->profile_records == (KernelProfileRecord *) NULL)
ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
device->profile_records[i]=profile_record;
device->profile_records[i+1]=(KernelProfileRecord) NULL;
}
if ((elapsed < profile_record->min) || (profile_record->count == 0))
profile_record->min=elapsed;
if (elapsed > profile_record->max)
profile_record->max=elapsed;
profile_record->total+=elapsed;
profile_record->count+=1;
UnlockSemaphoreInfo(device->lock);
return(MagickTrue);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ R e l e a s e O p e n C L C o m m a n d Q u e u e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% ReleaseOpenCLCommandQueue() releases the OpenCL command queue
%
% The format of the ReleaseOpenCLCommandQueue method is:
%
% void ReleaseOpenCLCommandQueue(MagickCLDevice device,
% cl_command_queue queue)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
%
% o queue: the OpenCL queue to be released.
*/
MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
cl_command_queue queue)
{
if (queue == (cl_command_queue) NULL)
return;
assert(device != (MagickCLDevice) NULL);
LockSemaphoreInfo(device->lock);
if ((device->profile_kernels != MagickFalse) ||
(device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
{
UnlockSemaphoreInfo(device->lock);
openCL_library->clFinish(queue);
(void) openCL_library->clReleaseCommandQueue(queue);
}
else
{
openCL_library->clFlush(queue);
device->command_queues[++device->command_queues_index]=queue;
UnlockSemaphoreInfo(device->lock);
}
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ R e l e a s e M a g i c k C L D e v i c e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% ReleaseOpenCLDevice() returns the OpenCL device to the environment
%
% The format of the ReleaseOpenCLDevice method is:
%
% void ReleaseOpenCLDevice(MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device to be released.
%
*/
MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
{
assert(device != (MagickCLDevice) NULL);
LockSemaphoreInfo(openCL_lock);
device->requested--;
UnlockSemaphoreInfo(openCL_lock);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ R e l i n q u i s h M a g i c k C L C a c h e I n f o %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% RelinquishMagickCLCacheInfo() frees memory acquired with
% AcquireMagickCLCacheInfo()
%
% The format of the RelinquishMagickCLCacheInfo method is:
%
% MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
% const MagickBooleanType relinquish_pixels)
%
% A description of each parameter follows:
%
% o info: the OpenCL cache info.
%
% o relinquish_pixels: the pixels will be relinquish when set to true.
%
*/
static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
cl_event magick_unused(event),
cl_int magick_unused(event_command_exec_status),void *user_data)
{
MagickCLCacheInfo
info;
Quantum
*pixels;
ssize_t
i;
magick_unreferenced(event);
magick_unreferenced(event_command_exec_status);
info=(MagickCLCacheInfo) user_data;
for (i=(ssize_t)info->event_count-1; i >= 0; i--)
{
cl_int
event_status;
cl_uint
status;
status=openCL_library->clGetEventInfo(info->events[i],
CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
NULL);
if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
{
openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
&DestroyMagickCLCacheInfoAndPixels,info);
return;
}
}
pixels=info->pixels;
RelinquishMagickResource(MemoryResource,info->length);
DestroyMagickCLCacheInfo(info);
(void) RelinquishAlignedMemory(pixels);
}
MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
{
if (info == (MagickCLCacheInfo) NULL)
return((MagickCLCacheInfo) NULL);
if (relinquish_pixels != MagickFalse)
DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
else
DestroyMagickCLCacheInfo(info);
return((MagickCLCacheInfo) NULL);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% R e l i n q u i s h M a g i c k C L D e v i c e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% RelinquishMagickCLDevice() releases the OpenCL device
%
% The format of the RelinquishMagickCLDevice method is:
%
% MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
%
% A description of each parameter follows:
%
% o device: the OpenCL device to be released.
%
*/
static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
{
if (device == (MagickCLDevice) NULL)
return((MagickCLDevice) NULL);
device->platform_name=RelinquishMagickMemory(device->platform_name);
device->vendor_name=RelinquishMagickMemory(device->vendor_name);
device->name=RelinquishMagickMemory(device->name);
device->version=RelinquishMagickMemory(device->version);
if (device->program != (cl_program) NULL)
(void) openCL_library->clReleaseProgram(device->program);
while (device->command_queues_index >= 0)
(void) openCL_library->clReleaseCommandQueue(
device->command_queues[device->command_queues_index--]);
RelinquishSemaphoreInfo(&device->lock);
return((MagickCLDevice) RelinquishMagickMemory(device));
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% R e l i n q u i s h M a g i c k C L E n v %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% RelinquishMagickCLEnv() releases the OpenCL environment
%
% The format of the RelinquishMagickCLEnv method is:
%
% MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
%
% A description of each parameter follows:
%
% o clEnv: the OpenCL environment to be released.
%
*/
static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
{
if (clEnv == (MagickCLEnv) NULL)
return((MagickCLEnv) NULL);
RelinquishSemaphoreInfo(&clEnv->lock);
RelinquishMagickCLDevices(clEnv);
if (clEnv->contexts != (cl_context *) NULL)
{
ssize_t
i;
for (i=0; i < clEnv->number_contexts; i++)
if (clEnv->contexts[i] != (cl_context) NULL)
(void) openCL_library->clReleaseContext(clEnv->contexts[i]);
clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
}
return((MagickCLEnv) RelinquishMagickMemory(clEnv));
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
+ R e q u e s t O p e n C L D e v i c e %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% RequestOpenCLDevice() returns one of the enabled OpenCL devices.
%
% The format of the RequestOpenCLDevice method is:
%
% MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
%
% A description of each parameter follows:
%
% o clEnv: the OpenCL environment.
*/
MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
{
MagickCLDevice
device;
double
score,
best_score;
size_t
i;
if (clEnv == (MagickCLEnv) NULL)
return((MagickCLDevice) NULL);
if (clEnv->number_devices == 1)
{
if (clEnv->devices[0]->enabled)
return(clEnv->devices[0]);
else
return((MagickCLDevice) NULL);
}
device=(MagickCLDevice) NULL;
best_score=0.0;
LockSemaphoreInfo(openCL_lock);
for (i = 0; i < clEnv->number_devices; i++)
{
if (clEnv->devices[i]->enabled == MagickFalse)
continue;
score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
clEnv->devices[i]->requested);
if ((device == (MagickCLDevice) NULL) || (score < best_score))
{
device=clEnv->devices[i];
best_score=score;
}
}
if (device != (MagickCLDevice)NULL)
device->requested++;
UnlockSemaphoreInfo(openCL_lock);
return(device);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% S e t O p e n C L D e v i c e E n a b l e d %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
%
% The format of the SetOpenCLDeviceEnabled method is:
%
% void SetOpenCLDeviceEnabled(MagickCLDevice device,
% MagickBooleanType value)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
%
% o value: determines if the device should be enabled or disabled.
*/
MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
const MagickBooleanType value)
{
if (device == (MagickCLDevice) NULL)
return;
device->enabled=value;
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% S e t O p e n C L K e r n e l P r o f i l e E n a b l e d %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
% kernel profiling of a device.
%
% The format of the SetOpenCLKernelProfileEnabled method is:
%
% void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
% MagickBooleanType value)
%
% A description of each parameter follows:
%
% o device: the OpenCL device.
%
% o value: determines if kernel profiling for the device should be enabled
% or disabled.
*/
MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
const MagickBooleanType value)
{
if (device == (MagickCLDevice) NULL)
return;
device->profile_kernels=value;
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% S e t O p e n C L E n a b l e d %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
%
% The format of the SetOpenCLEnabled method is:
%
% void SetOpenCLEnabled(MagickBooleanType)
%
% A description of each parameter follows:
%
% o value: specify true to enable OpenCL acceleration
*/
MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
{
MagickCLEnv
clEnv;
clEnv=GetCurrentOpenCLEnv();
if (clEnv == (MagickCLEnv) NULL)
return(MagickFalse);
clEnv->enabled=value;
return(clEnv->enabled);
}
#else
MagickExport double GetOpenCLDeviceBenchmarkScore(
const MagickCLDevice magick_unused(device))
{
magick_unreferenced(device);
return(0.0);
}
MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
const MagickCLDevice magick_unused(device))
{
magick_unreferenced(device);
return(MagickFalse);
}
MagickExport const char *GetOpenCLDeviceName(
const MagickCLDevice magick_unused(device))
{
magick_unreferenced(device);
return((const char *) NULL);
}
MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
ExceptionInfo *magick_unused(exception))
{
magick_unreferenced(exception);
if (length != (size_t *) NULL)
*length=0;
return((MagickCLDevice *) NULL);
}
MagickExport MagickCLDeviceType GetOpenCLDeviceType(
const MagickCLDevice magick_unused(device))
{
magick_unreferenced(device);
return(UndefinedCLDeviceType);
}
MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
const MagickCLDevice magick_unused(device),size_t *length)
{
magick_unreferenced(device);
if (length != (size_t *) NULL)
*length=0;
return((const KernelProfileRecord *) NULL);
}
MagickExport const char *GetOpenCLDeviceVersion(
const MagickCLDevice magick_unused(device))
{
magick_unreferenced(device);
return((const char *) NULL);
}
MagickExport MagickBooleanType GetOpenCLEnabled(void)
{
return(MagickFalse);
}
MagickExport void SetOpenCLDeviceEnabled(
MagickCLDevice magick_unused(device),
const MagickBooleanType magick_unused(value))
{
magick_unreferenced(device);
magick_unreferenced(value);
}
MagickExport MagickBooleanType SetOpenCLEnabled(
const MagickBooleanType magick_unused(value))
{
magick_unreferenced(value);
return(MagickFalse);
}
MagickExport void SetOpenCLKernelProfileEnabled(
MagickCLDevice magick_unused(device),
const MagickBooleanType magick_unused(value))
{
magick_unreferenced(device);
magick_unreferenced(value);
}
#endif
Reference in new issue View Git Blame Copy Permalink