From: dirk Date: Thu, 28 Aug 2014 20:13:59 +0000 (+0000) Subject: Removed AccelerateNegateImageChannel due to performance issues. X-Git-Tag: 7.0.1-0~2056 X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=bae8263c35cce0298e343817ec94331bc4b955a6;p=imagemagick Removed AccelerateNegateImageChannel due to performance issues. Also removed deprecated AccelerateConvoleImage method. --- diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index 2b4547b9c..fc7f0cfaa 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -1932,42 +1932,6 @@ const char* accelerateKernels = } ) - STRINGIFY( - __kernel void Negate(__global CLPixelType *im, - const ChannelType channel) - { - - const int x = get_global_id(0); - const int y = get_global_id(1); - const int columns = get_global_size(0); - const int c = x + y * columns; - - CLPixelType pixel = im[c]; - - CLQuantum - blue, - green, - red; - - red=getRed(pixel); - green=getGreen(pixel); - blue=getBlue(pixel); - - CLPixelType filteredPixel; - - if ((channel & RedChannel) !=0) - setRed(&filteredPixel, QuantumRange-red); - if ((channel & GreenChannel) !=0) - setGreen(&filteredPixel, QuantumRange-green); - if ((channel & BlueChannel) !=0) - setBlue(&filteredPixel, QuantumRange-blue); - - filteredPixel.w = pixel.w; - - im[c] = filteredPixel; - } - ) - STRINGIFY( __kernel void Grayscale(__global CLPixelType *im, const int method, const int colorspace) diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 54bcd8352..d61f5628f 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -4189,203 +4189,6 @@ MagickExport MagickBooleanType AccelerateModulateImage(Image *image, return(status); } -/* -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% % -% % -% % -% N e g a t e I m a g e w i t h O p e n C L % -% % -% % -% % -%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% -% -% -% A description of each parameter follows: -% -% o image: the image. -% -% o channel: the channel. -% -% o grayscale: If MagickTrue, only negate grayscale pixels within the image. -% -*/ - -MagickBooleanType ComputeNegateImageChannel(Image *image, - const ChannelType channel,const MagickBooleanType magick_unused(grayscale), - ExceptionInfo* exception) -{ - CacheView - *image_view; - - cl_context - context; - - cl_command_queue - queue; - - cl_int - clStatus; - - cl_kernel - negateKernel; - - cl_mem - imageBuffer; - - cl_mem_flags - mem_flags; - - MagickBooleanType - outputReady; - - MagickCLEnv - clEnv; - - MagickSizeType - length; - - register ssize_t - i; - - void - *inputPixels; - - magick_unreferenced(grayscale); - - inputPixels = NULL; - imageBuffer = NULL; - negateKernel = NULL; - - assert(image != (Image *) NULL); - assert(image->signature == MagickSignature); - if (image->debug != MagickFalse) - (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename); - - /* - * initialize opencl env - */ - clEnv = GetDefaultOpenCLEnv(); - context = GetOpenCLContext(clEnv); - queue = AcquireOpenCLCommandQueue(clEnv); - - outputReady = MagickFalse; - - /* Create and initialize OpenCL buffers. - inputPixels = AcquirePixelCachePixels(image, &length, exception); - assume this will get a writable image - */ - image_view=AcquireAuthenticCacheView(image,exception); - inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); - if (inputPixels == (void *) NULL) - { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); - goto cleanup; - } - - /* If the host pointer is aligned to the size of CLPixelPacket, - then use the host buffer directly from the GPU; otherwise, - create a buffer on the GPU and copy the data over - */ - if (ALIGNED(inputPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR; - } - else - { - mem_flags = CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR; - } - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - - negateKernel = AcquireOpenCLKernel(clEnv, MAGICK_OPENCL_ACCELERATE, "Negate"); - if (negateKernel == NULL) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", "."); - goto cleanup; - } - - i = 0; - clStatus=clEnv->library->clSetKernelArg(negateKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); - clStatus=clEnv->library->clSetKernelArg(negateKernel,i++,sizeof(ChannelType),(void *)&channel); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); - printf("no kernel\n"); - goto cleanup; - } - - { - size_t global_work_size[2]; - global_work_size[0] = image->columns; - global_work_size[1] = image->rows; - /* launch the kernel */ - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, negateKernel, 2, NULL, global_work_size, NULL, 0, NULL, NULL); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", "."); - goto cleanup; - } - clEnv->library->clFlush(queue); - } - - if (ALIGNED(inputPixels,CLPixelPacket)) - { - length = image->columns * image->rows; - clEnv->library->clEnqueueMapBuffer(queue, imageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); - } - else - { - length = image->columns * image->rows; - clStatus = clEnv->library->clEnqueueReadBuffer(queue, imageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), inputPixels, 0, NULL, NULL); - } - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); - goto cleanup; - } - - outputReady=SyncCacheViewAuthenticPixels(image_view,exception); - -cleanup: - OpenCLLogException(__FUNCTION__,__LINE__,exception); - - image_view=DestroyCacheView(image_view); - - if (imageBuffer!=NULL) - clEnv->library->clReleaseMemObject(imageBuffer); - if (negateKernel!=NULL) - RelinquishOpenCLKernel(clEnv, negateKernel); - if (queue != NULL) - RelinquishOpenCLCommandQueue(clEnv, queue); - - return(outputReady); -} - -MagickExport MagickBooleanType AccelerateNegateImageChannel(Image *image, - const ChannelType channel,const MagickBooleanType grayscale, - ExceptionInfo* exception) -{ - MagickBooleanType - status; - - assert(image != NULL); - assert(exception != (ExceptionInfo *) NULL); - - if ((checkOpenCLEnvironment(exception) == MagickFalse) || - (checkAccelerateCondition(image, channel) == MagickFalse)) - return(MagickFalse); - - status=ComputeNegateImageChannel(image,channel,grayscale,exception); - return(status); -} - /* %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%% % % @@ -7431,17 +7234,6 @@ MagickBooleanType AccelerateModulateImage( return(MagickFalse); } -MagickExport -MagickBooleanType AccelerateNegateImageChannel( - Image* image, const ChannelType channel, const MagickBooleanType grayscale, ExceptionInfo* exception) -{ - magick_unreferenced(image); - magick_unreferenced(channel); - magick_unreferenced(grayscale); - magick_unreferenced(exception); - return(MagickFalse); -} - MagickExport MagickBooleanType AccelerateGrayscaleImage( Image* image, const PixelIntensityMethod method, ExceptionInfo* exception) @@ -7485,18 +7277,4 @@ Image* AccelerateMotionBlurImage(const Image *image, const ChannelType channel, return NULL; } -#endif /* MAGICKCORE_OPENCL_SUPPORT */ - -MagickExport MagickBooleanType AccelerateConvolveImage( - const Image *magick_unused(image),const KernelInfo *magick_unused(kernel), - Image *magick_unused(convolve_image),ExceptionInfo *magick_unused(exception)) -{ - magick_unreferenced(image); - magick_unreferenced(kernel); - magick_unreferenced(convolve_image); - magick_unreferenced(exception); - - /* legacy, do not use */ - return(MagickFalse); -} - +#endif /* MAGICKCORE_OPENCL_SUPPORT */ \ No newline at end of file diff --git a/MagickCore/accelerate.h b/MagickCore/accelerate.h index 3916f1cf5..ba12d4f91 100644 --- a/MagickCore/accelerate.h +++ b/MagickCore/accelerate.h @@ -32,8 +32,6 @@ extern MagickExport MagickBooleanType AccelerateCompositeImage(Image *,const ChannelType,const CompositeOperator, const Image *,const ssize_t,const ssize_t,const float,const float,ExceptionInfo *), AccelerateContrastImage(Image *,const MagickBooleanType,ExceptionInfo *), - AccelerateConvolveImage(const Image *,const KernelInfo *,Image *, - ExceptionInfo *), AccelerateContrastStretchImageChannel(Image *, const ChannelType, const double, const double, ExceptionInfo*), AccelerateEqualizeImage(Image *,const ChannelType,ExceptionInfo *), @@ -43,8 +41,6 @@ extern MagickExport MagickBooleanType ExceptionInfo *), AccelerateModulateImage(Image*, double, double, double, ColorspaceType, ExceptionInfo*), - AccelerateNegateImageChannel(Image*, const ChannelType, const MagickBooleanType, - ExceptionInfo *), AccelerateRandomImage(Image*, ExceptionInfo*); extern MagickExport Image