From: dirk <dirk@git.imagemagick.org>
Date: Sat, 26 Mar 2016 17:57:23 +0000 (+0100)
Subject: AccelerateGrayscaleImage now also works for RGB images.
X-Git-Tag: 7.0.1-0~123
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=9b5639ff7016eb369ebe8b5e969997913d31e25b;p=imagemagick

AccelerateGrayscaleImage now also works for RGB images.
---

diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h
index a26e0a78a..77e1eacb8 100644
--- a/MagickCore/accelerate-private.h
+++ b/MagickCore/accelerate-private.h
@@ -385,8 +385,6 @@ OPENCL_ENDIF()
   inline float getAlphaF4(float4 p)                     { return p.w; }
   inline void setAlphaF4(float4* p, float value)        { (*p).w = value; }
 
-  inline void setGray(CLPixelType* p, CLQuantum value)  { (*p).z = value; (*p).y = value; (*p).x = value; }
-
   inline float GetPixelIntensity(const int method, const int colorspace, CLPixelType p)
   {
     float red = getRed(p);
@@ -2110,16 +2108,13 @@ uint MWC64X_NextUint(mwc64x_state_t *s)
 */
 
   STRINGIFY(
-  __kernel void Grayscale(__global CLPixelType *im, 
-    const int method, const int colorspace)
+  __kernel void Grayscale(__global CLQuantum *im,const int number_channels,
+    const int method,const int colorspace)
   {
-
-    const int x = get_global_id(0);  
+    const int x = get_global_id(0);
     const int y = get_global_id(1);
     const int columns = get_global_size(0);
-    const int c = x + y * columns;
-
-    CLPixelType pixel = im[c];
+    const int c = (x * number_channels) + (y * columns * number_channels);
 
     float
         blue,
@@ -2127,14 +2122,12 @@ uint MWC64X_NextUint(mwc64x_state_t *s)
         intensity,
         red;
 
-    red=(float)getRed(pixel);
-    green=(float)getGreen(pixel);
-    blue=(float)getBlue(pixel);
+    red=(float)im[c];
+    green=(float)im[c+1];
+    blue=(float)im[c+2];
 
     intensity=0.0;
 
-    CLPixelType filteredPixel;
- 
     switch (method)
     {
       case AveragePixelIntensityMethod:
@@ -2221,11 +2214,7 @@ uint MWC64X_NextUint(mwc64x_state_t *s)
 
     }
 
-    setGray(&filteredPixel, ClampToQuantum(intensity));
-
-    filteredPixel.w = pixel.w;
-
-    im[c] = filteredPixel;
+    im[c] = im[c+1] = im[c+2] = ClampToQuantum(intensity);
   }
   )
 
diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c
index 8ce7433e4..e1b31387a 100644
--- a/MagickCore/accelerate.c
+++ b/MagickCore/accelerate.c
@@ -3942,11 +3942,10 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,
 
   cl_int
     clStatus,
+    number_channels,
+    colorspace,
     intensityMethod;
 
-  cl_int
-    colorspace;
-
   cl_kernel
     grayscaleKernel;
 
@@ -4008,7 +4007,7 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,
    then use the host buffer directly from the GPU; otherwise, 
    create a buffer on the GPU and copy the data over
    */
-  if (ALIGNED(inputPixels,CLPixelPacket)) 
+  if (ALIGNED(inputPixels,CLQuantum))
   {
     mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR;
   }
@@ -4017,17 +4016,14 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,
     mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
   }
   /* create a CL buffer from image pixel buffer */
-  length = image->columns * image->rows;
-  imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus);
+  length = image->columns * image->rows * image->number_channels;
+  imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLQuantum), (void*)inputPixels, &clStatus);
   if (clStatus != CL_SUCCESS)
   {
     (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
     goto cleanup;
   }
 
-  intensityMethod = method;
-  colorspace = image->colorspace;
-
   grayscaleKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Grayscale");
   if (grayscaleKernel == NULL)
   {
@@ -4035,8 +4031,13 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,
     goto cleanup;
   }
 
+  number_channels = (cl_int) image->number_channels;
+  intensityMethod = (cl_int) method;
+  colorspace = (cl_int) image->colorspace;
+
   i = 0;
   clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&number_channels);
   clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&intensityMethod);
   clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_int),&colorspace);
   if (clStatus != CL_SUCCESS)
@@ -4062,15 +4063,15 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,
     clEnv->library->clReleaseEvent(event);
   }
 
-  if (ALIGNED(inputPixels,CLPixelPacket)) 
+  if (ALIGNED(inputPixels,CLQuantum))
   {
     length = image->columns * image->rows;
-    clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
+    clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLQuantum), 0, NULL, NULL, &clStatus);
   }
   else 
   {
     length = image->columns * image->rows;
-    clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
+    clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLQuantum), inputPixels, 0, NULL, NULL);
   }
   if (clStatus != CL_SUCCESS)
   {
@@ -4085,11 +4086,11 @@ cleanup:
 
   image_view=DestroyCacheView(image_view);
 
-  if (imageBuffer!=NULL)		      
+  if (imageBuffer!=NULL)
     clEnv->library->clReleaseMemObject(imageBuffer);
-  if (grayscaleKernel!=NULL)                     
+  if (grayscaleKernel!=NULL)
     RelinquishOpenCLKernel(clEnv, grayscaleKernel);
-  if (queue != NULL)                          
+  if (queue != NULL)
     RelinquishOpenCLCommandQueue(clEnv, queue);
 
   return( outputReady);
@@ -4104,7 +4105,7 @@ MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image,
   assert(image != NULL);
   assert(exception != (ExceptionInfo *) NULL);
 
-  if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
+  if ((checkAccelerateCondition(image) == MagickFalse) ||
       (checkOpenCLEnvironment(exception) == MagickFalse))
     return(MagickFalse);
 
@@ -4115,6 +4116,9 @@ MagickExport MagickBooleanType AccelerateGrayscaleImage(Image* image,
   if (image->colorspace != sRGBColorspace)
     return(MagickFalse);
 
+  if (image->number_channels < 3)
+    return(MagickFalse);
+
   status=ComputeGrayscaleImage(image,method,exception);
   return(status);
 }
diff --git a/MagickCore/opencl-private.h b/MagickCore/opencl-private.h
index 86c1e7be0..02cd7a215 100644
--- a/MagickCore/opencl-private.h
+++ b/MagickCore/opencl-private.h
@@ -318,6 +318,7 @@ struct _MagickCLEnv {
   "-DCLQuantum=float -DCLSignedQuantum=float -DCLPixelType=float4 -DQuantumRange=%f " \
   "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
   " -DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
+#define CLQuantum  cl_float
 #define CLPixelPacket  cl_float4
 #define CLCharQuantumScale 1.0f
 #elif (MAGICKCORE_QUANTUM_DEPTH == 8)
@@ -325,6 +326,7 @@ struct _MagickCLEnv {
   "-DCLQuantum=uchar -DCLSignedQuantum=char -DCLPixelType=uchar4 -DQuantumRange=%ff " \
   "-DQuantumScale=%ff -DCharQuantumScale=%ff -DMagickEpsilon=%ff -DMagickPI=%ff "\
   "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
+#define CLQuantum  cl_uchar
 #define CLPixelPacket  cl_uchar4
 #define CLCharQuantumScale 1.0f
 #elif (MAGICKCORE_QUANTUM_DEPTH == 16)
@@ -332,6 +334,7 @@ struct _MagickCLEnv {
   "-DCLQuantum=ushort -DCLSignedQuantum=short -DCLPixelType=ushort4 -DQuantumRange=%f "\
   "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
   "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
+#define CLQuantum  cl_ushort
 #define CLPixelPacket  cl_ushort4
 #define CLCharQuantumScale 257.0f
 #elif (MAGICKCORE_QUANTUM_DEPTH == 32)
@@ -339,6 +342,7 @@ struct _MagickCLEnv {
   "-DCLQuantum=uint -DCLSignedQuantum=int -DCLPixelType=uint4 -DQuantumRange=%f "\
   "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
   "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
+#define CLQuantum  cl_uint
 #define CLPixelPacket  cl_uint4
 #define CLCharQuantumScale 16843009.0f
 #elif (MAGICKCORE_QUANTUM_DEPTH == 64)
@@ -346,6 +350,7 @@ struct _MagickCLEnv {
   "-DCLQuantum=ulong -DCLSignedQuantum=long -DCLPixelType=ulong4 -DQuantumRange=%f "\
   "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
   "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
+#define CLQuantum  cl_ulong
 #define CLPixelPacket  cl_ulong4
 #define CLCharQuantumScale 72340172838076673.0f
 #endif