STRINGIFY(
+ inline __global CLQuantum *getPixel(__global CLQuantum *image, 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);
+ }
+
inline float getPixelRed(const __global CLQuantum *p) { return (float)*p; }
inline float getPixelGreen(const __global CLQuantum *p) { return (float)*(p+1); }
inline float getPixelBlue(const __global CLQuantum *p) { return (float)*(p+2); }
*alpha=getPixelAlpha(p);
}
+ inline float4 ReadFloat4(__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);
+
+ float red = 0.0f;
+ float green = 0.0f;
+ float blue = 0.0f;
+ float alpha = 0.0f;
+
+ ReadChannels(p, number_channels, channel, &red, &green, &blue, &alpha);
+ return (float4)(red, green, blue, alpha);
+ }
+
inline void WriteChannels(__global CLQuantum *p, const unsigned int number_channels,
const ChannelType channel, float red, float green, float blue, float alpha)
{
setPixelAlpha(p,alpha);
}
+ inline void WriteFloat4(__global CLQuantum *image, const unsigned int number_channels,
+ 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);
+ WriteChannels(p, number_channels, channel, pixel.x, pixel.y, pixel.z, pixel.w);
+ }
+
inline float GetPixelIntensity(const unsigned int colorspace,
const unsigned int method,float red,float green,float blue)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
const int columns = get_global_size(0);
- __global CLQuantum *p = image+(x * number_channels) + (y * columns * number_channels);
+ __global CLQuantum *p = getPixel(image, number_channels, columns, x, y);
float red;
float green;
const int x = get_global_id(0);
const int y = get_global_id(1);
const int columns = get_global_size(0);
- __global CLQuantum *p = image+(x * number_channels) + (y * columns * number_channels);
+ __global CLQuantum *p = getPixel(image, number_channels, columns, x, y);
float
blue,
STRINGIFY(
- __kernel void UnsharpMask(__global CLPixelType *im,
- __global CLPixelType *filtered_im,__constant float *filter,
- const unsigned int width,const unsigned int imageColumns,
- const unsigned int imageRows,__local float4 *pixels,const float gain,
- const float threshold, const unsigned int justBlur)
+ __kernel void UnsharpMask(__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,__local float4 *pixels,
+ const float gain,const float threshold, const unsigned int justBlur,
+ __global CLQuantum *filteredImage)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
int endRow = (get_group_id(1) + 1) * get_local_size(1) + radius;
while (row < endRow) {
- int srcy = (row < 0) ? -row : row; // mirror pad
+ int srcy = (row < 0) ? -row : row; // mirror pad
srcy = (srcy >= imageRows) ? (2 * imageRows - srcy - 1) : srcy;
float4 value = 0.0f;
int srcx = ix + j;
srcx = (srcx < 0) ? -srcx : srcx;
srcx = (srcx >= imageColumns) ? (2 * imageColumns - srcx - 1) : srcx;
- value += filter[i + j] * convert_float4(im[srcx + srcy * imageColumns]);
+ value += filter[i + j] * ReadFloat4(image, number_channels, imageColumns, srcx, srcy, channel);
}
ix += 8;
i += 8;
while (i < width) {
int srcx = (ix < 0) ? -ix : ix; // mirror pad
srcx = (srcx >= imageColumns) ? (2 * imageColumns - srcx - 1) : srcx;
- value += filter[i] * convert_float4(im[srcx + srcy * imageColumns]);
+ value += filter[i] * ReadFloat4(image, number_channels, imageColumns, srcx, srcy, channel);
++i;
++ix;
}
float4 value = (float4)(0.0f);
int i = 0;
- while (i + 7 < width) { // unrolled
- value += (float4)(filter[i]) * pixels[px + (py + i) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 1) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 2) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 3) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 4) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 5) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 6) * prp];
- value += (float4)(filter[i]) * pixels[px + (py + i + 7) * prp];
+ while (i + 7 < width) {
+ for (int j = 0; j < 8; ++j) // unrolled
+ value += (float4)(filter[i]) * pixels[px + (py + i + j) * prp];
i += 8;
}
while (i < width) {
}
if (justBlur == 0) { // apply sharpening
- float4 srcPixel = convert_float4(im[x + y * imageColumns]);
+ float4 srcPixel = ReadFloat4(image, number_channels, imageColumns, x, y, channel);
float4 diff = srcPixel - value;
float quantumThreshold = QuantumRange*threshold;
}
if ((x < imageColumns) && (y < imageRows))
- filtered_im[x + y * imageColumns] = (CLPixelType)(ClampToQuantum(value.s0), ClampToQuantum(value.s1), ClampToQuantum(value.s2), ClampToQuantum(value.s3));
+ WriteFloat4(filteredImage, number_channels, imageColumns, x, y, channel, value);
}
)
STRINGIFY(
- __kernel __attribute__((reqd_work_group_size(64, 4, 1)))
- void WaveletDenoise(__global CLQuantum *srcImage,__global CLQuantum *dstImage,
- const unsigned int number_channels,const unsigned int max_channels,
- const float threshold,const int passes,const unsigned int imageWidth,
- const unsigned int imageHeight)
+ __kernel __attribute__((reqd_work_group_size(64, 4, 1)))
+ void WaveletDenoise(__global CLQuantum *srcImage,__global CLQuantum *dstImage,
+ const unsigned int number_channels,const unsigned int max_channels,
+ const float threshold,const int passes,const unsigned int imageWidth,
+ const unsigned int imageHeight)
{
const int pad = (1 << (passes - 1));
const int tileSize = 64;
return(MagickTrue);
}
+static cl_mem createKernelInfo(MagickCLEnv clEnv,cl_context context,
+ cl_command_queue queue,const double radius,const double sigma,cl_uint *width,
+ ExceptionInfo *exception)
+{
+ char
+ geometry[MagickPathExtent];
+
+ cl_int
+ status;
+
+ cl_mem
+ imageKernelBuffer;
+
+ float
+ *kernelBufferPtr;
+
+ KernelInfo
+ *kernel;
+
+ size_t
+ i;
+
+ (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.",".");
+ return((cl_mem) NULL);
+ }
+
+ imageKernelBuffer=clEnv->library->clCreateBuffer(context,CL_MEM_READ_ONLY,
+ kernel->width*sizeof(float),NULL,&status);
+ if (status != CL_SUCCESS)
+ {
+ kernel=DestroyKernelInfo(kernel);
+ (void) OpenCLThrowMagickException(exception,GetMagickModule(),
+ ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.",".");
+ return((cl_mem) NULL);
+ }
+
+ kernelBufferPtr=(float*)clEnv->library->clEnqueueMapBuffer(queue,
+ imageKernelBuffer,CL_TRUE,CL_MAP_WRITE,0,kernel->width*sizeof(float),0,
+ NULL,NULL,&status);
+ if (status != CL_SUCCESS)
+ {
+ kernel=DestroyKernelInfo(kernel);
+ clEnv->library->clReleaseMemObject(imageKernelBuffer);
+ (void) OpenCLThrowMagickException(exception,GetMagickModule(),
+ ResourceLimitWarning,"clEnv->library->clEnqueueMapBuffer failed.",".");
+ return((cl_mem) NULL);
+ }
+ for (i = 0; i < kernel->width; i++)
+ {
+ kernelBufferPtr[i]=(float)kernel->values[i];
+ }
+
+ *width=(cl_uint) kernel->width;
+ kernel=DestroyKernelInfo(kernel);
+
+ status=clEnv->library->clEnqueueUnmapMemObject(queue,imageKernelBuffer,
+ kernelBufferPtr,0,NULL,NULL);
+ if (status != CL_SUCCESS)
+ {
+ clEnv->library->clReleaseMemObject(imageKernelBuffer);
+ (void) OpenCLThrowMagickException(exception,GetMagickModule(),
+ ResourceLimitWarning,"clEnv->library->clEnqueueUnmapMemObject failed.",
+ "'%s'",".");
+ return((cl_mem) NULL);
+ }
+ return(imageKernelBuffer);
+}
+
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
% %
imageBuffer,
parametersBuffer;
- cl_mem_flags
- mem_flags;
-
cl_uint
number_channels;
MagickCLEnv
clEnv;
- MagickSizeType
- length;
-
size_t
globalWorkSize[2];
MagickCLEnv
clEnv;
- MagickSizeType
- length;
-
register ssize_t
i;
static Image *ComputeUnsharpMaskImageSingle(const Image *image,
const double radius,const double sigma,const double gain,
- const double threshold,int blurOnly, ExceptionInfo *exception)
+ const double threshold,int blurOnly,ExceptionInfo *exception)
{
CacheView
*filteredImage_view,
*image_view;
- char
- geometry[MagickPathExtent];
-
cl_command_queue
queue;
imageBuffer,
imageKernelBuffer;
- cl_mem_flags
- mem_flags;
-
- const void
- *inputPixels;
+ cl_uint
+ i,
+ imageColumns,
+ imageRows,
+ kernelWidth,
+ number_channels;
float
fGain,
- fThreshold,
- *kernelBufferPtr;
+ fThreshold;
Image
*filteredImage;
- KernelInfo
- *kernel;
-
MagickBooleanType
outputReady;
MagickCLEnv
clEnv;
- MagickSizeType
- length;
-
void
- *filteredPixels,
- *hostPtr;
-
- unsigned int
- i,
- imageColumns,
- imageRows,
- kernelWidth;
+ *filteredPixels;
clEnv = NULL;
filteredImage = NULL;
filteredImage_view = NULL;
- kernel = NULL;
+ filteredPixels = NULL;
context = NULL;
imageBuffer = NULL;
filteredImageBuffer = 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;
- }
- }
-
- /* create output */
- {
- 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;
- }
- }
-
- /* 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;
- }
+ image_view=AcquireVirtualCacheView(image,exception);
+ imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception);
+ if (imageBuffer == (cl_mem) NULL)
+ 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;
- }
+ filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
+ if (filteredImage == (Image *) NULL)
+ goto cleanup;
+ filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
+ filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
+ context,filteredPixels,exception);
+ if (filteredImageBuffer == (void *) NULL)
+ 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);
{
/* get the opencl kernel */
}
{
- imageColumns = (unsigned int) image->columns;
- imageRows = (unsigned int) image->rows;
- kernelWidth = (unsigned int) kernel->width;
+ imageColumns = (cl_uint) image->columns;
+ imageRows = (cl_uint) image->rows;
+ number_channels = (cl_uint) image->number_channels;
fGain = (float) gain;
fThreshold = (float) threshold;
justBlur = blurOnly;
/* set the kernel arguments */
i = 0;
clStatus=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&number_channels);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&kernelWidth);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&imageColumns);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(unsigned int),(void *)&imageRows);
- clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernel->width)),(void *) NULL);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernelWidth)),(void *) NULL);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold);
clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&justBlur);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
wsize[0] = 8;
wsize[1] = 32;
- clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
+ clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 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)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+ if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
goto cleanup;
- }
outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
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 (imageKernelBuffer!=NULL) clEnv->library->clReleaseMemObject(imageKernelBuffer);
projects / imagemagick / commitdiff