*/
STRINGIFY(
- __kernel void RotationalBlur(const __global CLPixelType *im, __global CLPixelType *filtered_im,
- const float4 bias,
- const unsigned int channel, const unsigned int matte,
- const float2 blurCenter,
- __constant float *cos_theta, __constant float *sin_theta,
- const unsigned int cossin_theta_size)
+ __kernel void RotationalBlur(const __global CLQuantum *image,const unsigned int number_channels,
+ const unsigned int channel,const float4 bias,const float2 blurCenter,__constant float *cos_theta,
+ __constant float *sin_theta,const unsigned int cossin_theta_size,__global CLQuantum *filteredImage)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
step = cossin_theta_size-1;
}
- float4 result;
- result.x = (float)bias.x;
- result.y = (float)bias.y;
- result.z = (float)bias.z;
- result.w = (float)bias.w;
+ float4 result = bias;
+
float normalize = 0.0f;
+ float gamma = 0.0f;
+
+ for (unsigned int i=0; i<cossin_theta_size; i+=step)
+ {
+ int cx = ClampToCanvas(blurCenter.x+center_x*cos_theta[i]-center_y*sin_theta[i]+0.5f,columns);
+ int cy = ClampToCanvas(blurCenter.y+center_x*sin_theta[i]+center_y*cos_theta[i]+0.5f,rows);
+
+ float4 pixel = ReadFloat4(image, number_channels, columns, cx, cy, channel);
- if (((channel & AlphaChannel) == 0) || (matte == 0)) {
- for (unsigned int i=0; i<cossin_theta_size; i+=step)
+ if ((number_channels == 4) || (number_channels == 2))
{
- result += convert_float4(im[
- ClampToCanvas(blurCenter.x+center_x*cos_theta[i]-center_y*sin_theta[i]+0.5f,columns)+
- ClampToCanvas(blurCenter.y+center_x*sin_theta[i]+center_y*cos_theta[i]+0.5f, rows)*columns]);
- normalize += 1.0f;
+ float alpha = (float)(QuantumScale*pixel.w);
+
+ gamma += alpha;
+
+ result.x += alpha * pixel.x;
+ result.y += alpha * pixel.y;
+ result.z += alpha * pixel.z;
+ result.w += pixel.w;
}
- normalize = PerceptibleReciprocal(normalize);
- result = result * normalize;
+ else
+ result += pixel;
+
+ normalize += 1.0f;
}
- else {
- float gamma = 0.0f;
- for (unsigned int i=0; i<cossin_theta_size; i+=step)
- {
- float4 p = convert_float4(im[
- ClampToCanvas(blurCenter.x+center_x*cos_theta[i]-center_y*sin_theta[i]+0.5f,columns)+
- ClampToCanvas(blurCenter.y+center_x*sin_theta[i]+center_y*cos_theta[i]+0.5f, rows)*columns]);
-
- float alpha = (float)(QuantumScale*p.w);
- result.x += alpha * p.x;
- result.y += alpha * p.y;
- result.z += alpha * p.z;
- result.w += p.w;
- gamma+=alpha;
- normalize += 1.0f;
- }
+
+ normalize = PerceptibleReciprocal(normalize);
+
+ if ((number_channels == 4) || (number_channels == 2))
+ {
gamma = PerceptibleReciprocal(gamma);
- normalize = PerceptibleReciprocal(normalize);
- result.x = gamma*result.x;
- result.y = gamma*result.y;
- result.z = gamma*result.z;
- result.w = normalize*result.w;
+ result.x *= gamma;
+ result.y *= gamma;
+ result.z *= gamma;
+ result.w *= normalize;
}
- filtered_im[y * columns + x] = (CLPixelType) (ClampToQuantum(result.x), ClampToQuantum(result.y),
- ClampToQuantum(result.z), ClampToQuantum(result.w));
+ else
+ result *= normalize;
+
+ WriteFloat4(filteredImage, number_channels, columns, x, y, channel, result);
}
)
imageBuffer,
sinThetaBuffer;
- cl_mem_flags
- mem_flags;
-
cl_kernel
rotationalBlurKernel;
cl_event
event;
- const void
- *inputPixels;
+ cl_uint
+ cossin_theta_size,
+ number_channels;
float
blurRadius,
PixelInfo
bias;
- MagickSizeType
- length;
-
size_t
global_work_size[2];
unsigned int
- cossin_theta_size,
- i,
- matte;
+ i;
void
- *filteredPixels,
- *hostPtr;
+ *filteredPixels;
outputReady = MagickFalse;
context = NULL;
filteredImage = NULL;
filteredImage_view = NULL;
+ filteredPixels = NULL;
imageBuffer = NULL;
filteredImageBuffer = NULL;
sinThetaBuffer = NULL;
queue = NULL;
rotationalBlurKernel = NULL;
-
clEnv = GetDefaultOpenCLEnv();
context = GetOpenCLContext(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.",".");
+ image_view=AcquireVirtualCacheView(image, exception);
+ imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception);
+ if (imageBuffer == (cl_mem) NULL)
goto cleanup;
- }
-
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'", ".");
+ if (filteredImage == (Image *) NULL)
goto cleanup;
- }
- filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
- filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception);
- if (filteredPixels == (void *) NULL)
+ if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue)
{
- (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename);
+ (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
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.",".");
+ filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception);
+ filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
+ context,filteredPixels,exception);
+ if (filteredImageBuffer == (cl_mem) NULL)
goto cleanup;
- }
blurCenter.s[0] = (float) (image->columns-1)/2.0;
blurCenter.s[1] = (float) (image->rows-1)/2.0;
goto cleanup;
}
-
- /* set the kernel arguments */
- i = 0;
- clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
- clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
-
GetPixelInfo(image,&bias);
biasPixel.s[0] = bias.red;
biasPixel.s[1] = bias.green;
biasPixel.s[2] = bias.blue;
biasPixel.s[3] = bias.alpha;
- clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel);
- clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &image->channel_mask);
-
- matte = (image->alpha_trait > CopyPixelTrait)?1:0;
- clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &matte);
- clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter);
+ number_channels = image->number_channels;
+ /* set the kernel arguments */
+ i = 0;
+ clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+ clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_uint),&number_channels);
+ clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &image->channel_mask);
+ clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel);
+ clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter);
clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer);
clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer);
- clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size);
+ clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_uint), &cossin_theta_size);
+ clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
if (clStatus != CL_SUCCESS)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
goto cleanup;
}
-
global_work_size[0] = image->columns;
global_work_size[1] = image->rows;
/* launch the kernel */
RecordProfileData(clEnv,RotationalBlurKernel,event);
clEnv->library->clReleaseEvent(event);
- 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(filteredImage,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;
}
+
outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
cleanup:
assert(image != NULL);
assert(exception != (ExceptionInfo *) NULL);
- if ((checkAccelerateConditionRGBA(image) == MagickFalse) ||
+ if ((checkAccelerateCondition(image) == MagickFalse) ||
(checkOpenCLEnvironment(exception) == MagickFalse))
return NULL;
clEnv->library->clReleaseEvent(event);
}
- if (copyWriteBuffer(image,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
+ if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
{
(void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
goto cleanup;