(no commit message)

[imagemagick] / MagickCore / accelerate.c

diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c
index a2b5f4880b83d92e21b8bc634e8dff1175df1c51..c7b0b44c3fc940cc0f73a0925358c60ea576a25a 100644 (file)
--- a/MagickCore/accelerate.c
+++ b/MagickCore/accelerate.c
@@ -17,7 +17,7 @@
 %                               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  %
@@ -185,7 +185,7 @@ static const char
     "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"
@@ -195,6 +195,13 @@ static const char
     "#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"
@@ -295,7 +302,7 @@ static const char
     "      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"
@@ -316,7 +323,7 @@ static void ConvolveNotify(const char *message,const void *data,size_t length,
   (void) length;
   exception=(ExceptionInfo *) user_context;
   (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
-    "DelegateFailed","`%s'",message);
+    "DelegateFailed","'%s'",message);
 }
 
 static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
@@ -376,7 +383,8 @@ static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info,
     &convolve_info->height);
   if (status != CL_SUCCESS)
     return(MagickFalse);
-  convolve_info->matte=(cl_uint) image->matte;
+  convolve_info->matte=(cl_uint) image->alpha_trait == BlendPixelTrait ?
+    MagickTrue : MagickFalse;
   status=clSetKernelArg(convolve_info->kernel,i++,sizeof(cl_uint),(void *)
     &convolve_info->matte);
   if (status != CL_SUCCESS)
@@ -496,7 +504,7 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
   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));
@@ -509,7 +517,7 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
   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);
     }
@@ -527,7 +535,7 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
   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);
     }
@@ -545,7 +553,7 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
   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);
     }
@@ -556,6 +564,39 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
       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.
   */
@@ -597,7 +638,7 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name,
       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);
@@ -629,7 +670,7 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
   assert(convolve_image->signature == MagickSignature);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickSignature);
-  if ((image->storage_class != DirectClass) || 
+  if ((image->storage_class != DirectClass) ||
       (image->colorspace == CMYKColorspace))
     return(MagickFalse);
   if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) &&
@@ -670,7 +711,7 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
       {
         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);
@@ -678,7 +719,7 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
       {
         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*
@@ -688,7 +729,7 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image,
         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++)