averello/ImageMagick-mirror
1
0
Fork 0
mirror of https://github.com/ImageMagick/ImageMagick.git synced 2026-06-12 15:37:14 +02:00
Code Issues Packages Projects Releases Wiki Activity

Complete rewrite of the OpenCL API.

Browse Source
This commit is contained in:
dirk
2016-04-15 09:24:26 +02:00
parent 4b63e3e8d1
commit 5a2575faca
9 changed files with 3231 additions and 4284 deletions
Download Patch File Download Diff File Expand all files Collapse all files
Magick++/lib/Functions.cpp
+4 -16
View File
@@ -28,28 +28,16 @@ MagickPPExport void Magick::CloneString(char **destination_,
MagickPPExport void Magick::DisableOpenCL(void)
{
GetPPException;
MagickCore::InitImageMagickOpenCL(MagickCore::MAGICK_OPENCL_OFF,NULL,NULL,
exceptionInfo);
ThrowPPException(false);
MagickCore::SetOpenCLEnabled(MagickFalse);
}
MagickPPExport bool Magick::EnableOpenCL(const bool useCache_)
MagickPPExport bool Magick::EnableOpenCL(void)
{
bool
status;
GetPPException;
if (useCache_)
status=MagickCore::InitImageMagickOpenCL(
MagickCore::MAGICK_OPENCL_DEVICE_SELECT_AUTO,NULL,NULL,exceptionInfo) ==
MagickTrue;
else
status=MagickCore::InitImageMagickOpenCL(
MagickCore::MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE,NULL,NULL,
exceptionInfo) == MagickTrue;
ThrowPPException(false);
return(status);
status=MagickCore::SetOpenCLEnabled(MagickTrue) != MagickFalse;
return(status);
}
MagickPPExport void Magick::InitializeMagick(const char *path_)
Magick++/lib/Magick++/Functions.h
+1 -1
View File
@@ -23,7 +23,7 @@ namespace Magick
MagickPPExport void DisableOpenCL(void);
// Enable OpenCL acceleration (only works when build with OpenCL support)
MagickPPExport bool EnableOpenCL(const bool useCache_=true);
MagickPPExport bool EnableOpenCL(void);
// C library initialization routine
MagickPPExport void InitializeMagick(const char *path_);
MagickCore/accelerate-private.h
+24 -96
View File
@@ -49,7 +49,7 @@ typedef struct _FloatPixelPacket
black;
} FloatPixelPacket;
const char* accelerateKernels =
const char *accelerateKernels =
/*
Define declarations.
@@ -2387,79 +2387,6 @@ OPENCL_ENDIF()
}
)
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
% %
% %
% R a n d o m %
% %
% %
% %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
*/
STRINGIFY(
inline float GetPseudoRandomValue(uint4* seed, const float normalizeRand) {
uint4 s = *seed;
do {
unsigned int alpha = (unsigned int)(s.y ^ (s.y << 11));
s.y = s.z;
s.z = s.w;
s.w = s.x;
s.x = (s.x ^ (s.x >> 19)) ^ (alpha ^ (alpha >> 8));
} while (s.x == ~0UL);
*seed = s;
return (normalizeRand*s.x);
}
__kernel void RandomNumberGenerator(__global uint* seeds, const float normalizeRand
, __global float* randomNumbers, const uint init
, const uint numRandomNumbers) {
unsigned int id = get_global_id(0);
unsigned int seed[4];
if (init != 0) {
seed[0] = seeds[id * 4];
seed[1] = 0x50a7f451;
seed[2] = 0x5365417e;
seed[3] = 0xc3a4171a;
}
else {
seed[0] = seeds[id * 4];
seed[1] = seeds[id * 4 + 1];
seed[2] = seeds[id * 4 + 2];
seed[3] = seeds[id * 4 + 3];
}
unsigned int numRandomNumbersPerItem = (numRandomNumbers + get_global_size(0) - 1) / get_global_size(0);
for (unsigned int i = 0; i < numRandomNumbersPerItem; i++) {
do
{
unsigned int alpha = (unsigned int)(seed[1] ^ (seed[1] << 11));
seed[1] = seed[2];
seed[2] = seed[3];
seed[3] = seed[0];
seed[0] = (seed[0] ^ (seed[0] >> 19)) ^ (alpha ^ (alpha >> 8));
} while (seed[0] == ~0UL);
unsigned int pos = (get_group_id(0)*get_local_size(0)*numRandomNumbersPerItem)
+ get_local_size(0) * i + get_local_id(0);
if (pos >= numRandomNumbers)
break;
randomNumbers[pos] = normalizeRand*seed[0];
}
/* save the seeds for the time*/
seeds[id * 4] = seed[0];
seeds[id * 4 + 1] = seed[1];
seeds[id * 4 + 2] = seed[2];
seeds[id * 4 + 3] = seed[3];
}
)
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
@@ -2641,7 +2568,7 @@ STRINGIFY(
)
;
const char* accelerateKernels2 =
const char *accelerateKernels2 =
STRINGIFY(
@@ -2740,17 +2667,17 @@ STRINGIFY(
float4 cp = (float4) 0;
__local float *p = inputImageCache + (cacheIndex*number_channels);
cp.x = *(p);
__local CLQuantum *p = inputImageCache + (cacheIndex*number_channels);
cp.x = (float) *(p);
if (number_channels > 2)
{
cp.y = *(p + 1);
cp.z = *(p + 2);
cp.y = (float) *(p + 1);
cp.z = (float) *(p + 2);
}
if ((number_channels == 4) || (number_channels == 2))
{
cp.w = *(p + number_channels - 1);
cp.w = (float) *(p + number_channels - 1);
float alpha = weight * QuantumScale * cp.w;
@@ -2793,17 +2720,19 @@ STRINGIFY(
if (itemID < actualNumPixelInThisChunk)
{
float density = densityCache[itemID];
float gamma = gammaCache[itemID];
float4 filteredPixel = outputPixelCache[itemID];
if ((density != 0.0f) && (density != 1.0f))
{
density = PerceptibleReciprocal(density);
filteredPixel *= (float4) density;
gamma *= density;
}
if ((number_channels == 4) || (number_channels == 2))
{
float gamma = PerceptibleReciprocal(gammaCache[itemID]*density);
gamma = PerceptibleReciprocal(gamma);
filteredPixel.x *= gamma;
filteredPixel.y *= gamma;
filteredPixel.z *= gamma;
@@ -2899,17 +2828,17 @@ STRINGIFY(
float4 cp = (float4)0.0f;
__local float *p = inputImageCache + cacheIndex;
cp.x = *(p);
__local CLQuantum *p = inputImageCache + cacheIndex;
cp.x = (float) *(p);
if (number_channels > 2)
{
cp.y = *(p + rangeLength);
cp.z = *(p + (rangeLength * 2));
cp.y = (float) *(p + rangeLength);
cp.z = (float) *(p + (rangeLength * 2));
}
if ((number_channels == 4) || (number_channels == 2))
{
cp.w = *(p + (rangeLength * (number_channels - 1)));
cp.w = (float) *(p + (rangeLength * (number_channels - 1)));
float alpha = weight * QuantumScale * cp.w;
@@ -2952,17 +2881,19 @@ STRINGIFY(
if (itemID < actualNumPixelInThisChunk)
{
float density = densityCache[itemID];
float gamma = gammaCache[itemID];
float4 filteredPixel = outputPixelCache[itemID];
if ((density != 0.0f) && (density != 1.0f))
{
density = PerceptibleReciprocal(density);
filteredPixel *= (float4) density;
gamma *= density;
}
if ((number_channels == 4) || (number_channels == 2))
{
float gamma = PerceptibleReciprocal(gammaCache[itemID]*density);
gamma = PerceptibleReciprocal(gamma);
filteredPixel.x *= gamma;
filteredPixel.y *= gamma;
filteredPixel.z *= gamma;
@@ -3138,8 +3069,7 @@ STRINGIFY(
__kernel void UnsharpMask(const __global CLQuantum *image,const unsigned int number_channels,
const ChannelType channel,__constant float *filter,const unsigned int width,
const unsigned int columns,const unsigned int rows,__local float4 *pixels,
const float gain,const float threshold, const unsigned int justBlur,
__global CLQuantum *filteredImage)
const float gain,const float threshold,__global CLQuantum *filteredImage)
{
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
@@ -3199,15 +3129,13 @@ STRINGIFY(
++i;
}
if (justBlur == 0) { // apply sharpening
float4 srcPixel = ReadFloat4(image, number_channels, columns, x, y, channel);
float4 diff = srcPixel - value;
float4 srcPixel = ReadFloat4(image, number_channels, columns, x, y, channel);
float4 diff = srcPixel - value;
float quantumThreshold = QuantumRange*threshold;
float quantumThreshold = QuantumRange*threshold;
int4 mask = isless(fabs(2.0f * diff), (float4)quantumThreshold);
value = select(srcPixel + diff * gain, srcPixel, mask);
}
int4 mask = isless(fabs(2.0f * diff), (float4)quantumThreshold);
value = select(srcPixel + diff * gain, srcPixel, mask);
if ((x < columns) && (y < rows))
WriteFloat4(filteredImage, number_channels, columns, x, y, channel, value);
MagickCore/accelerate.c
+730 -1065
View File
File diff suppressed because it is too large Load Diff
MagickCore/accelerate.h
+1 -2
View File
@@ -56,8 +56,7 @@ extern MagickExport MagickBooleanType
AccelerateGrayscaleImage(Image *,const PixelIntensityMethod,
ExceptionInfo *),
AccelerateModulateImage(Image *,const double,const double,const double,
const ColorspaceType, ExceptionInfo*),
AccelerateRandomImage(Image*,ExceptionInfo*);
const ColorspaceType, ExceptionInfo*);
#if defined(__cplusplus) || defined(c_plusplus)
}
MagickCore/magick.c
+2
View File
@@ -1528,7 +1528,9 @@ MagickExport void MagickCoreTerminus(void)
AnnotateComponentTerminus();
MimeComponentTerminus();
TypeComponentTerminus();
#if defined(MAGICKCORE_OPENCL_SUPPORT)
OpenCLTerminus();
#endif
ColorComponentTerminus();
#if defined(MAGICKCORE_WINDOWS_SUPPORT)
NTWindowsTerminus();
MagickCore/opencl-private.h
+315 -325
View File
@@ -23,11 +23,41 @@ Include declarations.
*/
#include "MagickCore/studio.h"
#include "MagickCore/opencl.h"
#include "MagickCore/thread_.h"
#if defined(__cplusplus) || defined(c_plusplus)
extern "C" {
#endif
typedef enum
{
AddNoiseKernel,
BlurColumnKernel,
BlurRowKernel,
CompositeKernel,
ContrastKernel,
ContrastStretchKernel,
ConvolveKernel,
ConvolveOptimizedKernel,
ComputeFunctionKernel,
EqualizeKernel,
GrayScaleKernel,
HistogramKernel,
HullPass1Kernel,
HullPass2Kernel,
LocalContrastBlurApplyColumnKernel,
LocalContrastBlurRowKernel,
ModulateKernel,
MotionBlurKernel,
ResizeHorizontalKernel,
ResizeVerticalKernel,
RotationalBlurKernel,
UnsharpMaskKernel,
UnsharpMaskBlurColumnKernel,
WaveletDenoiseKernel,
KERNEL_COUNT
} ProfiledKernels;
#if !defined(MAGICKCORE_OPENCL_SUPPORT)
typedef void* cl_context;
typedef void* cl_command_queue;
@@ -36,279 +66,294 @@ extern "C" {
typedef void* cl_kernel;
typedef void* cl_mem;
typedef void* cl_platform_id;
typedef struct { unsigned char t[8]; } cl_device_type; /* 64-bit */
typedef void* cl_device_type;
#else
#define MAX_COMMAND_QUEUES 16
/*
*
* function pointer typedefs
*
*/
Define declarations.
*/
#define MAGICKCORE_OPENCL_UNDEFINED_SCORE -1.0
#define MAGICKCORE_OPENCL_PROFILE_KERNELS 0
#define MAGICKCORE_OPENCL_COMMAND_QUEUES 16
#if MAGICKCORE_OPENCL_PROFILE_KERNELS
typedef struct
{
cl_ulong min;
cl_ulong max;
cl_ulong total;
cl_ulong count;
} KernelProfileRecord;
#endif
/* Platform APIs */
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetPlatformIDs)(
cl_uint num_entries,
cl_platform_id * platforms,
cl_uint * num_platforms) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetPlatformIDs)(cl_uint num_entries,
cl_platform_id *platforms,cl_uint *num_platforms) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetPlatformInfo)(
cl_platform_id platform,
cl_platform_info param_name,
size_t param_value_size,
void * param_value,
size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
/* Device APIs */
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetDeviceIDs)(
cl_platform_id platform,
cl_device_type device_type,
cl_uint num_entries,
cl_device_id * devices,
cl_uint * num_devices) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetDeviceInfo)(
cl_device_id device,
cl_device_info param_name,
size_t param_value_size,
void * param_value,
size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
/* Context APIs */
typedef CL_API_ENTRY cl_context (CL_API_CALL *MAGICKpfn_clCreateContext)(
const cl_context_properties * properties,
cl_uint num_devices,
const cl_device_id * devices,
void (CL_CALLBACK *pfn_notify)(const char *, const void *, size_t, void *),
void * user_data,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseContext)(
cl_context context) CL_API_SUFFIX__VERSION_1_0;
/* Command Queue APIs */
typedef CL_API_ENTRY cl_command_queue (CL_API_CALL *MAGICKpfn_clCreateCommandQueue)(
cl_context context,
cl_device_id device,
cl_command_queue_properties properties,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseCommandQueue)(
cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
/* Memory Object APIs */
typedef CL_API_ENTRY cl_mem (CL_API_CALL *MAGICKpfn_clCreateBuffer)(
cl_context context,
cl_mem_flags flags,
size_t size,
void * host_ptr,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseMemObject)(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0;
/* Program Object APIs */
typedef CL_API_ENTRY cl_program (CL_API_CALL *MAGICKpfn_clCreateProgramWithSource)(
cl_context context,
cl_uint count,
const char ** strings,
const size_t * lengths,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_program (CL_API_CALL *MAGICKpfn_clCreateProgramWithBinary)(
cl_context context,
cl_uint num_devices,
const cl_device_id * device_list,
const size_t * lengths,
const unsigned char ** binaries,
cl_int * binary_status,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseProgram)(cl_program program) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clBuildProgram)(
cl_program program,
cl_uint num_devices,
const cl_device_id * device_list,
const char * options,
void (CL_CALLBACK *pfn_notify)(cl_program program, void * user_data),
void * user_data) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetProgramInfo)(
cl_program program,
cl_program_info param_name,
size_t param_value_size,
void * param_value,
size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetProgramBuildInfo)(
cl_program program,
cl_device_id device,
cl_program_build_info param_name,
size_t param_value_size,
void * param_value,
size_t * param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
/* Kernel Object APIs */
typedef CL_API_ENTRY cl_kernel (CL_API_CALL *MAGICKpfn_clCreateKernel)(
cl_program program,
const char * kernel_name,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseKernel)(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clSetKernelArg)(
cl_kernel kernel,
cl_uint arg_index,
size_t arg_size,
const void * arg_value) CL_API_SUFFIX__VERSION_1_0;
/* Flush and Finish APIs */
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clFlush)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clFinish)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
/* Enqueued Commands APIs */
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueReadBuffer)(
cl_command_queue command_queue,
cl_mem buffer,
cl_bool blocking_read,
size_t offset,
size_t cb,
void * ptr,
cl_uint num_events_in_wait_list,
const cl_event * event_wait_list,
cl_event * event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueWriteBuffer)(
cl_command_queue command_queue,
cl_mem buffer,
cl_bool blocking_write,
size_t offset,
size_t cb,
const void * ptr,
cl_uint num_events_in_wait_list,
const cl_event * event_wait_list,
cl_event * event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY void * (CL_API_CALL *MAGICKpfn_clEnqueueMapBuffer)(
cl_command_queue command_queue,
cl_mem buffer,
cl_bool blocking_map,
cl_map_flags map_flags,
size_t offset,
size_t cb,
cl_uint num_events_in_wait_list,
const cl_event * event_wait_list,
cl_event * event,
cl_int * errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueUnmapMemObject)(
cl_command_queue command_queue,
cl_mem memobj,
void * mapped_ptr,
cl_uint num_events_in_wait_list,
const cl_event * event_wait_list,
cl_event * event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueNDRangeKernel)(
cl_command_queue command_queue,
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,
cl_uint num_events_in_wait_list,
const cl_event * event_wait_list,
cl_event * event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)(
cl_event event,
cl_profiling_info param_name,
size_t param_value_size,
void *param_value,
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetPlatformInfo)(cl_platform_id platform,
cl_platform_info param_name,size_t param_value_size,void *param_value,
size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clWaitForEvents)(
cl_uint num_events,
/* Device APIs */
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetDeviceIDs)(cl_platform_id platform,
cl_device_type device_type,cl_uint num_entries,cl_device_id *devices,
cl_uint *num_devices) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetDeviceInfo)(cl_device_id device,
cl_device_info param_name,size_t param_value_size,void *param_value,
size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
/* Context APIs */
typedef CL_API_ENTRY cl_context
(CL_API_CALL *MAGICKpfn_clCreateContext)(
const cl_context_properties *properties,cl_uint num_devices,
const cl_device_id *devices,void (CL_CALLBACK *pfn_notify)(const char *,
const void *,size_t,void *),void *user_data,cl_int *errcode_ret)
CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clReleaseContext)(cl_context context)
CL_API_SUFFIX__VERSION_1_0;
/* Command Queue APIs */
typedef CL_API_ENTRY cl_command_queue
(CL_API_CALL *MAGICKpfn_clCreateCommandQueue)(cl_context context,
cl_device_id device,cl_command_queue_properties properties,
cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clReleaseCommandQueue)(
cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
/* Memory Object APIs */
typedef CL_API_ENTRY cl_mem
(CL_API_CALL *MAGICKpfn_clCreateBuffer)(cl_context context,
cl_mem_flags flags,size_t size,void *host_ptr,cl_int *errcode_ret)
CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clReleaseMemObject)(cl_mem memobj)
CL_API_SUFFIX__VERSION_1_0;
/* Program Object APIs */
typedef CL_API_ENTRY cl_program
(CL_API_CALL *MAGICKpfn_clCreateProgramWithSource)(cl_context context,
cl_uint count,const char **strings,const size_t *lengths,
cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_program
(CL_API_CALL *MAGICKpfn_clCreateProgramWithBinary)(cl_context context,
cl_uint num_devices,const cl_device_id *device_list,const size_t *lengths,
const unsigned char **binaries,cl_int *binary_status,cl_int *errcode_ret)
CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clReleaseProgram)(cl_program program)
CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clBuildProgram)(cl_program program,
cl_uint num_devices,const cl_device_id *device_list,const char *options,
void (CL_CALLBACK *pfn_notify)(cl_program program,void * user_data),
void *user_data) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetProgramBuildInfo)(cl_program program,
cl_device_id device,cl_program_build_info param_name,size_t param_value_size,
void *param_value,size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetProgramInfo)(cl_program program,
cl_program_info param_name,size_t param_value_size,void *param_value,
size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
/* Kernel Object APIs */
typedef CL_API_ENTRY cl_kernel
(CL_API_CALL *MAGICKpfn_clCreateKernel)(cl_program program,
const char *kernel_name,cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clReleaseKernel)(cl_kernel kernel)
CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clSetKernelArg)(cl_kernel kernel,cl_uint arg_index,
size_t arg_size,const void * arg_value) CL_API_SUFFIX__VERSION_1_0;
/* Enqueued Commands APIs */
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clEnqueueReadBuffer)(cl_command_queue command_queue,
cl_mem buffer,cl_bool blocking_read,size_t offset,size_t cb,void *ptr,
cl_uint num_events_in_wait_list,const cl_event *event_wait_list,
cl_event *event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY void
*(CL_API_CALL *MAGICKpfn_clEnqueueMapBuffer)(cl_command_queue command_queue,
cl_mem buffer,cl_bool blocking_map,cl_map_flags map_flags,size_t offset,
size_t cb,cl_uint num_events_in_wait_list,const cl_event *event_wait_list,
cl_event *event,cl_int *errcode_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clEnqueueUnmapMemObject)(
cl_command_queue command_queue,cl_mem memobj,void *mapped_ptr,
cl_uint num_events_in_wait_list,const cl_event *event_wait_list,
cl_event *event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clEnqueueNDRangeKernel)(
cl_command_queue command_queue,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,cl_uint num_events_in_wait_list,
const cl_event * event_wait_list,cl_event *event)
CL_API_SUFFIX__VERSION_1_0;
/* Profiling APIs */
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)(cl_event event,
cl_profiling_info param_name,size_t param_value_size,void *param_value,
size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clWaitForEvents)(cl_uint num_events,
const cl_event *event_list) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clReleaseEvent)(
cl_event event) CL_API_SUFFIX__VERSION_1_0;
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clReleaseEvent)(cl_event event)
CL_API_SUFFIX__VERSION_1_0;
/*
*
* vendor dispatch table structure
*
* note that the types in the structure KHRicdVendorDispatch mirror the function
* names listed in the string table khrIcdVendorDispatchFunctionNames
*
*/
/* Finish APIs, only here for GetAndLockRandSeedBuffer */
typedef CL_API_ENTRY cl_int
(CL_API_CALL *MAGICKpfn_clFinish)(cl_command_queue command_queue)
CL_API_SUFFIX__VERSION_1_0;
typedef struct MagickLibraryRec MagickLibrary;
struct MagickLibraryRec
{
void * base;
MAGICKpfn_clGetPlatformIDs clGetPlatformIDs;
MAGICKpfn_clGetPlatformInfo clGetPlatformInfo;
MAGICKpfn_clGetPlatformIDs clGetPlatformIDs;
MAGICKpfn_clGetPlatformInfo clGetPlatformInfo;
MAGICKpfn_clGetDeviceIDs clGetDeviceIDs;
MAGICKpfn_clGetDeviceInfo clGetDeviceInfo;
MAGICKpfn_clCreateContext clCreateContext;
MAGICKpfn_clCreateCommandQueue clCreateCommandQueue;
MAGICKpfn_clReleaseCommandQueue clReleaseCommandQueue;
MAGICKpfn_clCreateBuffer clCreateBuffer;
MAGICKpfn_clReleaseMemObject clReleaseMemObject;
MAGICKpfn_clCreateProgramWithSource clCreateProgramWithSource;
MAGICKpfn_clCreateProgramWithBinary clCreateProgramWithBinary;
MAGICKpfn_clReleaseProgram clReleaseProgram;
MAGICKpfn_clBuildProgram clBuildProgram;
MAGICKpfn_clGetProgramInfo clGetProgramInfo;
MAGICKpfn_clGetProgramBuildInfo clGetProgramBuildInfo;
MAGICKpfn_clCreateKernel clCreateKernel;
MAGICKpfn_clReleaseKernel clReleaseKernel;
MAGICKpfn_clSetKernelArg clSetKernelArg;
MAGICKpfn_clFlush clFlush;
MAGICKpfn_clFinish clFinish;
MAGICKpfn_clEnqueueReadBuffer clEnqueueReadBuffer;
MAGICKpfn_clEnqueueWriteBuffer clEnqueueWriteBuffer;
MAGICKpfn_clEnqueueMapBuffer clEnqueueMapBuffer;
MAGICKpfn_clEnqueueUnmapMemObject clEnqueueUnmapMemObject;
MAGICKpfn_clEnqueueNDRangeKernel clEnqueueNDRangeKernel;
MAGICKpfn_clGetEventProfilingInfo clGetEventProfilingInfo;
MAGICKpfn_clWaitForEvents clWaitForEvents;
MAGICKpfn_clReleaseEvent clReleaseEvent;
MAGICKpfn_clGetDeviceIDs clGetDeviceIDs;
MAGICKpfn_clGetDeviceInfo clGetDeviceInfo;
MAGICKpfn_clCreateContext clCreateContext;
MAGICKpfn_clReleaseContext clReleaseContext;
MAGICKpfn_clCreateCommandQueue clCreateCommandQueue;
MAGICKpfn_clReleaseCommandQueue clReleaseCommandQueue;
MAGICKpfn_clCreateBuffer clCreateBuffer;
MAGICKpfn_clReleaseMemObject clReleaseMemObject;
MAGICKpfn_clCreateProgramWithSource clCreateProgramWithSource;
MAGICKpfn_clCreateProgramWithBinary clCreateProgramWithBinary;
MAGICKpfn_clReleaseProgram clReleaseProgram;
MAGICKpfn_clBuildProgram clBuildProgram;
MAGICKpfn_clGetProgramBuildInfo clGetProgramBuildInfo;
MAGICKpfn_clGetProgramInfo clGetProgramInfo;
MAGICKpfn_clCreateKernel clCreateKernel;
MAGICKpfn_clReleaseKernel clReleaseKernel;
MAGICKpfn_clSetKernelArg clSetKernelArg;
MAGICKpfn_clEnqueueReadBuffer clEnqueueReadBuffer;
MAGICKpfn_clEnqueueMapBuffer clEnqueueMapBuffer;
MAGICKpfn_clEnqueueUnmapMemObject clEnqueueUnmapMemObject;
MAGICKpfn_clEnqueueNDRangeKernel clEnqueueNDRangeKernel;
MAGICKpfn_clGetEventProfilingInfo clGetEventProfilingInfo;
MAGICKpfn_clWaitForEvents clWaitForEvents;
MAGICKpfn_clReleaseEvent clReleaseEvent;
MAGICKpfn_clFinish clFinish;
};
struct _MagickCLEnv {
MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */
MagickBooleanType OpenCLDisabled; /* whether if OpenCL has been explicitely disabled. */
struct _MagickCLDevice
{
char
*name,
*version;
MagickLibrary * library;
cl_command_queue
command_queues[MAGICKCORE_OPENCL_COMMAND_QUEUES];
/*OpenCL objects */
cl_platform_id platform;
cl_device_type deviceType;
cl_device_id device;
cl_context context;
cl_context
context;
MagickBooleanType disableProgramCache; /* disable the OpenCL program cache */
cl_program programs[MAGICK_OPENCL_NUM_PROGRAMS]; /* one program object maps one kernel source file */
cl_device_id
deviceID;
MagickBooleanType regenerateProfile; /* re-run the microbenchmark in auto device selection mode */
cl_device_type
type;
/* Random number generator seeds */
unsigned int numGenerators;
float randNormalize;
cl_mem seeds;
SemaphoreInfo* seedsLock;
cl_platform_id
platform;
SemaphoreInfo* lock;
cl_program
program;
cl_command_queue commandQueues[MAX_COMMAND_QUEUES];
ssize_t commandQueuesPos;
SemaphoreInfo* commandQueuesLock;
cl_uint
max_clock_frequency,
max_compute_units;
cl_ulong
local_memory_size;
double
score;
MagickBooleanType
enabled;
SemaphoreInfo
*lock;
ssize_t
command_queues_index,
created_queues;
#if MAGICKCORE_OPENCL_PROFILE_KERNELS
KernelProfileRecord
profileRecords[KERNEL_COUNT];
#endif
};
struct _MagickCLEnv
{
double
cpu_score;
MagickBooleanType
enabled,
initialized;
MagickCLDevice
*devices;
MagickLibrary
*library;
MagickThreadType
benchmark_thread_id;
SemaphoreInfo
*lock;
size_t
number_devices;
};
#endif
@@ -355,93 +400,38 @@ struct _MagickCLEnv {
#define CLCharQuantumScale 72340172838076673.0f
#endif
typedef enum {
AddNoiseKernel,
BlurRowKernel,
BlurColumnKernel,
CompositeKernel,
ComputeFunctionKernel,
ContrastKernel,
ContrastStretchKernel,
ConvolveKernel,
EqualizeKernel,
GrayScaleKernel,
HistogramKernel,
HullPass1Kernel,
HullPass2Kernel,
LocalContrastBlurRowKernel,
LocalContrastBlurApplyColumnKernel,
ModulateKernel,
MotionBlurKernel,
RandomNumberGeneratorKernel,
ResizeHorizontalKernel,
ResizeVerticalKernel,
RotationalBlurKernel,
UnsharpMaskBlurColumnKernel,
UnsharpMaskKernel,
WaveletDenoiseKernel,
KERNEL_COUNT
} ProfiledKernels;
extern MagickPrivate cl_command_queue
AcquireOpenCLCommandQueue(MagickCLDevice);
extern MagickPrivate cl_context
GetOpenCLContext(MagickCLEnv);
extern MagickPrivate cl_kernel
AcquireOpenCLKernel(MagickCLDevice,const char *);
extern MagickPrivate cl_kernel
AcquireOpenCLKernel(MagickCLEnv, MagickOpenCLProgram, const char*);
extern MagickPrivate MagickBooleanType
InitializeOpenCL(MagickCLEnv,ExceptionInfo *),
OpenCLThrowMagickException(MagickCLDevice,ExceptionInfo *,
const char *,const char *,const size_t,const ExceptionType,const char *,
const char *,...);
extern MagickPrivate cl_command_queue
AcquireOpenCLCommandQueue(MagickCLEnv);
extern MagickPrivate MagickCLDevice
GetOpenCLDevice(MagickCLEnv);
extern MagickPrivate MagickBooleanType
OpenCLThrowMagickException(ExceptionInfo *,
const char *,const char *,const size_t,
const ExceptionType,const char *,const char *,...),
RelinquishOpenCLCommandQueue(MagickCLEnv, cl_command_queue),
RelinquishOpenCLKernel(MagickCLEnv, cl_kernel);
extern MagickPrivate MagickCLEnv
GetCurrentOpenCLEnv(void);
extern MagickPrivate unsigned long
GetOpenCLDeviceLocalMemorySize(MagickCLEnv),
GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv);
extern MagickPrivate const char*
GetOpenCLCachedFilesDirectory();
extern MagickPrivate void
OpenCLLog(const char*),
UnlockRandSeedBuffer(MagickCLEnv);
extern MagickPrivate cl_mem
GetAndLockRandSeedBuffer(MagickCLEnv);
extern MagickPrivate unsigned int
GetNumRandGenerators(MagickCLEnv);
extern MagickPrivate float
GetRandNormalize(MagickCLEnv);
extern MagickPrivate unsigned long
GetOpenCLDeviceLocalMemorySize(const MagickCLDevice);
extern MagickPrivate void
DumpOpenCLProfileData(),
OpenCLTerminus(),
RecordProfileData(MagickCLEnv,ProfiledKernels,cl_event);
/* #define OPENCLLOG_ENABLED 1 */
static inline void OpenCLLogException(const char* function,
const unsigned int line,
ExceptionInfo* exception) {
#ifdef OPENCLLOG_ENABLED
if (exception->severity!=0) {
char message[MagickPathExtent];
/* dump the source into a file */
(void) FormatLocaleString(message,MagickPathExtent,"%s:%d Exception(%d):%s "
,function,line,exception->severity,exception->reason);
OpenCLLog(message);
}
#else
magick_unreferenced(function);
magick_unreferenced(line);
magick_unreferenced(exception);
#endif
}
RecordProfileData(MagickCLDevice,ProfiledKernels,cl_event),
RelinquishOpenCLCommandQueue(MagickCLDevice,cl_command_queue),
RelinquishOpenCLKernel(cl_kernel);
extern MagickPrivate cl_mem
GetAndLockRandSeedBuffer(MagickCLEnv,MagickCLDevice);
extern MagickPrivate void
UnlockRandSeedBuffer(MagickCLEnv);
#if defined(__cplusplus) || defined(c_plusplus)
}
MagickCore/opencl.c
+2126 -2744
View File
File diff suppressed because it is too large Load Diff
MagickCore/opencl.h
+28 -35
View File
@@ -22,45 +22,38 @@
extern "C" {
#endif
/* OpenCL program modules */
typedef enum {
MAGICK_OPENCL_ACCELERATE = 0,
MAGICK_OPENCL_NUM_PROGRAMS /* !!! This has to be the last entry !!! */
} MagickOpenCLProgram;
typedef enum {
MAGICK_OPENCL_OFF = 0
, MAGICK_OPENCL_DEVICE_SELECT_AUTO = 1
, MAGICK_OPENCL_DEVICE_SELECT_USER = 2
, MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE = 3
} ImageMagickOpenCLMode;
/* Parameter type accepted by SetMagickOpenCLEnvParm and GetMagickOpenCLEnvParm */
typedef enum {
MAGICK_OPENCL_ENV_PARAM_DEVICE /* cl_device_id (from OpenCL) */
, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED /* MagickBooleanType */
, MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED /* MagickBooleanType */
, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED /* MagickBooleanType */
/* if true, disable the kernel binary cache */
, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE /* MagickBooleanType */
/* if true, rerun microbenchmark in auto device selection */
} MagickOpenCLEnvParam;
typedef enum
{
UndefinedCLDeviceType,
CpuCLDeviceType,
GpuCLDeviceType
} MagickCLDeviceType;
typedef struct _MagickCLDevice* MagickCLDevice;
typedef struct _MagickCLEnv* MagickCLEnv;
extern MagickExport MagickBooleanType
GetMagickOpenCLEnvParam(MagickCLEnv,MagickOpenCLEnvParam,size_t,void*,
ExceptionInfo*),
InitImageMagickOpenCL(ImageMagickOpenCLMode,void*,void*,ExceptionInfo*),
InitOpenCLEnv(MagickCLEnv,ExceptionInfo*),
RelinquishMagickOpenCLEnv(MagickCLEnv),
SetMagickOpenCLEnvParam(MagickCLEnv,MagickOpenCLEnvParam,size_t,void*,
ExceptionInfo*);
extern MagickExport const char
*GetOpenCLDeviceName(const MagickCLDevice),
*GetOpenCLDeviceVersion(const MagickCLDevice);
extern MagickExport const MagickCLDevice
*GetOpenCLDevices(size_t *);
extern MagickExport double
GetOpenCLDeviceBenchmarkScore(const MagickCLDevice);
extern MagickExport MagickCLDeviceType
GetOpenCLDeviceType(const MagickCLDevice);
extern MagickExport MagickBooleanType
GetOpenCLDeviceEnabled(const MagickCLDevice),
GetOpenCLEnabled(void),
SetOpenCLEnabled(const MagickBooleanType);
extern MagickExport void
SetOpenCLDeviceEnabled(const MagickCLDevice,
const MagickBooleanType);
extern MagickExport MagickCLEnv
AcquireMagickOpenCLEnv(void),
GetDefaultOpenCLEnv(void),
SetDefaultOpenCLEnv(MagickCLEnv);
#if defined(__cplusplus) || defined(c_plusplus)
}