From: dirk Date: Sat, 2 Apr 2016 20:28:46 +0000 (+0200) Subject: AccelerateUnsharpMaskImage now supports R/RA/RGB images. X-Git-Tag: 7.0.1-0~76 X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=9150168b2d6173f7d89da6c5cc3efc401f8f26a3;p=imagemagick AccelerateUnsharpMaskImage now supports R/RA/RGB images. --- diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index 5b8faa796..d2a6b1e49 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -358,10 +358,10 @@ OPENCL_ENDIF() STRINGIFY( - inline __global CLQuantum *getPixel(__global CLQuantum *image, const unsigned int number_channels, + inline unsigned int getPixelIndex(const unsigned int number_channels, const unsigned int columns, const unsigned int x, const unsigned int y) { - return image + (x * number_channels) + (y * columns * number_channels); + return (x * number_channels) + (y * columns * number_channels); } inline float getPixelRed(const __global CLQuantum *p) { return (float)*p; } @@ -414,10 +414,10 @@ OPENCL_ENDIF() *alpha=getPixelAlpha(p); } - inline float4 ReadFloat4(__global CLQuantum *image, const unsigned int number_channels, + inline float4 ReadFloat4(const __global CLQuantum *image, const unsigned int number_channels, const unsigned int columns, const unsigned int x, const unsigned int y, const ChannelType channel) { - const __global CLQuantum *p = getPixel(image, number_channels, columns, x, y); + const __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y); float red = 0.0f; float green = 0.0f; @@ -452,7 +452,7 @@ OPENCL_ENDIF() const unsigned int columns, const unsigned int x, const unsigned int y, const ChannelType channel, float4 pixel) { - __global CLQuantum *p = getPixel(image, number_channels, columns, x, y); + __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y); WriteChannels(p, number_channels, channel, pixel.x, pixel.y, pixel.z, pixel.w); } @@ -866,7 +866,7 @@ OPENCL_ENDIF() /* Reduce image noise and reduce detail levels by row */ - __kernel void BlurRow(__global CLQuantum *image, + __kernel void BlurRow(const __global CLQuantum *image, const unsigned int number_channels,const ChannelType channel, __constant float *filter,const unsigned int width, const unsigned int imageColumns,const unsigned int imageRows, @@ -1900,15 +1900,14 @@ OPENCL_ENDIF() number_parameters : numbers of parameters parameters : the parameter */ - __kernel void ComputeFunction(__global CLQuantum *image, - const unsigned int number_channels,const ChannelType channel, - const MagickFunction function,const unsigned int number_parameters, + __kernel void ComputeFunction(__global CLQuantum *image,const unsigned int number_channels, + const ChannelType channel,const MagickFunction function,const unsigned int number_parameters, __constant float *parameters) { const unsigned int x = get_global_id(0); const unsigned int y = get_global_id(1); const unsigned int columns = get_global_size(0); - __global CLQuantum *p = getPixel(image, number_channels, columns, x, y); + __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y); float red; float green; @@ -1956,7 +1955,7 @@ OPENCL_ENDIF() const unsigned int x = get_global_id(0); const unsigned int y = get_global_id(1); const unsigned int columns = get_global_size(0); - __global CLQuantum *p = getPixel(image, number_channels, columns, x, y); + __global CLQuantum *p = image + getPixelIndex(number_channels, columns, x, y); float blue, @@ -3101,12 +3100,13 @@ STRINGIFY( */ STRINGIFY( - __kernel void UnsharpMaskBlurColumn(const __global CLPixelType* inputImage, - const __global float4 *blurRowData, __global CLPixelType *filtered_im, - const unsigned int imageColumns, const unsigned int imageRows, - __local float4* cachedData, __local float* cachedFilter, - const ChannelType channel, const __global float *filter, const unsigned int width, - const float gain, const float threshold) + __kernel void UnsharpMaskBlurColumn(const __global CLQuantum* image, + const __global float4 *blurRowData,const unsigned int number_channels, + const ChannelType channel,const unsigned int columns, + const unsigned int rows,__local float4* cachedData, + __local float* cachedFilter,const __global float *filter, + const unsigned int width,const float gain, const float threshold, + __global CLQuantum *filteredImage) { const unsigned int radius = (width-1)/2; @@ -3115,17 +3115,17 @@ STRINGIFY( const int groupStartY = get_group_id(1)*get_local_size(1) - radius; const int groupStopY = (get_group_id(1)+1)*get_local_size(1) + radius; - if (groupStartY >= 0 - && groupStopY < imageRows) { - event_t e = async_work_group_strided_copy(cachedData - ,blurRowData+groupStartY*imageColumns+groupX - ,groupStopY-groupStartY,imageColumns,0); + if ((groupStartY >= 0) && (groupStopY < rows)) + { + event_t e = async_work_group_strided_copy(cachedData, + blurRowData+groupStartY*columns+groupX,groupStopY-groupStartY,columns,0); wait_group_events(1,&e); } - else { - for (int i = get_local_id(1); i < (groupStopY - groupStartY); i+=get_local_size(1)) { - cachedData[i] = blurRowData[ClampToCanvas(groupStartY+i,imageRows)*imageColumns+ groupX]; - } + else + { + for (int i = get_local_id(1); i < (groupStopY - groupStartY); i+=get_local_size(1)) + cachedData[i] = blurRowData[ClampToCanvas(groupStartY+i,rows)*columns + groupX]; + barrier(CLK_LOCAL_MEM_FENCE); } // cache the filter as well @@ -3133,36 +3133,24 @@ STRINGIFY( wait_group_events(1,&e); // only do the work if this is not a patched item - //const int cy = get_group_id(1)*get_local_size(1)+get_local_id(1); const int cy = get_global_id(1); - if (cy < imageRows) { + if (cy < rows) + { float4 blurredPixel = (float4) 0.0f; int i = 0; - \n #ifndef UFACTOR \n - \n #define UFACTOR 8 \n - \n #endif \n - - for ( ; i+UFACTOR < width; ) - { - \n #pragma unroll UFACTOR \n - for (int j=0; j < UFACTOR; j++, i++) - { - blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)]; - } - } - - for ( ; i < width; i++) + for ( ; i+7 < width; ) { - blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)]; + for (int j=0; j < 8; j++, i++) + blurredPixel+=cachedFilter[i+j]*cachedData[i+j+get_local_id(1)]; } - blurredPixel = floor((float4)(ClampToQuantum(blurredPixel.x), ClampToQuantum(blurredPixel.y) - ,ClampToQuantum(blurredPixel.z), ClampToQuantum(blurredPixel.w))); + for ( ; i < width; i++) + blurredPixel+=cachedFilter[i]*cachedData[i+get_local_id(1)]; - float4 inputImagePixel = convert_float4(inputImage[cy*imageColumns+groupX]); + float4 inputImagePixel = ReadFloat4(image,number_channels,columns,groupX,cy,channel); float4 outputPixel = inputImagePixel - blurredPixel; float quantumThreshold = QuantumRange*threshold; @@ -3171,15 +3159,13 @@ STRINGIFY( outputPixel = select(inputImagePixel + outputPixel * gain, inputImagePixel, mask); //write back - filtered_im[cy*imageColumns+groupX] = (CLPixelType) (ClampToQuantum(outputPixel.x), ClampToQuantum(outputPixel.y) - ,ClampToQuantum(outputPixel.z), ClampToQuantum(outputPixel.w)); - + WriteFloat4(filteredImage,number_channels,columns,groupX,cy,channel,outputPixel); } } ) STRINGIFY( - __kernel void UnsharpMask(__global CLQuantum *image,const unsigned int number_channels, + __kernel void UnsharpMask(const __global CLQuantum *image,const unsigned int number_channels, const ChannelType channel,__constant float *filter,const unsigned int width, const unsigned int columns,const unsigned int rows,__local float4 *pixels, const float gain,const float threshold, const unsigned int justBlur, diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 5e813adc3..889381692 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -5850,9 +5850,6 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, *filteredImage_view, *image_view; - char - geometry[MagickPathExtent]; - cl_command_queue queue; @@ -5875,16 +5872,15 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, imageKernelBuffer, tempImageBuffer; - cl_mem_flags - mem_flags; - - const void - *inputPixels; + cl_uint + imageColumns, + imageRows, + kernelWidth, + number_channels; float fGain, - fThreshold, - *kernelBufferPtr; + fThreshold; Image *filteredImage; @@ -5892,9 +5888,6 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, int chunkSize; - KernelInfo - *kernel; - MagickBooleanType outputReady; @@ -5905,22 +5898,18 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, length; void - *filteredPixels, - *hostPtr; + *filteredPixels; unsigned int - i, - imageColumns, - imageRows, - kernelWidth; + i; clEnv = NULL; filteredImage = NULL; filteredImage_view = NULL; - kernel = NULL; context = NULL; imageBuffer = NULL; filteredImageBuffer = NULL; + filteredPixels = NULL; tempImageBuffer = NULL; imageKernelBuffer = NULL; blurRowKernel = NULL; @@ -5932,116 +5921,33 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, context = GetOpenCLContext(clEnv); queue = AcquireOpenCLCommandQueue(clEnv); - /* Create and initialize OpenCL buffers. */ - { - image_view=AcquireVirtualCacheView(image,exception); - inputPixels=GetCacheViewVirtualPixels(image_view,0,0,image->columns,image->rows,exception); - if (inputPixels == (const void *) NULL) - { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); - goto cleanup; - } - - /* If the host pointer is aligned to the size of CLPixelPacket, - then use the host buffer directly from the GPU; otherwise, - create a buffer on the GPU and copy the data over */ - if (ALIGNED(inputPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; - } - else - { - mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; - } - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - } + image_view = AcquireVirtualCacheView(image, exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + if (imageBuffer == (cl_mem) NULL) + goto cleanup; - /* create output */ + filteredImage=CloneImage(image,0,0,MagickTrue,exception); + if (filteredImage == (Image *) NULL) + goto cleanup; + if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) { - filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); - assert(filteredImage != NULL); - if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); - goto cleanup; - } - filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); - filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); - if (filteredPixels == (void *) NULL) - { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); - goto cleanup; - } - - if (ALIGNED(filteredPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; - hostPtr = filteredPixels; - } - else - { - mem_flags = CL_MEM_WRITE_ONLY; - hostPtr = NULL; - } - - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } + (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + goto cleanup; } + filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); + filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, + context,filteredPixels,exception); + if (filteredImageBuffer == (cl_mem) NULL) + goto cleanup; - /* create the blur kernel */ - { - (void) FormatLocaleString(geometry,MagickPathExtent,"blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma); - kernel=AcquireKernelInfo(geometry,exception); - if (kernel == (KernelInfo *) NULL) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.","."); - goto cleanup; - } - - imageKernelBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - - - kernelBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.","."); - goto cleanup; - } - for (i = 0; i < kernel->width; i++) - { - kernelBufferPtr[i] = (float) kernel->values[i]; - } - clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", "."); - goto cleanup; - } - } + imageKernelBuffer=createKernelInfo(clEnv,context,queue,radius,sigma, + &kernelWidth,exception); { /* create temp buffer */ { length = image->columns * image->rows; - tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus); + tempImageBuffer = clEnv->library->clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus); if (clStatus != CL_SUCCESS) { (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); @@ -6066,24 +5972,24 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, }; } + number_channels = (cl_uint) image->number_channels; + imageColumns = (cl_uint) image->columns; + imageRows = (cl_uint) image->rows; + { chunkSize = 256; - imageColumns = (unsigned int) image->columns; - imageRows = (unsigned int) image->rows; - - kernelWidth = (unsigned int) kernel->width; - /* set the kernel arguments */ i = 0; clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); - clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); + clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels); clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask); clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); - clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&kernelWidth); - clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageColumns); - clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&imageRows); - clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *) NULL); + clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth); + clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns); + clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows); + clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL); + clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); if (clStatus != CL_SUCCESS) { (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); @@ -6101,7 +6007,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, wsize[0] = chunkSize; wsize[1] = 1; - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event); + clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); @@ -6115,25 +6021,23 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, { chunkSize = 256; - imageColumns = (unsigned int) image->columns; - imageRows = (unsigned int) image->rows; - kernelWidth = (unsigned int) kernel->width; fGain = (float) gain; fThreshold = (float) threshold; i = 0; clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageColumns); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&imageRows); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, (chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++, kernelWidth*sizeof(float),NULL); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),&number_channels); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,(chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,kernelWidth*sizeof(float),NULL); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer); - clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(unsigned int),(void *)&kernelWidth); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain); clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold); + clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); if (clStatus != CL_SUCCESS) { @@ -6152,7 +6056,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, wsize[0] = 1; wsize[1] = chunkSize; - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event); + clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event); if (clStatus != CL_SUCCESS) { (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); @@ -6166,19 +6070,9 @@ static Image *ComputeUnsharpMaskImage(const Image *image,const double radius, } /* get result */ - if (ALIGNED(filteredPixels,CLPixelPacket)) - { - length = image->columns * image->rows; - clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); - } - else - { - length = image->columns * image->rows; - clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); - } - if (clStatus != CL_SUCCESS) + if (copyWriteBuffer(image,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } @@ -6191,7 +6085,6 @@ cleanup: if (filteredImage_view != NULL) filteredImage_view=DestroyCacheView(filteredImage_view); - if (kernel != NULL) kernel=DestroyKernelInfo(kernel); if (imageBuffer!=NULL) clEnv->library->clReleaseMemObject(imageBuffer); if (filteredImageBuffer!=NULL) clEnv->library->clReleaseMemObject(filteredImageBuffer); if (tempImageBuffer!=NULL) clEnv->library->clReleaseMemObject(tempImageBuffer); @@ -6396,7 +6289,7 @@ MagickExport Image *AccelerateUnsharpMaskImage(const Image *image, assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || + if ((checkAccelerateCondition(image) == MagickFalse) || (checkOpenCLEnvironment(exception) == MagickFalse)) return NULL;