From 83683f454d1393bb6b568bfbbff50caffd9101f7 Mon Sep 17 00:00:00 2001 From: dirk Date: Fri, 8 Apr 2016 02:16:24 +0200 Subject: [PATCH] AccelerateRotationalBlurImage now supports R/RA/RGB images. --- MagickCore/accelerate-private.h | 81 ++++++++++---------- MagickCore/accelerate.c | 130 ++++++++------------------------ 2 files changed, 70 insertions(+), 141 deletions(-) diff --git a/MagickCore/accelerate-private.h b/MagickCore/accelerate-private.h index fa45cbe30..c72d08ee1 100644 --- a/MagickCore/accelerate-private.h +++ b/MagickCore/accelerate-private.h @@ -2987,12 +2987,9 @@ STRINGIFY( */ STRINGIFY( - __kernel void RotationalBlur(const __global CLPixelType *im, __global CLPixelType *filtered_im, - const float4 bias, - const unsigned int channel, const unsigned int matte, - const float2 blurCenter, - __constant float *cos_theta, __constant float *sin_theta, - const unsigned int cossin_theta_size) + __kernel void RotationalBlur(const __global CLQuantum *image,const unsigned int number_channels, + const unsigned int channel,const float4 bias,const float2 blurCenter,__constant float *cos_theta, + __constant float *sin_theta,const unsigned int cossin_theta_size,__global CLQuantum *filteredImage) { const int x = get_global_id(0); const int y = get_global_id(1); @@ -3014,49 +3011,49 @@ STRINGIFY( step = cossin_theta_size-1; } - float4 result; - result.x = (float)bias.x; - result.y = (float)bias.y; - result.z = (float)bias.z; - result.w = (float)bias.w; + float4 result = bias; + float normalize = 0.0f; + float gamma = 0.0f; + + for (unsigned int i=0; icolumns,image->rows,exception); - if (inputPixels == (const void *) NULL) - { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); - goto cleanup; - } - - /* If the host pointer is aligned to the size of CLPixelPacket, - then use the host buffer directly from the GPU; otherwise, - create a buffer on the GPU and copy the data over */ - if (ALIGNED(inputPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR; - } - else - { - mem_flags = CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR; - } - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - imageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + image_view=AcquireVirtualCacheView(image, exception); + imageBuffer=createReadBuffer(image,image_view,clEnv,context,exception); + if (imageBuffer == (cl_mem) NULL) goto cleanup; - } - filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception); - assert(filteredImage != NULL); - if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); + if (filteredImage == (Image *) NULL) goto cleanup; - } - filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception); - filteredPixels=GetCacheViewAuthenticPixels(filteredImage_view,0,0,filteredImage->columns,filteredImage->rows,exception); - if (filteredPixels == (void *) NULL) + if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue) { - (void) OpenCLThrowMagickException(exception,GetMagickModule(),CacheWarning, "UnableToReadPixelCache.","`%s'",filteredImage->filename); + (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", "."); goto cleanup; } - - if (ALIGNED(filteredPixels,CLPixelPacket)) - { - mem_flags = CL_MEM_WRITE_ONLY|CL_MEM_USE_HOST_PTR; - hostPtr = filteredPixels; - } - else - { - mem_flags = CL_MEM_WRITE_ONLY; - hostPtr = NULL; - } - /* create a CL buffer from image pixel buffer */ - length = image->columns * image->rows; - filteredImageBuffer = clEnv->library->clCreateBuffer(context, mem_flags, length * sizeof(CLPixelPacket), hostPtr, &clStatus); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); + filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception); + filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv, + context,filteredPixels,exception); + if (filteredImageBuffer == (cl_mem) NULL) goto cleanup; - } blurCenter.s[0] = (float) (image->columns-1)/2.0; blurCenter.s[1] = (float) (image->rows-1)/2.0; @@ -5662,35 +5607,31 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, goto cleanup; } - - /* set the kernel arguments */ - i = 0; - clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); - clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); - GetPixelInfo(image,&bias); biasPixel.s[0] = bias.red; biasPixel.s[1] = bias.green; biasPixel.s[2] = bias.blue; biasPixel.s[3] = bias.alpha; - clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel); - clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &image->channel_mask); - - matte = (image->alpha_trait > CopyPixelTrait)?1:0; - clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &matte); - clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter); + number_channels = image->number_channels; + /* set the kernel arguments */ + i = 0; + clStatus=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); + clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_uint),&number_channels); + clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &image->channel_mask); + clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel); + clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter); clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer); clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer); - clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(unsigned int), &cossin_theta_size); + clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_uint), &cossin_theta_size); + clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer); if (clStatus != CL_SUCCESS) { (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", "."); goto cleanup; } - global_work_size[0] = image->columns; global_work_size[1] = image->rows; /* launch the kernel */ @@ -5704,21 +5645,12 @@ static Image* ComputeRotationalBlurImage(const Image *image,const double angle, RecordProfileData(clEnv,RotationalBlurKernel,event); clEnv->library->clReleaseEvent(event); - if (ALIGNED(filteredPixels,CLPixelPacket)) - { - length = image->columns * image->rows; - clEnv->library->clEnqueueMapBuffer(queue, filteredImageBuffer, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, length * sizeof(CLPixelPacket), 0, NULL, NULL, &clStatus); - } - else - { - length = image->columns * image->rows; - clStatus = clEnv->library->clEnqueueReadBuffer(queue, filteredImageBuffer, CL_TRUE, 0, length * sizeof(CLPixelPacket), filteredPixels, 0, NULL, NULL); - } - if (clStatus != CL_SUCCESS) + if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) { - (void) OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); + (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; } + outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception); cleanup: @@ -5754,7 +5686,7 @@ MagickExport Image* AccelerateRotationalBlurImage(const Image *image, assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if ((checkAccelerateConditionRGBA(image) == MagickFalse) || + if ((checkAccelerateCondition(image) == MagickFalse) || (checkOpenCLEnvironment(exception) == MagickFalse)) return NULL; @@ -6366,7 +6298,7 @@ static Image *ComputeWaveletDenoiseImage(const Image *image, clEnv->library->clReleaseEvent(event); } - if (copyWriteBuffer(image,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) + if (copyWriteBuffer(filteredImage,clEnv,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse) { (void)OpenCLThrowMagickException(exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", "."); goto cleanup; -- 2.40.0