From 9b5639ff7016eb369ebe8b5e969997913d31e25b Mon Sep 17 00:00:00 2001 From: dirk Date: Sat, 26 Mar 2016 18:57:23 +0100 Subject: [PATCH] AccelerateGrayscaleImage now also works for RGB images. --- MagickCore/accelerate-private.h | 27 ++++++++----------------- MagickCore/accelerate.c | 36 ++++++++++++++++++--------------- MagickCore/opencl-private.h | 5 +++++ 3 files changed, 33 insertions(+), 35 deletions(-) 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 -- 2.40.0