From: cristy Date: Tue, 13 Sep 2011 00:31:55 +0000 (+0000) Subject: (no commit message) X-Git-Tag: 7.0.1-0~7016 X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=1a2e276c9e1be7859ac1c933ffdc00f9d91bbc02;p=imagemagick --- diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 0c0f394b9..141b606f0 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); @@ -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); @@ -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); @@ -640,6 +639,9 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, const void *pixels; + float + *filter; + ConvolveInfo *convolve_info; @@ -649,6 +651,9 @@ MagickExport MagickBooleanType AccelerateConvolveImage(const Image *image, MagickSizeType length; + register ssize_t + i; + void *convolve_pixels; @@ -658,29 +663,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);