X-Git-Url: https://granicus.if.org/sourcecode?a=blobdiff_plain;f=MagickCore%2Faccelerate.c;h=ccf9c7a24374b86992a6a1b9d24406d9eb3a656a;hb=9f027d12a04d07489f8d76f4c2229229130bd10e;hp=0c0f394b90508637da0d26ab4ac010ea1588ece3;hpb=c8523c11bbb8e0ac3f1f9bfae600a997d6dd04a4;p=imagemagick diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 0c0f394b9..ccf9c7a24 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -182,22 +182,21 @@ static char " return(offset);\n" "}\n" "\n" - "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n" - "static inline CLQuantum ClampToQuantum(const double value)\n" + "static inline CLQuantum ClampToQuantum(const float value)\n" "{\n" "#if defined(MAGICKCORE_HDRI_SUPPORT)\n" " return((CLQuantum) value)\n" "#else\n" " if (value < 0.0)\n" " return((CLQuantum) 0);\n" - " if (value >= (double) QuantumRange)\n" + " if (value >= (float) QuantumRange)\n" " return((CLQuantum) QuantumRange);\n" " return((CLQuantum) (value+0.5));\n" "#endif\n" "}\n" "\n" "__kernel void Convolve(const __global CLPixelType *input,\n" - " __constant double *filter,const unsigned long width,const unsigned long height,\n" + " __constant float *filter,const unsigned long width,const unsigned long height,\n" " const unsigned int matte,__global CLPixelType *output)\n" "{\n" " const unsigned long columns = get_global_size(0);\n" @@ -206,11 +205,11 @@ static char " const long x = get_global_id(0);\n" " const long y = get_global_id(1);\n" "\n" - " const double scale = (1.0/QuantumRange);\n" + " const float scale = (1.0/QuantumRange);\n" " const long mid_width = (width-1)/2;\n" " const long mid_height = (height-1)/2;\n" - " double4 sum = { 0.0, 0.0, 0.0, 0.0 };\n" - " double gamma = 0.0;\n" + " float4 sum = { 0.0, 0.0, 0.0, 0.0 };\n" + " float gamma = 0.0;\n" " register unsigned long i = 0;\n" "\n" " int method = 0;\n" @@ -250,7 +249,7 @@ static char " {\n" " const unsigned long index=ClampToCanvas(y+v,rows)*columns+\n" " ClampToCanvas(x+u,columns);\n" - " const double alpha=scale*input[index].w;\n" + " const float alpha=scale*input[index].w;\n" " sum.x+=alpha*filter[i]*input[index].x;\n" " sum.y+=alpha*filter[i]*input[index].y;\n" " sum.z+=alpha*filter[i]*input[index].z;\n" @@ -284,7 +283,7 @@ static char " for (long u=(-mid_width); u <= mid_width; u++)\n" " {\n" " const unsigned long index=(y+v)*columns+(x+u);\n" - " const double alpha=scale*input[index].w;\n" + " const float alpha=scale*input[index].w;\n" " sum.x+=alpha*filter[i]*input[index].x;\n" " sum.y+=alpha*filter[i]*input[index].y;\n" " sum.z+=alpha*filter[i]*input[index].z;\n" @@ -321,7 +320,7 @@ static void ConvolveNotify(const char *message,const void *data,size_t length, } static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info, - const Image *image,const void *pixels,double *filter,const size_t width, + const Image *image,const void *pixels,float *filter,const size_t width, const size_t height,void *convolve_pixels) { cl_int @@ -344,7 +343,7 @@ static MagickBooleanType BindConvolveParameters(ConvolveInfo *convolve_info, return(MagickFalse); length=width*height; convolve_info->filter=clCreateBuffer(convolve_info->context,(cl_mem_flags) - (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(cl_double),filter, + (CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR),length*sizeof(cl_float),filter, &status); if ((convolve_info->filter == (cl_mem) NULL) || (status != CL_SUCCESS)) return(MagickFalse); @@ -397,7 +396,7 @@ static void DestroyConvolveBuffers(ConvolveInfo *convolve_info) cl_int status; - (void) status; + status=0; if (convolve_info->convolve_pixels != (cl_mem) NULL) status=clReleaseMemObject(convolve_info->convolve_pixels); if (convolve_info->pixels != (cl_mem) NULL) @@ -411,7 +410,7 @@ static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info) cl_int status; - (void) status; + status=0; if (convolve_info->kernel != (cl_kernel) NULL) status=clReleaseKernel(convolve_info->kernel); if (convolve_info->program != (cl_program) NULL) @@ -425,7 +424,7 @@ static ConvolveInfo *DestroyConvolveInfo(ConvolveInfo *convolve_info) } static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info, - const Image *image,const void *pixels,double *filter,const size_t width, + const Image *image,const void *pixels,float *filter,const size_t width, const size_t height,void *convolve_pixels) { cl_int @@ -441,7 +440,7 @@ static MagickBooleanType EnqueueConvolveKernel(ConvolveInfo *convolve_info, NULL); length=width*height; status=clEnqueueWriteBuffer(convolve_info->command_queue, - convolve_info->filter,CL_TRUE,0,length*sizeof(cl_double),filter,0,NULL, + convolve_info->filter,CL_TRUE,0,length*sizeof(cl_float),filter,0,NULL, NULL); if (status != CL_SUCCESS) return(MagickFalse); @@ -502,8 +501,8 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name, /* Create OpenCL context. */ - status=clGetPlatformIDs(0,NULL,&number_platforms); - if (status == CL_SUCCESS) + status=clGetPlatformIDs(0,(cl_platform_id *) NULL,&number_platforms); + if ((status == CL_SUCCESS) && (number_platforms > 0)) status=clGetPlatformIDs(1,platforms,NULL); if (status != CL_SUCCESS) { @@ -576,7 +575,7 @@ static ConvolveInfo *GetConvolveInfo(const Image *image,const char *name, convolve_info=DestroyConvolveInfo(convolve_info); return((ConvolveInfo *) NULL); } - (void) FormatLocaleString(options,MaxTextExtent,CLOptions,(double) + (void) FormatLocaleString(options,MaxTextExtent,CLOptions,(float) QuantumRange,MagickEpsilon); status=clBuildProgram(convolve_info->program,1,convolve_info->devices,options, NULL,NULL); @@ -630,9 +629,12 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, assert(exception->signature == MagickSignature); if ((image->storage_class != DirectClass) || (image->colorspace == CMYKColorspace)) + return(MagickFalse); if ((GetImageVirtualPixelMethod(image) != UndefinedVirtualPixelMethod) && (GetImageVirtualPixelMethod(image) != EdgeVirtualPixelMethod)) return(MagickFalse); + if (GetPixelChannels(image) != 4) + return(MagickFalse); #if !defined(MAGICKCORE_OPENCL_SUPPORT) return(MagickFalse); #else @@ -640,6 +642,9 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, const void *pixels; + float + *filter; + ConvolveInfo *convolve_info; @@ -649,6 +654,9 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, MagickSizeType length; + register ssize_t + i; + void *convolve_pixels; @@ -658,29 +666,43 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, pixels=AcquirePixelCachePixels(image,&length,exception); if (pixels == (const void *) NULL) { + convolve_info=DestroyConvolveInfo(convolve_info); (void) ThrowMagickException(exception,GetMagickModule(),CacheError, "UnableToReadPixelCache","`%s'",image->filename); - convolve_info=DestroyConvolveInfo(convolve_info); return(MagickFalse); } convolve_pixels=GetPixelCachePixels(convolve_image,&length,exception); if (convolve_pixels == (void *) NULL) { + convolve_info=DestroyConvolveInfo(convolve_info); (void) ThrowMagickException(exception,GetMagickModule(),CacheError, "UnableToReadPixelCache","`%s'",image->filename); + return(MagickFalse); + } + filter=(float *) AcquireQuantumMemory(kernel->width,kernel->height* + sizeof(*filter)); + if (filter == (float *) NULL) + { + DestroyConvolveBuffers(convolve_info); convolve_info=DestroyConvolveInfo(convolve_info); + (void) ThrowMagickException(exception,GetMagickModule(), + ResourceLimitError,"MemoryAllocationFailed","`%s'",image->filename); return(MagickFalse); } - status=BindConvolveParameters(convolve_info,image,pixels,kernel->values, + for (i=0; i < (ssize_t) (kernel->width*kernel->height); i++) + filter[i]=(float) kernel->values[i]; + status=BindConvolveParameters(convolve_info,image,pixels,filter, kernel->width,kernel->height,convolve_pixels); if (status == MagickFalse) { + filter=(float *) RelinquishMagickMemory(filter); DestroyConvolveBuffers(convolve_info); convolve_info=DestroyConvolveInfo(convolve_info); return(MagickFalse); } - status=EnqueueConvolveKernel(convolve_info,image,pixels,kernel->values, + status=EnqueueConvolveKernel(convolve_info,image,pixels,filter, kernel->width,kernel->height,convolve_pixels); + filter=(float *) RelinquishMagickMemory(filter); if (status == MagickFalse) { DestroyConvolveBuffers(convolve_info);