% January 2010 %
% %
% %
-% Copyright 1999-2011 ImageMagick Studio LLC, a non-profit organization %
+% Copyright 1999-2012 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 %
#if defined(MAGICKCORE_HDRI_SUPPORT)
#define CLOptions "-DMAGICKCORE_HDRI_SUPPORT=1 -DCLQuantum=float " \
"-DCLPixelType=float4 -DQuantumRange=%g -DMagickEpsilon=%g"
-#define CLPixelPacket cl_float4
+#define CLPixelInfo cl_float4
#else
#if (MAGICKCORE_QUANTUM_DEPTH == 8)
#define CLOptions "-DCLQuantum=uchar -DCLPixelType=uchar4 " \
"-DQuantumRange=%g -DMagickEpsilon=%g"
-#define CLPixelPacket cl_uchar4
+#define CLPixelInfo cl_uchar4
#elif (MAGICKCORE_QUANTUM_DEPTH == 16)
#define CLOptions "-DCLQuantum=ushort -DCLPixelType=ushort4 " \
"-DQuantumRange=%g -DMagickEpsilon=%g"
-#define CLPixelPacket cl_ushort4
+#define CLPixelInfo cl_ushort4
#elif (MAGICKCORE_QUANTUM_DEPTH == 32)
#define CLOptions "-DCLQuantum=uint -DCLPixelType=uint4 " \
"-DQuantumRange=%g -DMagickEpsilon=%g"
-#define CLPixelPacket cl_uint4
+#define CLPixelInfo cl_uint4
#elif (MAGICKCORE_QUANTUM_DEPTH == 64)
#define CLOptions "-DCLQuantum=ussize_t -DCLPixelType=ussize_t4 " \
"-DQuantumRange=%g -DMagickEpsilon=%g"
-#define CLPixelPacket cl_ulong4
+#define CLPixelInfo cl_ulong4
#endif
#endif
filter;
} ConvolveInfo;
-static char
+static const char
*ConvolveKernel =
"static inline long ClampToCanvas(const long offset,const unsigned long range)\n"
"{\n"
"static inline CLQuantum ClampToQuantum(const float value)\n"
"{\n"
"#if defined(MAGICKCORE_HDRI_SUPPORT)\n"
- " return((CLQuantum) value)\n"
+ " return((CLQuantum) value);\n"
"#else\n"
" if (value < 0.0)\n"
" return((CLQuantum) 0);\n"
"#endif\n"
"}\n"
"\n"
+ "static inline float MagickEpsilonReciprocal(const float x)\n"
+ "{\n"
+ " float sign = x < (float) 0.0 ? (float) -1.0 : (float) 1.0;\n"
+ " return((sign*x) >= MagickEpsilon ? (float) 1.0/x : sign*((float) 1.0/\n"
+ " MagickEpsilon));\n"
+ "}\n"
+ "\n"
"__kernel void Convolve(const __global CLPixelType *input,\n"
" __constant float *filter,const unsigned long width,const unsigned long height,\n"
" const unsigned int matte,__global CLPixelType *output)\n"
" break;\n"
" }\n"
" }\n"
- " gamma=1.0/(fabs(gamma) <= MagickEpsilon ? 1.0 : gamma);\n"
+ " gamma=MagickEpsilonReciprocal(gamma);\n"
" const unsigned long index = y*columns+x;\n"
" output[index].x=ClampToQuantum(gamma*sum.x);\n"
" output[index].y=ClampToQuantum(gamma*sum.y);\n"
(void) length;
exception=(ExceptionInfo *) user_context;
(void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
- "DelegateFailed","`%s'",message);
+ "DelegateFailed","'%s'",message);
}
static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
*/
length=image->columns*image->rows;
convolve_info->pixels=clCreateBuffer(convolve_info->context,(cl_mem_flags)
- (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(CLPixelPacket),
+ (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(CLPixelInfo),
(void *) pixels,&status);
if ((convolve_info->pixels == (cl_mem) NULL) || (status != CL_SUCCESS))
return(MagickFalse);
length=image->columns*image->rows;
convolve_info->convolve_pixels=clCreateBuffer(convolve_info->context,
(cl_mem_flags) (CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR),length*
- sizeof(CLPixelPacket),convolve_pixels,&status);
+ sizeof(CLPixelInfo),convolve_pixels,&status);
if ((convolve_info->convolve_pixels == (cl_mem) NULL) ||
(status != CL_SUCCESS))
return(MagickFalse);
cl_int
status;
- (void) status;
+ status=0;
if (convolve_info->convolve_pixels != (cl_mem) NULL)
status=clReleaseMemObject(convolve_info->convolve_pixels);
if (convolve_info->pixels != (cl_mem) NULL)
status=clReleaseMemObject(convolve_info->pixels);
if (convolve_info->filter != (cl_mem) NULL)
status=clReleaseMemObject(convolve_info->filter);
+ (void) status;
}
static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info)
cl_int
status;
- (void) status;
+ status=0;
if (convolve_info->kernel != (cl_kernel) NULL)
status=clReleaseKernel(convolve_info->kernel);
if (convolve_info->program != (cl_program) NULL)
status=clReleaseCommandQueue(convolve_info->command_queue);
if (convolve_info->context != (cl_context) NULL)
status=clReleaseContext(convolve_info->context);
+ (void) status;
convolve_info=(ConvolveInfo *) RelinquishMagickMemory(convolve_info);
return(convolve_info);
}
length=image->columns*image->rows;
status=clEnqueueWriteBuffer(convolve_info->command_queue,
- convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),pixels,0,NULL,
+ convolve_info->pixels,CL_TRUE,0,length*sizeof(CLPixelInfo),pixels,0,NULL,
NULL);
length=width*height;
status=clEnqueueWriteBuffer(convolve_info->command_queue,
return(MagickFalse);
length=image->columns*image->rows;
status=clEnqueueReadBuffer(convolve_info->command_queue,
- convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelPacket),
+ convolve_info->convolve_pixels,CL_TRUE,0,length*sizeof(CLPixelInfo),
convolve_pixels,0,NULL,NULL);
if (status != CL_SUCCESS)
return(MagickFalse);
if (convolve_info == (ConvolveInfo *) NULL)
{
(void) ThrowMagickException(exception,GetMagickModule(),
- ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
+ ResourceLimitError,"MemoryAllocationFailed","'%s'",image->filename);
return((ConvolveInfo *) NULL);
}
(void) ResetMagickMemory(convolve_info,0,sizeof(*convolve_info));
if (status != CL_SUCCESS)
{
(void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
- "failed to create OpenCL context","`%s' (%d)",image->filename,status);
+ "failed to create OpenCL context","'%s' (%d)",image->filename,status);
convolve_info=DestroyConvolveInfo(convolve_info);
return((ConvolveInfo *) NULL);
}
if ((convolve_info->context == (cl_context) NULL) || (status != CL_SUCCESS))
{
(void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
- "failed to create OpenCL context","`%s' (%d)",image->filename,status);
+ "failed to create OpenCL context","'%s' (%d)",image->filename,status);
convolve_info=DestroyConvolveInfo(convolve_info);
return((ConvolveInfo *) NULL);
}
if (convolve_info->devices == (cl_device_id *) NULL)
{
(void) ThrowMagickException(exception,GetMagickModule(),
- ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
+ ResourceLimitError,"MemoryAllocationFailed","'%s'",image->filename);
convolve_info=DestroyConvolveInfo(convolve_info);
return((ConvolveInfo *) NULL);
}
convolve_info=DestroyConvolveInfo(convolve_info);
return((ConvolveInfo *) NULL);
}
+ if (image->debug != MagickFalse)
+ {
+ char
+ attribute[MaxTextExtent];
+
+ size_t
+ length;
+
+ clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_NAME,
+ sizeof(attribute),attribute,&length);
+ (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Name: %s",
+ attribute);
+ clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VENDOR,
+ sizeof(attribute),attribute,&length);
+ (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Vendor: %s",
+ attribute);
+ clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_VERSION,
+ sizeof(attribute),attribute,&length);
+ (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
+ "Driver Version: %s",attribute);
+ clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_PROFILE,
+ sizeof(attribute),attribute,&length);
+ (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Profile: %s",
+ attribute);
+ clGetDeviceInfo(convolve_info->devices[0],CL_DRIVER_VERSION,
+ sizeof(attribute),attribute,&length);
+ (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Driver: %s",
+ attribute);
+ clGetDeviceInfo(convolve_info->devices[0],CL_DEVICE_EXTENSIONS,
+ sizeof(attribute),attribute,&length);
+ (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Extensions: %s",
+ attribute);
+ }
/*
Create OpenCL command queue.
*/
status=clGetProgramBuildInfo(convolve_info->program,
convolve_info->devices[0],CL_PROGRAM_BUILD_LOG,length,log,&length);
(void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
- "failed to build OpenCL program","`%s' (%s)",image->filename,log);
+ "failed to build OpenCL program","'%s' (%s)",image->filename,log);
log=DestroyString(log);
convolve_info=DestroyConvolveInfo(convolve_info);
return((ConvolveInfo *) NULL);
{
convolve_info=DestroyConvolveInfo(convolve_info);
(void) ThrowMagickException(exception,GetMagickModule(),CacheError,
- "UnableToReadPixelCache","`%s'",image->filename);
+ "UnableToReadPixelCache","'%s'",image->filename);
return(MagickFalse);
}
convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception);
{
convolve_info=DestroyConvolveInfo(convolve_info);
(void) ThrowMagickException(exception,GetMagickModule(),CacheError,
- "UnableToReadPixelCache","`%s'",image->filename);
+ "UnableToReadPixelCache","'%s'",image->filename);
return(MagickFalse);
}
filter=(float *) AcquireQuantumMemory(kernel->width,kernel->height*
DestroyConvolveBuffers(convolve_info);
convolve_info=DestroyConvolveInfo(convolve_info);
(void) ThrowMagickException(exception,GetMagickModule(),
- ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename);
+ ResourceLimitError,"MemoryAllocationFailed","'%s'",image->filename);
return(MagickFalse);
}
for (i=0; i < (ssize_t) (kernel->width*kernel->height); i++)
projects / imagemagick / blobdiff