STRINGIFY(
- inline __global CLQuantum *getPixel(__global CLQuantum *image, const unsigned int number_channels,
+ inline unsigned int getPixelIndex(const unsigned int number_channels,
const unsigned int columns, const unsigned int x, const unsigned int y)
{
- return image + (x * number_channels) + (y * columns * number_channels);
+ return (x * number_channels) + (y * columns * number_channels);
}
inline float getPixelRed(const __global CLQuantum *p) { return (float)*p; }
*alpha=getPixelAlpha(p);
}
- inline float4 ReadFloat4(__global CLQuantum *image, const unsigned int number_channels,
+ inline float4 ReadFloat4(const __global CLQuantum *image, const unsigned int number_channels,
const unsigned int columns, const unsigned int x, const unsigned int y, const ChannelType channel)
{
- const __global CLQuantum *p = getPixel(image, number_channels, columns, x, y);
+ const __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
float red = 0.0f;
float green = 0.0f;
const unsigned int columns, const unsigned int x, const unsigned int y, const ChannelType channel,
float4 pixel)
{
- __global CLQuantum *p = getPixel(image, number_channels, columns, x, y);
+ __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
WriteChannels(p, number_channels, channel, pixel.x, pixel.y, pixel.z, pixel.w);
}
/*
Reduce image noise and reduce detail levels by row
*/
- __kernel void BlurRow(__global CLQuantum *image,
+ __kernel void BlurRow(const __global CLQuantum *image,
const unsigned int number_channels,const ChannelType channel,
__constant float *filter,const unsigned int width,
const unsigned int imageColumns,const unsigned int imageRows,
number_parameters : numbers of parameters
parameters : the parameter
*/
- __kernel void ComputeFunction(__global CLQuantum *image,
- const unsigned int number_channels,const ChannelType channel,
- const MagickFunction function,const unsigned int number_parameters,
+ __kernel void ComputeFunction(__global CLQuantum *image,const unsigned int number_channels,
+ const ChannelType channel,const MagickFunction function,const unsigned int number_parameters,
__constant float *parameters)
{
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
const unsigned int columns = get_global_size(0);
- __global CLQuantum *p = getPixel(image, number_channels, columns, x, y);
+ __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
float red;
float green;
const unsigned int x = get_global_id(0);
const unsigned int y = get_global_id(1);
const unsigned int columns = get_global_size(0);
- __global CLQuantum *p = getPixel(image, number_channels, columns, x, y);
+ __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y);
float
blue,
*/
STRINGIFY(
- __kernel void UnsharpMaskBlurColumn(const __global CLPixelType* inputImage,
- const __global float4 *blurRowData, __global CLPixelType *filtered_im,
- const unsigned int imageColumns, const unsigned int imageRows,
- __local float4* cachedData, __local float* cachedFilter,
- const ChannelType channel, const __global float *filter, const unsigned int width,
- const float gain, const float threshold)
+ __kernel void UnsharpMaskBlurColumn(const __global CLQuantum* image,
+ const __global float4 *blurRowData,const unsigned int number_channels,
+ const ChannelType channel,const unsigned int columns,
+ const unsigned int rows,__local float4* cachedData,
+ __local float* cachedFilter,const __global float *filter,
+ const unsigned int width,const float gain, const float threshold,
+ __global CLQuantum *filteredImage)
{
const unsigned int radius = (width-1)/2;
const int groupStartY = get_group_id(1)*get_local_size(1) - radius;
const int groupStopY = (get_group_id(1)+1)*get_local_size(1) + radius;
- if (groupStartY >= 0
- && groupStopY < imageRows) {
- event_t e = async_work_group_strided_copy(cachedData
- ,blurRowData+groupStartY*imageColumns+groupX
- ,groupStopY-groupStartY,imageColumns,0);
+ if ((groupStartY >= 0) && (groupStopY < rows))
+ {
+ event_t e = async_work_group_strided_copy(cachedData,
+ blurRowData+groupStartY*columns+groupX,groupStopY-groupStartY,columns,0);
wait_group_events(1,&e);
}
- else {
- for (int i = get_local_id(1); i < (groupStopY - groupStartY); i+=get_local_size(1)) {
- cachedData[i] = blurRowData[ClampToCanvas(groupStartY+i,imageRows)*imageColumns+ groupX];
- }
+ else
+ {
+ for (int i = get_local_id(1); i < (groupStopY - groupStartY); i+=get_local_size(1))
+ cachedData[i] = blurRowData[ClampToCanvas(groupStartY+i,rows)*columns + groupX];
+
barrier(CLK_LOCAL_MEM_FENCE);
}
// cache the filter as well
wait_group_events(1,&e);
// only do the work if this is not a patched item
- //const int cy = get_group_id(1)*get_local_size(1)+get_local_id(1);
const int cy = get_global_id(1);
- if (cy < imageRows) {
+ if (cy < rows)
+ {
float4 blurredPixel = (float4) 0.0f;
int i = 0;
- \n #ifndef UFACTOR \n
- \n #define UFACTOR 8 \n
- \n #endif \n
-
- for ( ; i+UFACTOR < width; )
- {
- \n #pragma unroll UFACTOR \n
- for (int j=0; j < UFACTOR; j++, i++)
- {
- blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)];
- }
- }
-
- for ( ; i < width; i++)
+ for ( ; i+7 < width; )
{
- blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)];
+ for (int j=0; j < 8; j++, i++)
+ blurredPixel+=cachedFilter[i+j]*cachedData[i+j+get_local_id(1)];
}
- blurredPixel = floor((float4)(ClampToQuantum(blurredPixel.x), ClampToQuantum(blurredPixel.y)
- ,ClampToQuantum(blurredPixel.z), ClampToQuantum(blurredPixel.w)));
+ for ( ; i < width; i++)
+ blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)];
- float4 inputImagePixel = convert_float4(inputImage[cy*imageColumns+groupX]);
+ float4 inputImagePixel = ReadFloat4(image,number_channels,columns,groupX,cy,channel);
float4 outputPixel = inputImagePixel - blurredPixel;
float quantumThreshold = QuantumRange*threshold;
outputPixel = select(inputImagePixel + outputPixel * gain, inputImagePixel, mask);
//write back
- filtered_im[cy*imageColumns+groupX] = (CLPixelType) (ClampToQuantum(outputPixel.x), ClampToQuantum(outputPixel.y)
- ,ClampToQuantum(outputPixel.z), ClampToQuantum(outputPixel.w));
-
+ WriteFloat4(filteredImage,number_channels,columns,groupX,cy,channel,outputPixel);
}
}
)
STRINGIFY(
- __kernel void UnsharpMask(__global CLQuantum *image,const unsigned int number_channels,
+ __kernel void UnsharpMask(const __global CLQuantum *image,const unsigned int number_channels,
const ChannelType channel,__constant float *filter,const unsigned int width,
const unsigned int columns,const unsigned int rows,__local float4 *pixels,
const float gain,const float threshold, const unsigned int justBlur,
*filteredImage_view,
*image_view;
- char
- geometry[MagickPathExtent];
-
cl_command_queue
queue;
imageKernelBuffer,
tempImageBuffer;
- cl_mem_flags
- mem_flags;
-
- const void
- *inputPixels;
+ cl_uint
+ imageColumns,
+ imageRows,
+ kernelWidth,
+ number_channels;
float
fGain,
- fThreshold,
- *kernelBufferPtr;
+ fThreshold;
Image
*filteredImage;
int
chunkSize;
- KernelInfo
- *kernel;
-
MagickBooleanType
outputReady;
length;
void
- *filteredPixels,
- *hostPtr;
+ *filteredPixels;
unsigned int
- i,
- imageColumns,
- imageRows,
- kernelWidth;
+ i;
clEnv = NULL;
filteredImage = NULL;
filteredImage_view = NULL;
- kernel = NULL;
context = NULL;
imageBuffer = NULL;
filteredImageBuffer = NULL;
+ filteredPixels = NULL;
tempImageBuffer = NULL;
imageKernelBuffer = NULL;
blurRowKernel = NULL;
context = GetOpenCLContext(clEnv);
queue = AcquireOpenCLCommandQueue(clEnv);
- /* Create and initialize OpenCL buffers. */
- {
- image_view=AcquireVirtualCacheView(image,exception);
- inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
- if (inputPixels == (const void *) NULL)
- {
- (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
- goto cleanup;
- }
-
- /* If the host pointer is aligned to the size of CLPixelPacket,
- 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))
- {
- mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR;
- }
- else
- {
- mem_flags = CL_MEM_READ_ONLY|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);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
- goto cleanup;
- }
- }
+ image_view = AcquireVirtualCacheView(image, exception);
+ imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception);
+ if (imageBuffer == (cl_mem) NULL)
+ goto cleanup;
- /* create output */
+ filteredImage=CloneImage(image,0,0,MagickTrue,exception);
+ if (filteredImage == (Image *) NULL)
+ goto cleanup;
+ if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue)
{
- filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
- assert(filteredImage != NULL);
- if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
- goto cleanup;
- }
- filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
- filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
- if (filteredPixels == (void *) NULL)
- {
- (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
- goto cleanup;
- }
-
- if (ALIGNED(filteredPixels,CLPixelPacket))
- {
- mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR;
- hostPtr = filteredPixels;
- }
- else
- {
- mem_flags = CL_MEM_WRITE_ONLY;
- hostPtr = NULL;
- }
-
- /* create a CL buffer from image pixel buffer */
- length = image->columns * image->rows;
- filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
- goto cleanup;
- }
+ (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+ goto cleanup;
}
+ filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception);
+ filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
+ context,filteredPixels,exception);
+ if (filteredImageBuffer == (cl_mem) NULL)
+ goto cleanup;
- /* create the blur kernel */
- {
- (void) FormatLocaleString(geometry,MagickPathExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
- kernel=AcquireKernelInfo(geometry,exception);
- if (kernel == (KernelInfo *) NULL)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
- goto cleanup;
- }
-
- imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
- goto cleanup;
- }
-
-
- kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
- goto cleanup;
- }
- for (i = 0; i < kernel->width; i++)
- {
- kernelBufferPtr[i] = (float) kernel->values[i];
- }
- clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
- goto cleanup;
- }
- }
+ imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma,
+ &kernelWidth,exception);
{
/* create temp buffer */
{
length = image->columns * image->rows;
- tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus);
+ tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus);
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
};
}
+ number_channels = (cl_uint) image->number_channels;
+ imageColumns = (cl_uint) image->columns;
+ imageRows = (cl_uint) image->rows;
+
{
chunkSize = 256;
- imageColumns = (unsigned int) image->columns;
- imageRows = (unsigned int) image->rows;
-
- kernelWidth = (unsigned int) kernel->width;
-
/* set the kernel arguments */
i = 0;
clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
- clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
+ clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels);
clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask);
clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
- clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
- clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
- clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows);
- clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL);
+ clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+ clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+ clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+ clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
+ clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
wsize[0] = chunkSize;
wsize[1] = 1;
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
{
chunkSize = 256;
- imageColumns = (unsigned int) image->columns;
- imageRows = (unsigned int) image->rows;
- kernelWidth = (unsigned int) kernel->width;
fGain = (float) gain;
fThreshold = (float) threshold;
i = 0;
clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),&number_channels);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,(chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,kernelWidth*sizeof(float),NULL);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
if (clStatus != CL_SUCCESS)
{
wsize[0] = 1;
wsize[1] = chunkSize;
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
}
/* get result */
- if (ALIGNED(filteredPixels,CLPixelPacket))
- {
- length = image->columns * image->rows;
- clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus);
- }
- else
- {
- length = image->columns * image->rows;
- clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL);
- }
- if (clStatus != CL_SUCCESS)
+ if (copyWriteBuffer(image,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
{
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+ (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
goto cleanup;
}
if (filteredImage_view != NULL)
filteredImage_view=DestroyCacheView(filteredImage_view);
- if (kernel != NULL) kernel=DestroyKernelInfo(kernel);
if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer);
assert(image != NULL);
assert(exception != (ExceptionInfo *) NULL);
- if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
+ if ((checkAccelerateCondition(image) == MagickFalse) ||
(checkOpenCLEnvironment(exception) == MagickFalse))
return NULL;