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);
*/
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,
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:
}
- setGray(&filteredPixel, ClampToQuantum(intensity));
-
- filteredPixel.w = pixel.w;
-
- im[c] = filteredPixel;
+ im[c] = im[c+1] = im[c+2] = ClampToQuantum(intensity);
}
)
cl_int
clStatus,
+ number_channels,
+ colorspace,
intensityMethod;
- cl_int
- colorspace;
-
cl_kernel
grayscaleKernel;
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;
}
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)
{
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)
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)
{
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);
assert(image != NULL);
assert(exception != (ExceptionInfo *) NULL);
- if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
+ if ((checkAccelerateCondition(image) == MagickFalse) ||
(checkOpenCLEnvironment(exception) == MagickFalse))
return(MagickFalse);
if (image->colorspace != sRGBColorspace)
return(MagickFalse);
+ if (image->number_channels < 3)
+ return(MagickFalse);
+
status=ComputeGrayscaleImage(image,method,exception);
return(status);
}
"-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)
"-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)
"-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)
"-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)
"-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