From a22457d1cb46b3681b0758006816ffc2a7fcf1f7 Mon Sep 17 00:00:00 2001 From: cristy Date: Sat, 7 Dec 2013 14:03:06 +0000 Subject: [PATCH] --- MagickCore/accelerate-private.h | 100 +++++++- MagickCore/accelerate.c | 438 ++++++++++++++++---------------- MagickCore/opencl.c | 99 +++++++- 3 files changed, 418 insertions(+), 219 deletions(-) diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index 739dd3ad5..fda847282 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -35,11 +35,20 @@ extern "C" { typedef struct _FloatPixelPacket { +#ifdef MAGICK_PIXEL_RGBA MagickRealType red, green, blue, - alpha; + opacity; +#endif +#ifdef MAGICK_PIXEL_BGRA + MagickRealType + blue, + green, + red, + opacity; +#endif } FloatPixelPacket; const char* accelerateKernels = @@ -184,7 +193,7 @@ const char* accelerateKernels = STRINGIFY( __kernel - void Convolve(const __global CLPixelType *input, __global CLPixelType *output, + void ConvolveOptimized(const __global CLPixelType *input, __global CLPixelType *output, const unsigned int imageWidth, const unsigned int imageHeight, __constant float *filter, const unsigned int filterWidth, const unsigned int filterHeight, const uint matte, const ChannelType channel, __local CLPixelType *pixelLocalCache, __local float* filterCache) { @@ -305,6 +314,93 @@ const char* accelerateKernels = } ) + STRINGIFY( + __kernel + void Convolve(const __global CLPixelType *input, __global CLPixelType *output, + __constant float *filter, const unsigned int filterWidth, const unsigned int filterHeight, + const uint matte, const ChannelType channel) { + + int2 imageIndex; + imageIndex.x = get_global_id(0); + imageIndex.y = get_global_id(1); + + unsigned int imageWidth = get_global_size(0); + unsigned int imageHeight = get_global_size(1); + + if (imageIndex.x >= imageWidth + || imageIndex.y >= imageHeight) + return; + + int2 midFilterDimen; + midFilterDimen.x = (filterWidth-1)/2; + midFilterDimen.y = (filterHeight-1)/2; + + int filterIndex = 0; + float4 sum = (float4)0.0f; + float gamma = 0.0f; + if (((channel & OpacityChannel) == 0) || (matte == 0)) { + for (int j = 0; j < filterHeight; j++) { + int2 inputPixelIndex; + inputPixelIndex.y = imageIndex.y - midFilterDimen.y + j; + inputPixelIndex.y = ClampToCanvas(inputPixelIndex.y, imageHeight); + for (int i = 0; i < filterWidth; i++) { + inputPixelIndex.x = imageIndex.x - midFilterDimen.x + i; + inputPixelIndex.x = ClampToCanvas(inputPixelIndex.x, imageWidth); + + CLPixelType p = input[inputPixelIndex.y * imageWidth + inputPixelIndex.x]; + float f = filter[filterIndex]; + + sum.x += f * p.x; + sum.y += f * p.y; + sum.z += f * p.z; + sum.w += f * p.w; + + gamma += f; + + filterIndex++; + } + } + } + else { + + for (int j = 0; j < filterHeight; j++) { + int2 inputPixelIndex; + inputPixelIndex.y = imageIndex.y - midFilterDimen.y + j; + inputPixelIndex.y = ClampToCanvas(inputPixelIndex.y, imageHeight); + for (int i = 0; i < filterWidth; i++) { + inputPixelIndex.x = imageIndex.x - midFilterDimen.x + i; + inputPixelIndex.x = ClampToCanvas(inputPixelIndex.x, imageWidth); + + CLPixelType p = input[inputPixelIndex.y * imageWidth + inputPixelIndex.x]; + float alpha = QuantumScale*(QuantumRange-p.w); + float f = filter[filterIndex]; + float g = alpha * f; + + sum.x += g*p.x; + sum.y += g*p.y; + sum.z += g*p.z; + sum.w += f*p.w; + + gamma += g; + + + filterIndex++; + } + } + gamma = PerceptibleReciprocal(gamma); + sum.xyz = gamma*sum.xyz; + } + + CLPixelType outputPixel; + outputPixel.x = ClampToQuantum(sum.x); + outputPixel.y = ClampToQuantum(sum.y); + outputPixel.z = ClampToQuantum(sum.z); + outputPixel.w = ((channel & OpacityChannel)!=0)?ClampToQuantum(sum.w):input[imageIndex.y * imageWidth + imageIndex.x].w; + + output[imageIndex.y * imageWidth + imageIndex.x] = outputPixel; + } + ) + STRINGIFY( typedef enum { diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 5f022d4ae..4ce01e553 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -95,7 +95,7 @@ static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception) MagickCLEnv clEnv; clEnv = GetDefaultOpenCLEnv(); - + GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED , sizeof(MagickBooleanType), &flag, exception); if (flag == MagickTrue) @@ -107,6 +107,11 @@ static MagickBooleanType checkOpenCLEnvironment(ExceptionInfo* exception) { if(InitOpenCLEnv(clEnv, exception) == MagickFalse) return MagickFalse; + + GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED + , sizeof(MagickBooleanType), &flag, exception); + if (flag == MagickTrue) + return MagickFalse; } return MagickTrue; @@ -185,7 +190,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -207,7 +212,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -215,13 +220,13 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -240,7 +245,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -248,7 +253,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch convolutionKernel = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernelSize * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -258,7 +263,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch , 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } for (i = 0; i < kernelSize; i++) @@ -268,7 +273,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch clStatus = clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -295,10 +300,10 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch if (localMemoryRequirement <= deviceLocalMemorySize) { /* get the OpenCL kernel */ - clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve"); + clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "ConvolveOptimized"); if (clkernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -322,7 +327,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch clStatus|=clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -334,7 +339,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch clStatus = clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } } @@ -344,7 +349,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Convolve"); if (clkernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -362,7 +367,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch clStatus|=clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&channel); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -373,7 +378,7 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch clStatus = clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } } @@ -391,15 +396,15 @@ static Image* ComputeConvolveImage(const Image* inputImage, const ChannelType ch } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } /* everything is fine! :) */ outputReady = MagickTrue; - cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); if (inputImageBuffer != NULL) clReleaseMemObject(inputImageBuffer); @@ -478,7 +483,6 @@ MagickExport Image* AccelerateConvolveImageChannel(const Image *image, const Cha return NULL; filteredImage = ComputeConvolveImage(image, channel, kernel, exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return filteredImage; } @@ -518,7 +522,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch pixels = GetPixelCachePixels(image, &length, exception); if (pixels == (void *) NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), CacheWarning, + (void) OpenCLThrowMagickException(exception, GetMagickModule(), CacheWarning, "GetPixelCachePixels failed.", "'%s'", image->filename); goto cleanup; @@ -538,14 +542,14 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch imageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)pixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } parametersBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -555,7 +559,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch , 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } for (i = 0; i < number_parameters; i++) @@ -565,7 +569,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch clStatus = clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -573,7 +577,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch clkernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "FunctionImage"); if (clkernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -586,7 +590,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch clStatus|=clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)¶metersBuffer); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -596,7 +600,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch clStatus = clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -614,12 +618,13 @@ static MagickBooleanType ComputeFunctionImage(Image *image, const ChannelType ch } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } status = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); if (clkernel != NULL) RelinquishOpenCLKernel(clEnv, clkernel); if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); @@ -649,7 +654,6 @@ MagickExport MagickBooleanType if (status == MagickTrue) { status = ComputeFunctionImage(image, channel, function, number_parameters, parameters, exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); } } return status; @@ -729,7 +733,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } /* If the host pointer is aligned to the size of CLPixelPacket, @@ -748,7 +752,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -759,13 +763,13 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -784,7 +788,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -795,20 +799,20 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe kernel=AcquireKernelInfo(geometry); if (kernel == (KernelInfo *) NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.","."); goto cleanup; } imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } @@ -820,7 +824,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -833,7 +837,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -843,14 +847,14 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow"); if (blurRowKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumn"); if (blurColumnKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -876,7 +880,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -894,7 +898,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -922,7 +926,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernel->width),(void *)NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -940,7 +944,7 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe clStatus = clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -962,13 +966,15 @@ static Image* ComputeBlurImage(const Image* inputImage, const ChannelType channe } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer); if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer); @@ -1038,7 +1044,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } /* If the host pointer is aligned to the size of CLPixelPacket, @@ -1057,7 +1063,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -1068,13 +1074,13 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -1093,7 +1099,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -1104,20 +1110,20 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType kernel=AcquireKernelInfo(geometry); if (kernel == (KernelInfo *) NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "MemoryAllocationFailed.","."); goto cleanup; } imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, kernel->width * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } @@ -1129,7 +1135,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -1144,7 +1150,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -1154,14 +1160,14 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection"); if (blurRowKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; blurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurColumnSection"); if (blurColumnKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -1197,7 +1203,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -1215,7 +1221,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -1251,7 +1257,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType clStatus|=clSetKernelArg(blurColumnKernel,i++,sizeof(unsigned int),(void *)&sec); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -1269,7 +1275,7 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType clStatus = clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -1292,13 +1298,15 @@ static Image* ComputeBlurImageSection(const Image* inputImage, const ChannelType } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer); if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer); @@ -1378,7 +1386,6 @@ Image* AccelerateBlurImage(const Image *image, const ChannelType channel, const else filteredImage = ComputeBlurImage(image, channel, radius, sigma, exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return filteredImage; } @@ -1436,7 +1443,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -1456,7 +1463,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -1465,13 +1472,13 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -1490,7 +1497,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -1503,13 +1510,13 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType sinThetaBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } cosThetaBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -1518,14 +1525,14 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType sinThetaPtr = (float*) clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); goto cleanup; } cosThetaPtr = (float*) clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.","."); goto cleanup; } @@ -1541,7 +1548,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType clStatus |= clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } @@ -1549,7 +1556,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType radialBlurKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "RadialBlur"); if (radialBlurKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -1577,7 +1584,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType clStatus|=clSetKernelArg(radialBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -1588,7 +1595,7 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType clStatus = clEnqueueNDRangeKernel(queue, radialBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -1605,12 +1612,14 @@ static Image* ComputeRadialBlurImage(const Image *inputImage, const ChannelType } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer); if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (sinThetaBuffer!=NULL) clReleaseMemObject(sinThetaBuffer); @@ -1681,7 +1690,6 @@ Image* AccelerateRadialBlurImage(const Image *image, const ChannelType channel, return NULL; filteredImage = ComputeRadialBlurImage(image, channel, angle, exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return filteredImage; } @@ -1732,7 +1740,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -1752,7 +1760,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -1763,13 +1771,13 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -1789,7 +1797,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -1800,14 +1808,14 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType kernel=AcquireKernelInfo(geometry); if (kernel == (KernelInfo *) NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.","."); goto cleanup; } imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -1815,7 +1823,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } for (i = 0; i < kernel->width; i++) @@ -1825,7 +1833,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -1837,7 +1845,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -1847,14 +1855,14 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRow"); if (blurRowKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumn"); if (unsharpMaskBlurColumnKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -1879,7 +1887,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(CLPixelPacket)*(chunkSize+kernel->width),(void *)NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -1897,7 +1905,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -1928,7 +1936,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -1946,7 +1954,7 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType clStatus = clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -1967,13 +1975,15 @@ static Image* ComputeUnsharpMaskImage(const Image *inputImage, const ChannelType } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (kernel != NULL) kernel=DestroyKernelInfo(kernel); if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer); @@ -2039,7 +2049,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -2059,7 +2069,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -2070,13 +2080,13 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -2096,7 +2106,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -2107,14 +2117,14 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan kernel=AcquireKernelInfo(geometry); if (kernel == (KernelInfo *) NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireKernelInfo failed.","."); goto cleanup; } imageKernelBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY, kernel->width * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -2122,7 +2132,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan kernelBufferPtr = (float*)clEnqueueMapBuffer(queue, imageKernelBuffer, CL_TRUE, CL_MAP_WRITE, 0, kernel->width * sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } for (i = 0; i < kernel->width; i++) @@ -2132,7 +2142,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan clStatus = clEnqueueUnmapMemObject(queue, imageKernelBuffer, kernelBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -2147,7 +2157,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length * 4 * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -2157,14 +2167,14 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan blurRowKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "BlurRowSection"); if (blurRowKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "UnsharpMaskBlurColumnSection"); if (unsharpMaskBlurColumnKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; }; } @@ -2198,7 +2208,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan clStatus|=clSetKernelArg(blurRowKernel,i++,sizeof(unsigned int),(void *)&sec); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -2215,7 +2225,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan clStatus = clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -2227,7 +2237,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan imageColumns = inputImage->columns; if (sec == 0) - imageRows = inputImage->rows / 2 + (kernel->width-1) / 2; + imageRows = inputImage->rows / 2; else imageRows = (inputImage->rows - inputImage->rows / 2); @@ -2256,7 +2266,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } } @@ -2274,7 +2284,7 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan clStatus = clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -2295,13 +2305,15 @@ static Image* ComputeUnsharpMaskImageSection(const Image *inputImage, const Chan } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (kernel != NULL) kernel=DestroyKernelInfo(kernel); if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer); @@ -2391,7 +2403,6 @@ Image* AccelerateUnsharpMaskImage(const Image *image, const ChannelType channel, filteredImage = ComputeUnsharpMaskImageSection(image,channel,radius,sigma,gain,threshold,exception); else filteredImage = ComputeUnsharpMaskImage(image,channel,radius,sigma,gain,threshold,exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return filteredImage; } @@ -2452,9 +2463,7 @@ static MagickBooleanType resizeHorizontalFilter(cl_mem inputImage /* 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)/xFactor+MagickEpsilon)-support+0.5); @@ -2508,7 +2517,7 @@ RestoreMSCWarning } if (horizontalKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -2552,7 +2561,7 @@ RestoreMSCWarning if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -2564,7 +2573,7 @@ RestoreMSCWarning clStatus = clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -2572,6 +2581,8 @@ RestoreMSCWarning cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel); return status; @@ -2634,9 +2645,7 @@ static MagickBooleanType resizeVerticalFilter(cl_mem inputImage /* 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); @@ -2686,7 +2695,7 @@ RestoreMSCWarning if (horizontalKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -2730,7 +2739,7 @@ RestoreMSCWarning if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -2742,7 +2751,7 @@ RestoreMSCWarning clStatus = clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -2750,6 +2759,8 @@ RestoreMSCWarning cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (horizontalKernel != NULL) RelinquishOpenCLKernel(clEnv, horizontalKernel); return status; @@ -2793,7 +2804,7 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (const void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -2813,14 +2824,14 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } cubicCoefficientsBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } queue = AcquireOpenCLCommandQueue(clEnv); @@ -2828,7 +2839,7 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo , 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter); @@ -2839,7 +2850,7 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo clStatus = clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } @@ -2849,13 +2860,13 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -2875,7 +2886,7 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -2888,7 +2899,7 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -2912,7 +2923,7 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo tempImageBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE, length*sizeof(CLPixelPacket), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -2941,12 +2952,14 @@ static Image* ComputeResizeImage(const Image* inputImage, const size_t resizedCo } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (tempImageBuffer!=NULL) clReleaseMemObject(tempImageBuffer); if (filteredImageBuffer!=NULL) clReleaseMemObject(filteredImageBuffer); @@ -3060,7 +3073,6 @@ Image* AccelerateResizeImage(const Image* image, const size_t resizedColumns, co return NULL; filteredImage = ComputeResizeImage(image,resizedColumns,resizedRows,resizeFilter,exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return filteredImage; } @@ -3092,7 +3104,7 @@ static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBoo inputPixels = GetPixelCachePixels(inputImage, &length, exception); if (inputPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -3112,14 +3124,14 @@ static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBoo inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } filterKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Contrast"); if (filterKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3130,7 +3142,7 @@ static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBoo clStatus|=clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -3141,7 +3153,7 @@ static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBoo clStatus = clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -3158,12 +3170,13 @@ static MagickBooleanType ComputeContrastImage(Image *inputImage, const MagickBoo } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); if (filterKernel!=NULL) RelinquishOpenCLKernel(clEnv, filterKernel); @@ -3216,7 +3229,6 @@ MagickBooleanType AccelerateContrastImage(Image* image, const MagickBooleanType return MagickFalse; status = ComputeContrastImage(image,sharpen,exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return status; } @@ -3277,7 +3289,7 @@ MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, inputPixels = GetPixelCachePixels(inputImage, &length, exception); if (inputPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -3298,14 +3310,14 @@ MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } modulateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Modulate"); if (modulateKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3322,7 +3334,7 @@ MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, clStatus|=clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); printf("no kernel\n"); goto cleanup; } @@ -3335,7 +3347,7 @@ MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, clStatus = clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -3353,13 +3365,14 @@ MagickBooleanType ComputeModulateImage(Image* image, double percent_brightness, } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); if (inputPixels) { //ReleasePixelCachePixels(); @@ -3428,7 +3441,6 @@ MagickBooleanType AccelerateModulateImage(Image* image, double percent_brightnes status = ComputeModulateImage(image,percent_brightness, percent_hue, percent_saturation, colorspace, exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return status; } @@ -3517,7 +3529,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha if (inputPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } /* If the host pointer is aligned to the size of CLPixelPacket, @@ -3536,7 +3548,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -3558,7 +3570,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha histogramBuffer = clCreateBuffer(context, mem_flags, length * sizeof(cl_uint4), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -3580,7 +3592,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha histogramKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Histogram"); if (histogramKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3592,7 +3604,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha clStatus|=clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -3604,7 +3616,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -3622,7 +3634,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } @@ -3632,7 +3644,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha clStatus = clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueUnmapMemObject failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.", "'%s'", "."); goto cleanup; } } @@ -3644,7 +3656,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha printf("histogram %d: red %d\n", i, histogram[i].s[2]); printf("histogram %d: green %d\n", i, histogram[i].s[1]); printf("histogram %d: blue %d\n", i, histogram[i].s[0]); - printf("histogram %d: opacity %d\n", i, histogram[i].s[3]); + printf("histogram %d: alpha %d\n", i, histogram[i].s[3]); } } @@ -3787,7 +3799,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -3807,7 +3819,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha equalizeMapBuffer = clCreateBuffer(context, mem_flags, length * sizeof(PixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -3815,7 +3827,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha equalizeKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Equalize"); if (equalizeKernel == NULL) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "AcquireOpenCLKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); goto cleanup; } @@ -3828,7 +3840,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha clStatus|=clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -3840,7 +3852,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } clFlush(queue); @@ -3858,7 +3870,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } @@ -3867,6 +3879,7 @@ MagickExport MagickBooleanType ComputeEqualizeImage(Image *inputImage, const Cha equalize_map=(PixelPacket *) RelinquishMagickMemory(equalize_map); cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); if (inputPixels) { /*ReleasePixelCachePixels();*/ @@ -3937,7 +3950,6 @@ MagickBooleanType AccelerateEqualizeImage(Image* image, const ChannelType channe return MagickFalse; status = ComputeEqualizeImage(image,channel,exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return status; } @@ -3982,7 +3994,7 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -3999,7 +4011,7 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -4010,7 +4022,7 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce tempImageBuffer[k] = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } } @@ -4019,13 +4031,13 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -4044,7 +4056,7 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -4057,11 +4069,11 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce clStatus |=clSetKernelArg(hullPass1,2,sizeof(unsigned int),(void *)&imageWidth); imageHeight = inputImage->rows; clStatus |=clSetKernelArg(hullPass1,3,sizeof(unsigned int),(void *)&imageHeight); - matte = (inputImage->alpha_trait == BlendPixelTrait)?1:0; + matte = (inputImage->matte==MagickFalse)?0:1; clStatus |=clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -4075,7 +4087,7 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce clStatus |=clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } @@ -4099,21 +4111,21 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } @@ -4129,21 +4141,21 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } @@ -4156,21 +4168,21 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce clStatus|=clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } @@ -4187,21 +4199,21 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clSetKernelArg failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clSetKernelArg failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } /* launch the kernel */ clStatus = clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clEnqueueNDRangeKernel failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueNDRangeKernel failed.", "'%s'", "."); goto cleanup; } } @@ -4218,13 +4230,15 @@ static Image* ComputeDespeckleImage(const Image* inputImage, ExceptionInfo* exce } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } outputReady = MagickTrue; cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (queue != NULL) RelinquishOpenCLCommandQueue(clEnv, queue); if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); for (k = 0; k < 2; k++) @@ -4293,7 +4307,6 @@ Image* AccelerateDespeckleImage(const Image* image, ExceptionInfo* exception) return NULL; newImage = ComputeDespeckleImage(image,exception); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return newImage; } @@ -4344,7 +4357,7 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -4361,7 +4374,7 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -4370,13 +4383,13 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -4395,7 +4408,7 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -4473,7 +4486,7 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, , numRandomNumberPerBuffer*sizeof(float), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } @@ -4490,7 +4503,7 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, clStatus = clEnqueueUnmapMemObject(queue, randomNumberBuffer, randomNumberBufferPtr, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.","."); goto cleanup; } @@ -4512,13 +4525,15 @@ static Image* ComputeAddNoiseImage(const Image* inputImage, } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } - outputReady = MagickTrue; + cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue); if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel); if (inputImageBuffer!=NULL) clReleaseMemObject(inputImageBuffer); @@ -4581,7 +4596,7 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, inputPixels = AcquirePixelCachePixels(inputImage, &length, exception); if (inputPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",inputImage->filename); goto cleanup; } @@ -4598,7 +4613,7 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, inputImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -4607,13 +4622,13 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, assert(filteredImage != NULL); if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } filteredPixels = GetPixelCachePixels(filteredImage, &length, exception); if (filteredPixels == (void *) NULL) { - (void) ThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); goto cleanup; } @@ -4632,7 +4647,7 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, filteredImageBuffer = clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } @@ -4681,14 +4696,14 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, , numRandomNumberGenerators * 4 * sizeof(unsigned long), NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clCreateBuffer failed.","."); goto cleanup; } seeds = (unsigned long*) clEnqueueMapBuffer(queue, randomNumberSeedsBuffer, CL_TRUE, CL_MAP_WRITE, 0 , numRandomNumberGenerators*4*sizeof(unsigned long), 0, NULL, NULL, &clStatus); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueMapBuffer failed.","."); goto cleanup; } @@ -4706,7 +4721,7 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, clStatus = clEnqueueUnmapMemObject(queue, randomNumberSeedsBuffer, seeds, 0, NULL, NULL); if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.","."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnqueueUnmapMemObject failed.","."); goto cleanup; } @@ -4774,13 +4789,15 @@ static Image* ComputeAddNoiseImageOptRandomNum(const Image* inputImage, } if (clStatus != CL_SUCCESS) { - (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Reading output image from CL buffer failed.", "'%s'", "."); + (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } - outputReady = MagickTrue; + cleanup: + OpenCLLogException(__FUNCTION__,__LINE__,exception); + if (queue!=NULL) RelinquishOpenCLCommandQueue(clEnv, queue); if (addNoiseKernel!=NULL) RelinquishOpenCLKernel(clEnv, addNoiseKernel); if (randomNumberGeneratorKernel!=NULL) RelinquishOpenCLKernel(clEnv, randomNumberGeneratorKernel); @@ -4817,14 +4834,11 @@ Image* AccelerateAddNoiseImage(const Image *image, const ChannelType channel, if (status == 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); - OpenCLLogException(__FUNCTION__,__LINE__,exception); return filteredImage; } diff --git a/MagickCore/opencl.c b/MagickCore/opencl.c index 34ce74213..10ac9a685 100644 --- a/MagickCore/opencl.c +++ b/MagickCore/opencl.c @@ -384,7 +384,7 @@ MagickExport MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param , size_t dataSize, void* data, ExceptionInfo* exception) { - MagickBooleanType + MagickBooleanType status; magick_unreferenced(exception); @@ -1870,6 +1870,8 @@ cleanup: return status; } + +#if 0 static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) { unsigned int i; if (profile == NULL || num==NULL) @@ -1877,11 +1879,12 @@ static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* n *num=0; for (i = 0; i < profile->numDevices; i++) { if (profile->devices[i].score == NULL) { - *num++; + (*num)++; } } return DS_SUCCESS; } +#endif /* End of the OpenCL device selection infrastructure @@ -2108,10 +2111,15 @@ static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exce unsigned int bestDeviceIndex; AccelerateScoreType bestScore; char path[MaxTextExtent]; - + MagickBooleanType flag; LockDefaultOpenCLEnv(); + /* Initially, just set OpenCL to off */ + flag = MagickTrue; + SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED + , sizeof(MagickBooleanType), &flag, exception); + status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION); if (status!=DS_SUCCESS) { (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", "."); @@ -2149,12 +2157,15 @@ static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exce /* set up clEnv with the best device */ if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) { /* CPU device */ - MagickBooleanType flag = MagickTrue; + flag = MagickTrue; SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED , sizeof(MagickBooleanType), &flag, exception); } else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) { /* OpenCL device */ + flag = MagickFalse; + SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED + , sizeof(MagickBooleanType), &flag, exception); SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception); } @@ -2272,6 +2283,58 @@ MagickExport MagickBooleanType InitImageMagickOpenCL( } +MagickExport +MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception, + const char *module,const char *function,const size_t line, + const ExceptionType severity,const char *tag,const char *format,...) { + MagickBooleanType + status; + + MagickCLEnv clEnv; + + status = MagickTrue; + + clEnv = GetDefaultOpenCLEnv(); + + assert(exception != (ExceptionInfo *) NULL); + assert(exception->signature == MagickSignature); + + if (severity!=0) { + cl_device_type dType; + clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL); + if (dType == CL_DEVICE_TYPE_CPU) { + char buffer[MaxTextExtent]; + clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL); + + /* Workaround for Intel OpenCL CPU runtime bug */ + /* Turn off OpenCL when a problem is detected! */ + if (strncmp(buffer, "Intel",5) == 0) { + + InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception); + } + } + } + +#ifdef OPENCLLOG_ENABLED + { + va_list + operands; + va_start(operands,format); + status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands); + va_end(operands); + } +#else + magick_unreferenced(module); + magick_unreferenced(function); + magick_unreferenced(line); + magick_unreferenced(tag); + magick_unreferenced(format); +#endif + + return(status); +} + + #else struct _MagickCLEnv { @@ -2404,6 +2467,21 @@ MagickExport MagickBooleanType InitImageMagickOpenCL( return MagickFalse; } + +MagickExport +MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception, + const char *module,const char *function,const size_t line, + const ExceptionType severity,const char *tag,const char *format,...) +{ + magick_unreferenced(exception); + magick_unreferenced(module); + magick_unreferenced(function); + magick_unreferenced(line); + magick_unreferenced(severity); + magick_unreferenced(tag); + magick_unreferenced(format); + return(MagickFalse); +} #endif /* MAGICKCORE_OPENCL_SUPPORT */ char* openclCachedFilesDirectory; @@ -2471,6 +2549,11 @@ void OpenCLLog(const char* message) { { if (message) { char path[MaxTextExtent]; + unsigned long allocSize; + + MagickCLEnv clEnv; + + clEnv = GetDefaultOpenCLEnv(); /* dump the source into a file */ (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s" @@ -2481,6 +2564,13 @@ void OpenCLLog(const char* message) { log = fopen(path, "ab"); fwrite(message, sizeof(char), strlen(message), log); fwrite("\n", sizeof(char), 1, log); + + if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled) + { + allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv); + fprintf(log, "Devic Max Memory Alloc Size: %ld\n", allocSize); + } + fclose(log); } } @@ -2488,4 +2578,3 @@ void OpenCLLog(const char* message) { magick_unreferenced(message); #endif } - -- 2.40.0