return filteredImage;
}
+static Image *ComputeUnsharpMaskImageSingle(const Image *image,
+ const ChannelType channel,const double radius,const double sigma,
+ const double gain,const double threshold,int blurOnly, ExceptionInfo *exception);
+
+static Image* ComputeBlurImageSingle(const Image* image,
+ const ChannelType channel,const double radius,const double sigma,
+ ExceptionInfo *exception)
+{
+ return ComputeUnsharpMaskImageSingle(image, channel, radius, sigma, 0.0, 0.0, 1, exception);
+}
+
MagickExport Image* AccelerateBlurImage(const Image *image,
const ChannelType channel,const double radius,const double sigma,
ExceptionInfo *exception)
(checkAccelerateCondition(image, channel) == MagickFalse))
return NULL;
- if (splitImage(image) && (image->rows / 2 > radius))
+ if (radius < 12.1)
+ filteredImage=ComputeBlurImageSingle(image, channel, radius, sigma, exception);
+ else if (splitImage(image) && (image->rows / 2 > radius))
filteredImage=ComputeBlurImageSection(image, channel, radius, sigma, exception);
else
filteredImage=ComputeBlurImage(image, channel, radius, sigma, exception);
return filteredImage;
}
-MagickExport Image *AccelerateUnsharpMaskImage(const Image *image,
+static Image *ComputeUnsharpMaskImageSingle(const Image *image,
const ChannelType channel,const double radius,const double sigma,
- const double gain,const double threshold,ExceptionInfo *exception)
+ const double gain,const double threshold,int blurOnly, ExceptionInfo *exception)
{
- Image
- *filteredImage;
+ CacheView
+ *filteredImage_view,
+ *image_view;
- assert(image != NULL);
- assert(exception != (ExceptionInfo *) NULL);
+ char
+ geometry[MaxTextExtent];
- if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
- (checkAccelerateCondition(image, channel) == MagickFalse))
- return NULL;
+ cl_command_queue
+ queue;
- if (splitImage(image) && (image->rows / 2 > radius))
- filteredImage = ComputeUnsharpMaskImageSection(image,channel,radius,sigma,gain,threshold,exception);
- else
- filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
- return(filteredImage);
-}
+ cl_context
+ context;
-/*
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-% %
-% %
-% %
-% A c c e l e r a t e R e s i z e I m a g e %
-% %
-% %
-% %
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-%
-% AccelerateResizeImage() is an OpenCL implementation of ResizeImage()
-%
-% AccelerateResizeImage() scales an image to the desired dimensions, using the given
-% filter (see AcquireFilterInfo()).
-%
-% If an undefined filter is given the filter defaults to Mitchell for a
-% colormapped image, a image with a matte channel, or if the image is
-% enlarged. Otherwise the filter defaults to a Lanczos.
-%
-% AccelerateResizeImage() was inspired by Paul Heckbert's "zoom" program.
-%
-% The format of the AccelerateResizeImage method is:
-%
-% Image *ResizeImage(Image *image,const size_t columns,
-% const size_t rows, const ResizeFilter* filter,
-% ExceptionInfo *exception)
-%
-% A description of each parameter follows:
-%
-% o image: the image.
-%
-% o columns: the number of columns in the scaled image.
-%
-% o rows: the number of rows in the scaled image.
-%
-% o filter: Image filter to use.
-%
-% o exception: return any errors or warnings in this structure.
-%
-*/
+ cl_int
+ justBlur,
+ clStatus;
-static MagickBooleanType resizeHorizontalFilter(cl_mem image,
- const unsigned int imageColumns,const unsigned int imageRows,
- const unsigned int matte,cl_mem resizedImage,
- const unsigned int resizedColumns,const unsigned int resizedRows,
- const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
- const float xFactor,MagickCLEnv clEnv,cl_command_queue queue,
- ExceptionInfo *exception)
-{
cl_kernel
- horizontalKernel;
+ unsharpMaskKernel;
- cl_int clStatus;
+ cl_mem
+ filteredImageBuffer,
+ imageBuffer,
+ imageKernelBuffer;
- const unsigned int
- workgroupSize = 256;
+ cl_mem_flags
+ mem_flags;
+
+ const void
+ *inputPixels;
float
- resizeFilterScale,
- resizeFilterSupport,
- resizeFilterWindowSupport,
- resizeFilterBlur,
- scale,
- support;
+ fGain,
+ fThreshold,
+ *kernelBufferPtr;
- int
- cacheRangeStart,
- cacheRangeEnd,
- numCachedPixels,
- resizeFilterType,
- resizeWindowType;
+ Image
+ *filteredImage;
+
+ KernelInfo
+ *kernel;
MagickBooleanType
- status = MagickFalse;
+ outputReady;
- size_t
- deviceLocalMemorySize,
- gammaAccumulatorLocalMemorySize,
- global_work_size[2],
- imageCacheLocalMemorySize,
- pixelAccumulatorLocalMemorySize,
- local_work_size[2],
- totalLocalMemorySize,
- weightAccumulatorLocalMemorySize;
+ MagickCLEnv
+ clEnv;
- unsigned int
- chunkSize,
- i,
- pixelPerWorkgroup;
+ MagickSizeType
+ length;
- horizontalKernel = NULL;
- status = MagickFalse;
+ void
+ *filteredPixels,
+ *hostPtr;
- /*
- Apply filter to resize vertically from image to resize image.
- */
- scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
- support=scale*GetResizeFilterSupport(resizeFilter);
- if (support < 0.5)
- {
- /*
- Support too small even for nearest neighbour: Reduce to point
- sampling.
- */
- support=(MagickRealType) 0.5;
- scale=1.0;
- }
- scale=PerceptibleReciprocal(scale);
+ unsigned int
+ i,
+ imageColumns,
+ imageRows,
+ kernelWidth;
- if (resizedColumns < workgroupSize)
- {
- chunkSize = 32;
- pixelPerWorkgroup = 32;
- }
- else
- {
- chunkSize = workgroupSize;
- pixelPerWorkgroup = workgroupSize;
- }
+ clEnv = NULL;
+ filteredImage = NULL;
+ filteredImage_view = NULL;
+ kernel = NULL;
+ context = NULL;
+ imageBuffer = NULL;
+ filteredImageBuffer = NULL;
+ imageKernelBuffer = NULL;
+ unsharpMaskKernel = NULL;
+ queue = NULL;
+ outputReady = MagickFalse;
- /* get the local memory size supported by the device */
- deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
+ clEnv = GetDefaultOpenCLEnv();
+ context = GetOpenCLContext(clEnv);
+ queue = AcquireOpenCLCommandQueue(clEnv);
-DisableMSCWarning(4127)
- while(1)
-RestoreMSCWarning
+ /* Create and initialize OpenCL buffers. */
{
- /* calculate the local memory size needed per workgroup */
- cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
- cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
- numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
- imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
- totalLocalMemorySize = imageCacheLocalMemorySize;
-
- /* local size for the pixel accumulator */
- pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
- totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
-
- /* local memory size for the weight accumulator */
- weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
- totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
-
- /* local memory size for the gamma accumulator */
- if (matte == 0)
- gammaAccumulatorLocalMemorySize = sizeof(float);
- else
- gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
- totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
+ 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 (totalLocalMemorySize <= deviceLocalMemorySize)
- break;
- else
+ /* 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))
{
- pixelPerWorkgroup = pixelPerWorkgroup/2;
- chunkSize = chunkSize/2;
- if (pixelPerWorkgroup == 0
- || chunkSize == 0)
- {
- /* quit, fallback to CPU */
- goto cleanup;
- }
+ 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;
}
}
- resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
- resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
-
-
- if (resizeFilterType == SincFastWeightingFunction
- && resizeWindowType == SincFastWeightingFunction)
- {
- horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilterSinc");
- }
- else
+ /* create output */
{
- horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter");
- }
- if (horizontalKernel == NULL)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
- 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'", ".");
+ 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;
+ }
}
- i = 0;
- clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&image);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
+ /* create the blur kernel */
+ {
+ (void) FormatLocaleString(geometry,MaxTextExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
+ kernel=AcquireKernelInfo(geometry);
+ if (kernel == (KernelInfo *) NULL)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.",".");
+ goto cleanup;
+ }
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
+ 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;
+ }
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
- resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
+ 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;
+ }
+ }
- resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
+ {
+ /* get the opencl kernel */
+ {
+ unsharpMaskKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMask");
+ if (unsharpMaskKernel == NULL)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+ goto cleanup;
+ };
+ }
- resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
+ {
+ imageColumns = image->columns;
+ imageRows = image->rows;
+ kernelWidth = kernel->width;
+ fGain = (float)gain;
+ fThreshold = (float)threshold;
+ justBlur = blurOnly;
- resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
+ /* 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_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(float),(void *)&fGain);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold);
+ clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&justBlur);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ }
+ /* launch the kernel */
+ {
+ size_t gsize[2];
+ size_t wsize[2];
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
-
+ gsize[0] = ((image->columns + 7) / 8) * 8;
+ gsize[1] = ((image->rows + 31) / 32) * 32;
+ wsize[0] = 8;
+ wsize[1] = 32;
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
- clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
+ clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ clEnv->library->clFlush(queue);
+ }
+ }
- if (clStatus != CL_SUCCESS)
+ /* get result */
+ if (ALIGNED(filteredPixels,CLPixelPacket))
{
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
+ 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);
}
-
- global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
- global_work_size[1] = resizedRows;
-
- local_work_size[0] = workgroupSize;
- local_work_size[1] = 1;
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
if (clStatus != CL_SUCCESS)
{
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
goto cleanup;
}
- clEnv->library->clFlush(queue);
- status = MagickTrue;
+ outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
cleanup:
OpenCLLogException(__FUNCTION__,__LINE__,exception);
- if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
+ image_view=DestroyCacheView(image_view);
+ if (filteredImage_view != NULL)
+ filteredImage_view=DestroyCacheView(filteredImage_view);
- return(status);
+ 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);
+ if (unsharpMaskKernel!=NULL) RelinquishOpenCLKernel(clEnv, unsharpMaskKernel);
+ if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
+ if (outputReady == MagickFalse)
+ {
+ if (filteredImage != NULL)
+ {
+ DestroyImage(filteredImage);
+ filteredImage = NULL;
+ }
+ }
+ return(filteredImage);
}
-static MagickBooleanType resizeVerticalFilter(cl_mem image,
+
+MagickExport Image *AccelerateUnsharpMaskImage(const Image *image,
+ const ChannelType channel,const double radius,const double sigma,
+ const double gain,const double threshold,ExceptionInfo *exception)
+{
+ Image
+ *filteredImage;
+
+ assert(image != NULL);
+ assert(exception != (ExceptionInfo *) NULL);
+
+ if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
+ (checkAccelerateCondition(image, channel) == MagickFalse))
+ return NULL;
+
+ if (radius < 12.1)
+ filteredImage = ComputeUnsharpMaskImageSingle(image,channel,radius,sigma,gain,threshold, 0, exception);
+ else if (splitImage(image) && (image->rows / 2 > radius))
+ filteredImage = ComputeUnsharpMaskImageSection(image,channel,radius,sigma,gain,threshold,exception);
+ else
+ filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception);
+ return(filteredImage);
+}
+
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+% %
+% %
+% %
+% A c c e l e r a t e R e s i z e I m a g e %
+% %
+% %
+% %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+% AccelerateResizeImage() is an OpenCL implementation of ResizeImage()
+%
+% AccelerateResizeImage() scales an image to the desired dimensions, using the given
+% filter (see AcquireFilterInfo()).
+%
+% If an undefined filter is given the filter defaults to Mitchell for a
+% colormapped image, a image with a matte channel, or if the image is
+% enlarged. Otherwise the filter defaults to a Lanczos.
+%
+% AccelerateResizeImage() was inspired by Paul Heckbert's "zoom" program.
+%
+% The format of the AccelerateResizeImage method is:
+%
+% Image *ResizeImage(Image *image,const size_t columns,
+% const size_t rows, const ResizeFilter* filter,
+% ExceptionInfo *exception)
+%
+% A description of each parameter follows:
+%
+% o image: the image.
+%
+% o columns: the number of columns in the scaled image.
+%
+% o rows: the number of rows in the scaled image.
+%
+% o filter: Image filter to use.
+%
+% o exception: return any errors or warnings in this structure.
+%
+*/
+
+static MagickBooleanType resizeHorizontalFilter(cl_mem image,
const unsigned int imageColumns,const unsigned int imageRows,
const unsigned int matte,cl_mem resizedImage,
const unsigned int resizedColumns,const unsigned int resizedRows,
const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
- const float yFactor,MagickCLEnv clEnv,cl_command_queue queue,
+ const float xFactor,MagickCLEnv clEnv,cl_command_queue queue,
ExceptionInfo *exception)
{
cl_kernel
/*
Apply filter to resize vertically from image to resize image.
*/
- scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
+ scale=MAGICK_MAX(1.0/xFactor+MagickEpsilon,1.0);
support=scale*GetResizeFilterSupport(resizeFilter);
if (support < 0.5)
{
}
scale=PerceptibleReciprocal(scale);
- if (resizedRows < workgroupSize)
+ if (resizedColumns < workgroupSize)
{
chunkSize = 32;
pixelPerWorkgroup = 32;
RestoreMSCWarning
{
/* calculate the local memory size needed per workgroup */
- cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
- cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
+ cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
+ cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
totalLocalMemorySize = imageCacheLocalMemorySize;
resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
+
if (resizeFilterType == SincFastWeightingFunction
&& resizeWindowType == SincFastWeightingFunction)
- horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilterSinc");
- else
- horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter");
-
- if (horizontalKernel == NULL)
+ {
+ horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilterSinc");
+ }
+ else
+ {
+ horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeHorizontalFilter");
+ }
+ if (horizontalKernel == NULL)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+
+ i = 0;
+ clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&image);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageColumns);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&imageRows);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&matte);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
+
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedColumns);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), (void*)&resizedRows);
+
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
+
+ resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
+
+ resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
+
+ resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
+
+ resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
+
+
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
+
+
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
+ clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
+
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
+ }
+
+ global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
+ global_work_size[1] = resizedRows;
+
+ local_work_size[0] = workgroupSize;
+ local_work_size[1] = 1;
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ clEnv->library->clFlush(queue);
+ status = MagickTrue;
+
+
+cleanup:
+ OpenCLLogException(__FUNCTION__,__LINE__,exception);
+
+ if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel);
+
+ return(status);
+}
+
+static MagickBooleanType resizeVerticalFilter(cl_mem image,
+ const unsigned int imageColumns,const unsigned int imageRows,
+ const unsigned int matte,cl_mem resizedImage,
+ const unsigned int resizedColumns,const unsigned int resizedRows,
+ const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
+ const float yFactor,MagickCLEnv clEnv,cl_command_queue queue,
+ ExceptionInfo *exception)
+{
+ cl_kernel
+ horizontalKernel;
+
+ cl_int clStatus;
+
+ const unsigned int
+ workgroupSize = 256;
+
+ float
+ resizeFilterScale,
+ resizeFilterSupport,
+ resizeFilterWindowSupport,
+ resizeFilterBlur,
+ scale,
+ support;
+
+ int
+ cacheRangeStart,
+ cacheRangeEnd,
+ numCachedPixels,
+ resizeFilterType,
+ resizeWindowType;
+
+ MagickBooleanType
+ status = MagickFalse;
+
+ size_t
+ deviceLocalMemorySize,
+ gammaAccumulatorLocalMemorySize,
+ global_work_size[2],
+ imageCacheLocalMemorySize,
+ pixelAccumulatorLocalMemorySize,
+ local_work_size[2],
+ totalLocalMemorySize,
+ weightAccumulatorLocalMemorySize;
+
+ unsigned int
+ chunkSize,
+ i,
+ pixelPerWorkgroup;
+
+ horizontalKernel = NULL;
+ status = MagickFalse;
+
+ /*
+ Apply filter to resize vertically from image to resize image.
+ */
+ scale=MAGICK_MAX(1.0/yFactor+MagickEpsilon,1.0);
+ support=scale*GetResizeFilterSupport(resizeFilter);
+ if (support < 0.5)
+ {
+ /*
+ Support too small even for nearest neighbour: Reduce to point
+ sampling.
+ */
+ support=(MagickRealType) 0.5;
+ scale=1.0;
+ }
+ scale=PerceptibleReciprocal(scale);
+
+ if (resizedRows < workgroupSize)
+ {
+ chunkSize = 32;
+ pixelPerWorkgroup = 32;
+ }
+ else
+ {
+ chunkSize = workgroupSize;
+ pixelPerWorkgroup = workgroupSize;
+ }
+
+ /* get the local memory size supported by the device */
+ deviceLocalMemorySize = GetOpenCLDeviceLocalMemorySize(clEnv);
+
+DisableMSCWarning(4127)
+ while(1)
+RestoreMSCWarning
+ {
+ /* calculate the local memory size needed per workgroup */
+ cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
+ cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
+ numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
+ imageCacheLocalMemorySize = numCachedPixels * sizeof(CLPixelPacket);
+ totalLocalMemorySize = imageCacheLocalMemorySize;
+
+ /* local size for the pixel accumulator */
+ pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
+ totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
+
+ /* local memory size for the weight accumulator */
+ weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
+ totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
+
+ /* local memory size for the gamma accumulator */
+ if (matte == 0)
+ gammaAccumulatorLocalMemorySize = sizeof(float);
+ else
+ gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
+ totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
+
+ if (totalLocalMemorySize <= deviceLocalMemorySize)
+ break;
+ else
+ {
+ pixelPerWorkgroup = pixelPerWorkgroup/2;
+ chunkSize = chunkSize/2;
+ if (pixelPerWorkgroup == 0
+ || chunkSize == 0)
+ {
+ /* quit, fallback to CPU */
+ goto cleanup;
+ }
+ }
+ }
+
+ resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
+ resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
+
+ if (resizeFilterType == SincFastWeightingFunction
+ && resizeWindowType == SincFastWeightingFunction)
+ horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilterSinc");
+ else
+ horizontalKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ResizeVerticalFilter");
+
+ if (horizontalKernel == NULL)
{
(void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
goto cleanup;
}
else
{
- mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
- hostPtr = stretch_map;
- }
- /* create a CL buffer for stretch_map */
- length = (MaxMap+1);
- stretchMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
- goto cleanup;
- }
-
- /* get the OpenCL kernel */
- stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Stretch");
- if (stretchKernel == NULL)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
- goto cleanup;
- }
-
- /* set the kernel arguments */
- i = 0;
- clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
- clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(ChannelType),&channel);
- clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&stretchMapBuffer);
- clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&white);
- clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
- }
-
- /* launch the kernel */
- global_work_size[0] = image->columns;
- global_work_size[1] = image->rows;
-
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
-
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
- clEnv->library->clFlush(queue);
-
- /* read the data back */
- if (ALIGNED(inputPixels,CLPixelPacket))
- {
- 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);
- }
- else
- {
- length = image->columns * image->rows;
- clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
- }
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
- goto cleanup;
- }
-
- outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
-
-cleanup:
- OpenCLLogException(__FUNCTION__,__LINE__,exception);
-
- image_view=DestroyCacheView(image_view);
-
- if (imageBuffer!=NULL)
- clEnv->library->clReleaseMemObject(imageBuffer);
-
- if (stretchMapBuffer!=NULL)
- clEnv->library->clReleaseMemObject(stretchMapBuffer);
- if (stretch_map!=NULL)
- stretch_map=(PixelPacket *) RelinquishMagickMemory(stretch_map);
-
-
- if (histogramBuffer!=NULL)
- clEnv->library->clReleaseMemObject(histogramBuffer);
- if (histogram!=NULL)
- histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
-
-
- if (histogramKernel!=NULL)
- RelinquishOpenCLKernel(clEnv, histogramKernel);
- if (stretchKernel!=NULL)
- RelinquishOpenCLKernel(clEnv, stretchKernel);
-
- if (queue != NULL)
- RelinquishOpenCLCommandQueue(clEnv, queue);
-
- return(outputReady);
-}
-
-MagickExport MagickBooleanType AccelerateContrastStretchImageChannel(
- Image *image,const ChannelType channel,const double black_point,
- const double white_point,ExceptionInfo *exception)
-{
- MagickBooleanType
- status;
-
- assert(image != NULL);
- assert(exception != (ExceptionInfo *) NULL);
-
- if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
- (checkAccelerateCondition(image, channel) == MagickFalse) ||
- (checkHistogramCondition(image, channel) == MagickFalse))
- return(MagickFalse);
-
- status=ComputeContrastStretchImageChannel(image,channel, black_point, white_point, exception);
- return(status);
-}
-
-/*
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-% %
-% %
-% %
-% D e s p e c k l e I m a g e w i t h O p e n C L %
-% %
-% %
-% %
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-%
-% DespeckleImage() reduces the speckle noise in an image while perserving the
-% edges of the original image. A speckle removing filter uses a complementary
-% hulling technique (raising pixels that are darker than their surrounding
-% neighbors, then complementarily lowering pixels that are brighter than their
-% surrounding neighbors) to reduce the speckle index of that image (reference
-% Crimmins speckle removal).
-%
-% The format of the DespeckleImage method is:
-%
-% Image *DespeckleImage(const Image *image,ExceptionInfo *exception)
-%
-% A description of each parameter follows:
-%
-% o image: the image.
-%
-% o exception: return any errors or warnings in this structure.
-%
-*/
-
-static Image *ComputeDespeckleImage(const Image *image,
- ExceptionInfo*exception)
-{
- static const int
- X[4] = {0, 1, 1,-1},
- Y[4] = {1, 0, 1, 1};
-
- CacheView
- *filteredImage_view,
- *image_view;
-
- cl_command_queue
- queue;
-
- cl_context
- context;
-
- cl_int
- clStatus;
-
- cl_kernel
- hullPass1,
- hullPass2;
-
- cl_mem_flags
- mem_flags;
-
- cl_mem
- filteredImageBuffer,
- imageBuffer,
- tempImageBuffer[2];
-
- const void
- *inputPixels;
-
- Image
- *filteredImage;
-
- int
- k,
- matte;
-
- MagickBooleanType
- outputReady;
-
- MagickCLEnv
- clEnv;
-
- MagickSizeType
- length;
-
- size_t
- global_work_size[2];
-
- unsigned int
- imageHeight,
- imageWidth;
-
- void
- *filteredPixels,
- *hostPtr;
-
- outputReady = MagickFalse;
- clEnv = NULL;
- inputPixels = NULL;
- filteredImage = NULL;
- filteredImage_view = NULL;
- filteredPixels = NULL;
- context = NULL;
- imageBuffer = NULL;
- filteredImageBuffer = NULL;
- hullPass1 = NULL;
- hullPass2 = NULL;
- queue = NULL;
- tempImageBuffer[0] = tempImageBuffer[1] = NULL;
- clEnv = GetDefaultOpenCLEnv();
- context = GetOpenCLContext(clEnv);
- queue = AcquireOpenCLCommandQueue(clEnv);
-
- image_view=AcquireVirtualCacheView(image,exception);
- inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception);
- if (inputPixels == (void *) NULL)
- {
- (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
- goto cleanup;
- }
-
- 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;
- }
-
- mem_flags = CL_MEM_READ_WRITE;
- length = image->columns * image->rows;
- for (k = 0; k < 2; k++)
- {
- tempImageBuffer[k] = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), 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);
- 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;
- }
-
- hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1");
- hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2");
-
- clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&imageBuffer);
- clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1));
- imageWidth = image->columns;
- clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth);
- imageHeight = image->rows;
- clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight);
- matte = (image->alpha_trait==BlendPixelTrait)?0:1;
- clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
- }
-
- clStatus = clEnv->library->clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1));
- clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer);
- imageWidth = image->columns;
- clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth);
- imageHeight = image->rows;
- clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight);
- matte = (image->alpha_trait==BlendPixelTrait)?0:1;
- clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
- 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;
-
-
- for (k = 0; k < 4; k++)
- {
- cl_int2 offset;
- int polarity;
-
-
- offset.s[0] = X[k];
- offset.s[1] = Y[k];
- polarity = 1;
- clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
- clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
-
-
- if (k == 0)
- clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
- offset.s[0] = -X[k];
- offset.s[1] = -Y[k];
- polarity = 1;
- clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
- clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
-
- offset.s[0] = -X[k];
- offset.s[1] = -Y[k];
- polarity = -1;
- clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
- clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
+ mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR;
+ hostPtr = stretch_map;
+ }
+ /* create a CL buffer for stretch_map */
+ length = (MaxMap+1);
+ stretchMapBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
+ goto cleanup;
+ }
- offset.s[0] = X[k];
- offset.s[1] = Y[k];
- polarity = -1;
- clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
- clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
- clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
+ /* get the OpenCL kernel */
+ stretchKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Stretch");
+ if (stretchKernel == NULL)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
- if (k == 3)
- clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer);
+ /* set the kernel arguments */
+ i = 0;
+ clStatus=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+ clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(ChannelType),&channel);
+ clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(cl_mem),(void *)&stretchMapBuffer);
+ clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&white);
+ clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
+ }
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
- /* launch the kernel */
- clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
- goto cleanup;
- }
+ /* launch the kernel */
+ global_work_size[0] = image->columns;
+ global_work_size[1] = image->rows;
+
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, stretchKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
}
+ clEnv->library->clFlush(queue);
- if (ALIGNED(filteredPixels,CLPixelPacket))
+ /* read the data back */
+ if (ALIGNED(inputPixels,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);
+ clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, 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);
+ clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL);
}
if (clStatus != CL_SUCCESS)
{
goto cleanup;
}
- outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
+ outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
cleanup:
OpenCLLogException(__FUNCTION__,__LINE__,exception);
image_view=DestroyCacheView(image_view);
- if (filteredImage_view != NULL)
- filteredImage_view=DestroyCacheView(filteredImage_view);
- if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
- if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
- for (k = 0; k < 2; k++)
- {
- if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
- }
- if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
- if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
- if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
- if (outputReady == MagickFalse && filteredImage != NULL)
- filteredImage=DestroyImage(filteredImage);
- return(filteredImage);
+ if (imageBuffer!=NULL)
+ clEnv->library->clReleaseMemObject(imageBuffer);
+
+ if (stretchMapBuffer!=NULL)
+ clEnv->library->clReleaseMemObject(stretchMapBuffer);
+ if (stretch_map!=NULL)
+ stretch_map=(PixelPacket *) RelinquishMagickMemory(stretch_map);
+
+
+ if (histogramBuffer!=NULL)
+ clEnv->library->clReleaseMemObject(histogramBuffer);
+ if (histogram!=NULL)
+ histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
+
+
+ if (histogramKernel!=NULL)
+ RelinquishOpenCLKernel(clEnv, histogramKernel);
+ if (stretchKernel!=NULL)
+ RelinquishOpenCLKernel(clEnv, stretchKernel);
+
+ if (queue != NULL)
+ RelinquishOpenCLCommandQueue(clEnv, queue);
+
+ return(outputReady);
}
-MagickExport Image *AccelerateDespeckleImage(const Image* image,
- ExceptionInfo* exception)
+MagickExport MagickBooleanType AccelerateContrastStretchImageChannel(
+ Image *image,const ChannelType channel,const double black_point,
+ const double white_point,ExceptionInfo *exception)
{
- Image
- *filteredImage;
+ MagickBooleanType
+ status;
assert(image != NULL);
assert(exception != (ExceptionInfo *) NULL);
if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
- (checkAccelerateCondition(image, AllChannels) == MagickFalse))
- return NULL;
+ (checkAccelerateCondition(image, channel) == MagickFalse) ||
+ (checkHistogramCondition(image, channel) == MagickFalse))
+ return(MagickFalse);
- filteredImage=ComputeDespeckleImage(image,exception);
- return(filteredImage);
+ status=ComputeContrastStretchImageChannel(image,channel, black_point, white_point, exception);
+ return(status);
}
-static Image *ComputeAddNoiseImage(const Image *image,
- const ChannelType channel,const NoiseType noise_type,
- ExceptionInfo *exception)
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+% %
+% %
+% %
+% D e s p e c k l e I m a g e w i t h O p e n C L %
+% %
+% %
+% %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+% DespeckleImage() reduces the speckle noise in an image while perserving the
+% edges of the original image. A speckle removing filter uses a complementary
+% hulling technique (raising pixels that are darker than their surrounding
+% neighbors, then complementarily lowering pixels that are brighter than their
+% surrounding neighbors) to reduce the speckle index of that image (reference
+% Crimmins speckle removal).
+%
+% The format of the DespeckleImage method is:
+%
+% Image *DespeckleImage(const Image *image,ExceptionInfo *exception)
+%
+% A description of each parameter follows:
+%
+% o image: the image.
+%
+% o exception: return any errors or warnings in this structure.
+%
+*/
+
+static Image *ComputeDespeckleImage(const Image *image,
+ ExceptionInfo*exception)
{
+ static const int
+ X[4] = {0, 1, 1,-1},
+ Y[4] = {1, 0, 1, 1};
+
CacheView
*filteredImage_view,
*image_view;
clStatus;
cl_kernel
- addNoiseKernel;
+ hullPass1,
+ hullPass2;
cl_mem_flags
mem_flags;
cl_mem
filteredImageBuffer,
imageBuffer,
- randomNumberBuffer;
-
- const char
- *option;
+ tempImageBuffer[2];
const void
*inputPixels;
- float
- attenuate,
- *randomNumberBufferPtr;
+ Image
+ *filteredImage;
+
+ int
+ k,
+ matte;
MagickBooleanType
outputReady;
MagickSizeType
length;
- Image
- *filteredImage;
-
- int
- i;
-
- RandomInfo
- **restrict random_info;
-
size_t
global_work_size[2];
unsigned int
- inputColumns,
- inputRows,
- k,
- numRandomNumberPerBuffer,
- numRandomNumberPerPixel,
- numRowsPerKernelLaunch,
- r;
-
-#if defined(MAGICKCORE_OPENMP_SUPPORT)
- unsigned long
- key;
-#endif
+ imageHeight,
+ imageWidth;
void
*filteredPixels,
filteredImage = NULL;
filteredImage_view = NULL;
filteredPixels = NULL;
- randomNumberBufferPtr = NULL;
context = NULL;
imageBuffer = NULL;
- randomNumberBuffer = NULL;
filteredImageBuffer = NULL;
+ hullPass1 = NULL;
+ hullPass2 = NULL;
queue = NULL;
- addNoiseKernel = NULL;
-
+ tempImageBuffer[0] = tempImageBuffer[1] = NULL;
clEnv = GetDefaultOpenCLEnv();
context = GetOpenCLContext(clEnv);
queue = AcquireOpenCLCommandQueue(clEnv);
goto cleanup;
}
+ mem_flags = CL_MEM_READ_WRITE;
+ length = image->columns * image->rows;
+ for (k = 0; k < 2; k++)
+ {
+ tempImageBuffer[k] = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), 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);
assert(filteredImage != NULL);
goto cleanup;
}
- /* find out how many random numbers needed by pixel */
- numRandomNumberPerPixel = 0;
- {
- unsigned int numRandPerChannel = 0;
- switch (noise_type)
- {
- case UniformNoise:
- case ImpulseNoise:
- case LaplacianNoise:
- case RandomNoise:
- default:
- numRandPerChannel = 1;
- break;
- case GaussianNoise:
- case MultiplicativeGaussianNoise:
- case PoissonNoise:
- numRandPerChannel = 2;
- break;
- };
+ hullPass1 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass1");
+ hullPass2 = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "HullPass2");
- if ((channel & RedChannel) != 0)
- numRandomNumberPerPixel+=numRandPerChannel;
- if ((channel & GreenChannel) != 0)
- numRandomNumberPerPixel+=numRandPerChannel;
- if ((channel & BlueChannel) != 0)
- numRandomNumberPerPixel+=numRandPerChannel;
- if ((channel & OpacityChannel) != 0)
- numRandomNumberPerPixel+=numRandPerChannel;
+ clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)&imageBuffer);
+ clStatus |=clEnv->library->clSetKernelArg(hullPass1,1,sizeof(cl_mem),(void *)(tempImageBuffer+1));
+ imageWidth = image->columns;
+ clStatus |=clEnv->library->clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth);
+ imageHeight = image->rows;
+ clStatus |=clEnv->library->clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight);
+ matte = (image->alpha_trait==BlendPixelTrait)?0:1;
+ clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
}
- numRowsPerKernelLaunch = 512;
- /* create a buffer for random numbers */
- numRandomNumberPerBuffer = (image->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
- randomNumberBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, numRandomNumberPerBuffer*sizeof(float)
- , NULL, &clStatus);
-
-
- /* set up the random number generators */
- attenuate=1.0;
- option=GetImageArtifact(image,"attenuate");
- if (option != (char *) NULL)
- attenuate=StringToDouble(option,(char **) NULL);
- random_info=AcquireRandomInfoThreadSet();
-#if defined(MAGICKCORE_OPENMP_SUPPORT)
- key=GetRandomSecretKey(random_info[0]);
-#endif
+ clStatus = clEnv->library->clSetKernelArg(hullPass2,0,sizeof(cl_mem),(void *)(tempImageBuffer+1));
+ clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)tempImageBuffer);
+ imageWidth = image->columns;
+ clStatus |=clEnv->library->clSetKernelArg(hullPass2,2,sizeof(unsigned int),(void *)&imageWidth);
+ imageHeight = image->rows;
+ clStatus |=clEnv->library->clSetKernelArg(hullPass2,3,sizeof(unsigned int),(void *)&imageHeight);
+ matte = (image->alpha_trait==BlendPixelTrait)?0:1;
+ clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
+ }
- addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
- k = 0;
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&imageBuffer);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
- inputColumns = image->columns;
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
- inputRows = image->rows;
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
- attenuate=1.0f;
- option=GetImageArtifact(image,"attenuate");
- if (option != (char *) NULL)
- attenuate=(float)StringToDouble(option,(char **) NULL);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
+ global_work_size[0] = image->columns;
+ global_work_size[1] = image->rows;
- global_work_size[0] = inputColumns;
- for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
+
+ for (k = 0; k < 4; k++)
{
- /* Generate random numbers in the buffer */
- randomNumberBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, randomNumberBuffer, CL_TRUE, CL_MAP_WRITE, 0
- , numRandomNumberPerBuffer*sizeof(float), 0, NULL, NULL, &clStatus);
+ cl_int2 offset;
+ int polarity;
+
+
+ offset.s[0] = X[k];
+ offset.s[1] = Y[k];
+ polarity = 1;
+ clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
if (clStatus != CL_SUCCESS)
{
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
goto cleanup;
}
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
-#if defined(MAGICKCORE_OPENMP_SUPPORT)
- #pragma omp parallel for schedule(static,4) \
- num_threads((key == ~0UL) == 0 ? 1 : (size_t) GetMagickResourceLimit(ThreadResource))
-#endif
- for (i = 0; i < numRandomNumberPerBuffer; i++)
+
+ if (k == 0)
+ clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
+ offset.s[0] = -X[k];
+ offset.s[1] = -Y[k];
+ polarity = 1;
+ clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
+ if (clStatus != CL_SUCCESS)
{
- const int id = GetOpenMPThreadId();
- randomNumberBufferPtr[i] = (float)GetPseudoRandomValue(random_info[id]);
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
}
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
- clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, randomNumberBuffer, randomNumberBufferPtr, 0, NULL, NULL);
+ offset.s[0] = -X[k];
+ offset.s[1] = -Y[k];
+ polarity = -1;
+ clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
if (clStatus != CL_SUCCESS)
{
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.",".");
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
goto cleanup;
}
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+
+ offset.s[0] = X[k];
+ offset.s[1] = Y[k];
+ polarity = -1;
+ clStatus = clEnv->library->clSetKernelArg(hullPass1,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|= clEnv->library->clSetKernelArg(hullPass1,5,sizeof(int),(void *)&polarity);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,4,sizeof(cl_int2),(void *)&offset);
+ clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
+
+ if (k == 3)
+ clStatus |=clEnv->library->clSetKernelArg(hullPass2,1,sizeof(cl_mem),(void *)&filteredImageBuffer);
- /* set the row offset */
- clEnv->library->clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
- global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
- clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
+ /* launch the kernel */
+ clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL);
+ if (clStatus != CL_SUCCESS)
+ {
+ (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+ goto cleanup;
+ }
}
if (ALIGNED(filteredPixels,CLPixelPacket))
if (filteredImage_view != NULL)
filteredImage_view=DestroyCacheView(filteredImage_view);
- if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
- if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
- if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
- if (randomNumberBuffer!=NULL) clEnv->library->clReleaseMemObject(randomNumberBuffer);
- if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
- if (outputReady == MagickFalse && filteredImage != NULL)
+ if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
+ if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
+ for (k = 0; k < 2; k++)
+ {
+ if (tempImageBuffer[k]!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer[k]);
+ }
+ if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
+ if (hullPass1!=NULL) RelinquishOpenCLKernel(clEnv, hullPass1);
+ if (hullPass2!=NULL) RelinquishOpenCLKernel(clEnv, hullPass2);
+ if (outputReady == MagickFalse && filteredImage != NULL)
filteredImage=DestroyImage(filteredImage);
+ return(filteredImage);
+}
+
+MagickExport Image *AccelerateDespeckleImage(const Image* image,
+ ExceptionInfo* exception)
+{
+ Image
+ *filteredImage;
+
+ assert(image != NULL);
+ assert(exception != (ExceptionInfo *) NULL);
+
+ if ((checkOpenCLEnvironment(exception) == MagickFalse) ||
+ (checkAccelerateCondition(image, AllChannels) == MagickFalse))
+ return NULL;
+ filteredImage=ComputeDespeckleImage(image,exception);
return(filteredImage);
}
-static Image *ComputeAddNoiseImageOptRandomNum(const Image*image,
+static Image *ComputeAddNoiseImage(const Image *image,
const ChannelType channel,const NoiseType noise_type,
- ExceptionInfo *exception)
+ ExceptionInfo *exception)
{
CacheView
*filteredImage_view,
context;
cl_int
+ inputPixelCount,
+ pixelsPerWorkitem,
clStatus;
- cl_kernel
- addNoiseKernel,
- randomNumberGeneratorKernel;
+ cl_uint
+ seed0,
+ seed1;
- cl_mem
- filteredImageBuffer,
- imageBuffer,
- randomNumberBuffer,
- randomNumberSeedsBuffer;
+ cl_kernel
+ addNoiseKernel;
cl_mem_flags
mem_flags;
+ cl_mem
+ filteredImageBuffer,
+ imageBuffer;
+
const char
*option;
*inputPixels;
float
- attenuate,
- fNormalize;
-
- Image
- *filteredImage;
-
- int
- i;
+ attenuate;
MagickBooleanType
outputReady;
MagickSizeType
length;
+ Image
+ *filteredImage;
+
+ RandomInfo
+ **restrict random_info;
+
size_t
- global_work_size[2],
- random_work_size;
+ global_work_size[1],
+ local_work_size[1];
unsigned int
- initRandom,
- inputColumns,
- inputRows,
k,
- numRandomNumberGenerators,
- numRandomNumberPerBuffer,
- numRandomNumberPerPixel,
- numRowsPerKernelLaunch,
- r;
+ numRandomNumberPerPixel;
+
+#if defined(MAGICKCORE_OPENMP_SUPPORT)
+ unsigned long
+ key;
+#endif
void
*filteredPixels,
filteredPixels = NULL;
context = NULL;
imageBuffer = NULL;
- randomNumberBuffer = NULL;
filteredImageBuffer = NULL;
- randomNumberSeedsBuffer = NULL;
queue = NULL;
addNoiseKernel = NULL;
- randomNumberGeneratorKernel = NULL;
clEnv = GetDefaultOpenCLEnv();
context = GetOpenCLContext(clEnv);
numRandomNumberPerPixel+=numRandPerChannel;
}
- numRowsPerKernelLaunch = 512;
+ /* set up the random number generators */
+ attenuate=1.0;
+ option=GetImageArtifact(image,"attenuate");
+ if (option != (char *) NULL)
+ attenuate=StringToDouble(option,(char **) NULL);
+ random_info=AcquireRandomInfoThreadSet();
+#if defined(MAGICKCORE_OPENMP_SUPPORT)
+ key=GetRandomSecretKey(random_info[0]);
+#endif
- /* create a buffer for random numbers */
- numRandomNumberPerBuffer = (image->columns*numRowsPerKernelLaunch)*numRandomNumberPerPixel;
- randomNumberBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, numRandomNumberPerBuffer*sizeof(float)
- , NULL, &clStatus);
+ addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"GenerateNoiseImage");
{
- /* setup the random number generators */
- unsigned long* seeds;
- numRandomNumberGenerators = 512;
- randomNumberSeedsBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR|CL_MEM_READ_WRITE
- , numRandomNumberGenerators * 4 * sizeof(unsigned long), NULL, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
- goto cleanup;
- }
- seeds = (unsigned long*) clEnv->library->clEnqueueMapBuffer(queue, randomNumberSeedsBuffer, CL_TRUE, CL_MAP_WRITE, 0
- , numRandomNumberGenerators*4*sizeof(unsigned long), 0, NULL, NULL, &clStatus);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
- goto cleanup;
- }
-
- for (i = 0; i < numRandomNumberGenerators; i++) {
- RandomInfo* randomInfo = AcquireRandomInfo();
- const unsigned long* s = GetRandomInfoSeed(randomInfo);
-
- if (i == 0)
- fNormalize = GetRandomInfoNormalize(randomInfo);
-
- seeds[i*4] = s[0];
- randomInfo = DestroyRandomInfo(randomInfo);
- }
-
- clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, randomNumberSeedsBuffer, seeds, 0, NULL, NULL);
- if (clStatus != CL_SUCCESS)
- {
- (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.",".");
- goto cleanup;
- }
-
- randomNumberGeneratorKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE
- ,"randomNumberGeneratorKernel");
-
- k = 0;
- clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberSeedsBuffer);
- clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(float),(void *)&fNormalize);
- clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
- initRandom = 1;
- clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&initRandom);
- clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerBuffer);
+ cl_uint computeUnitCount;
+ cl_uint workItemCount;
+ clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &computeUnitCount, NULL);
+ workItemCount = computeUnitCount * 2 * 256; // 256 work items per group, 2 groups per CU
+ inputPixelCount = image->columns * image->rows;
+ pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount;
+ pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4;
- random_work_size = numRandomNumberGenerators;
+ local_work_size[0] = 256;
+ global_work_size[0] = workItemCount;
+ }
+ {
+ RandomInfo* randomInfo = AcquireRandomInfo();
+ const unsigned long* s = GetRandomInfoSeed(randomInfo);
+ seed0 = s[0];
+ GetPseudoRandomValue(randomInfo);
+ seed1 = s[0];
+ randomInfo = DestroyRandomInfo(randomInfo);
}
- addNoiseKernel = AcquireOpenCLKernel(clEnv,MAGICK_OPENCL_ACCELERATE,"AddNoiseImage");
k = 0;
clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&imageBuffer);
clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
- inputColumns = image->columns;
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputColumns);
- inputRows = image->rows;
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&inputRows);
+ clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&inputPixelCount);
+ clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&pixelsPerWorkitem);
clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&channel);
clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
attenuate=1.0f;
if (option != (char *) NULL)
attenuate=(float)StringToDouble(option,(char **) NULL);
clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(float),(void *)&attenuate);
- clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&randomNumberBuffer);
+ clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed0);
+ clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed1);
clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(unsigned int),(void *)&numRandomNumberPerPixel);
- global_work_size[0] = inputColumns;
- for (r = 0; r < inputRows; r+=numRowsPerKernelLaunch)
- {
- size_t generator_local_size = 64;
- /* Generate random numbers in the buffer */
- clEnv->library->clEnqueueNDRangeKernel(queue,randomNumberGeneratorKernel,1,NULL
- ,&random_work_size,&generator_local_size,0,NULL,NULL);
- if (initRandom != 0)
- {
- /* make sure we only do init once */
- initRandom = 0;
- clEnv->library->clSetKernelArg(randomNumberGeneratorKernel,3,sizeof(unsigned int),(void *)&initRandom);
- }
-
- /* set the row offset */
- clEnv->library->clSetKernelArg(addNoiseKernel,k,sizeof(unsigned int),(void *)&r);
- global_work_size[1] = MAGICK_MIN(numRowsPerKernelLaunch, inputRows - r);
- clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,2,NULL,global_work_size,NULL,0,NULL,NULL);
- }
+ clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,NULL,0,NULL,NULL);
if (ALIGNED(filteredPixels,CLPixelPacket))
{
if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue);
if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel);
- if (randomNumberGeneratorKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomNumberGeneratorKernel);
if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer);
- if (randomNumberBuffer!=NULL) clEnv->library->clReleaseMemObject(randomNumberBuffer);
if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer);
- if (randomNumberSeedsBuffer!=NULL) clEnv->library->clReleaseMemObject(randomNumberSeedsBuffer);
if (outputReady == MagickFalse && filteredImage != NULL)
filteredImage=DestroyImage(filteredImage);
return(filteredImage);
}
+
MagickExport Image *AccelerateAddNoiseImage(const Image *image,
const ChannelType channel,const NoiseType noise_type,
ExceptionInfo *exception)
(checkAccelerateCondition(image, channel) == MagickFalse))
return NULL;
-DisableMSCWarning(4127)
- if (sizeof(unsigned long) == 4)
-RestoreMSCWarning
- filteredImage = ComputeAddNoiseImageOptRandomNum(image,channel,noise_type,exception);
- else
- filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
+ filteredImage = ComputeAddNoiseImage(image,channel,noise_type,exception);
return(filteredImage);
}