From: dirk Date: Sun, 24 Jul 2016 20:46:29 +0000 (+0200) Subject: AccelerateContrastImage now supports R/RA/RGB images. X-Git-Tag: 7.0.2-6~8 X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=4dc3fbb71651342ec1f9aeb1fce2f93f7e70f9f4;p=imagemagick AccelerateContrastImage now supports R/RA/RGB images. --- diff --git a/MagickCore/accelerate-kernels-private.h b/MagickCore/accelerate-kernels-private.h index a97c75a70..5ed4da4c3 100644 --- a/MagickCore/accelerate-kernels-private.h +++ b/MagickCore/accelerate-kernels-private.h @@ -1025,97 +1025,110 @@ OPENCL_ENDIF() STRINGIFY( - inline float3 ConvertRGBToHSB(CLPixelType pixel) { - float3 HueSaturationBrightness; - HueSaturationBrightness.x = 0.0f; // Hue - HueSaturationBrightness.y = 0.0f; // Saturation - HueSaturationBrightness.z = 0.0f; // Brightness - - float r=(float) getRed(pixel); - float g=(float) getGreen(pixel); - float b=(float) getBlue(pixel); - - float tmin=MagickMin(MagickMin(r,g),b); - float tmax=MagickMax(MagickMax(r,g),b); - - if (tmax!=0.0f) { + inline float4 ConvertRGBToHSB(const float4 pixel) + { + float4 result=0.0f; + result.w=pixel.w; + float tmax=MagickMax(MagickMax(pixel.x,pixel.y),pixel.z); + if (tmax != 0.0f) + { + float tmin=MagickMin(MagickMin(pixel.x,pixel.y),pixel.z); float delta=tmax-tmin; - HueSaturationBrightness.y=delta/tmax; - HueSaturationBrightness.z=QuantumScale*tmax; - - if (delta != 0.0f) { - HueSaturationBrightness.x = ((r == tmax)?0.0f:((g == tmax)?2.0f:4.0f)); - HueSaturationBrightness.x += ((r == tmax)?(g-b):((g == tmax)?(b-r):(r-g)))/delta; - HueSaturationBrightness.x/=6.0f; - HueSaturationBrightness.x += (HueSaturationBrightness.x < 0.0f)?0.0f:1.0f; + + result.y=delta/tmax; + result.z=QuantumScale*tmax; + if (delta != 0.0f) + { + result.x =((pixel.x == tmax) ? 0.0f : ((pixel.y == tmax) ? 2.0f : 4.0f)); + result.x+=((pixel.x == tmax) ? (pixel.y-pixel.z) : ((pixel.y == tmax) ? + (pixel.z-pixel.x) : (pixel.x-pixel.y)))/delta; + result.x/=6.0f; + result.x+=(result.x < 0.0f) ? 0.0f : 1.0f; } } - return HueSaturationBrightness; + return(result); } - inline CLPixelType ConvertHSBToRGB(float3 HueSaturationBrightness) { + inline float4 ConvertHSBToRGB(const float4 pixel) + { + float hue=pixel.x; + float brightness=pixel.z; + float saturation=pixel.y; - float hue = HueSaturationBrightness.x; - float brightness = HueSaturationBrightness.z; - float saturation = HueSaturationBrightness.y; - - CLPixelType rgb; + float4 result=pixel; - if (saturation == 0.0f) { - setRed(&rgb,ClampToQuantum(QuantumRange*brightness)); - setGreen(&rgb,getRed(rgb)); - setBlue(&rgb,getRed(rgb)); + if (saturation == 0.0f) + { + result.x=result.y=result.z=ClampToQuantum(QuantumRange*brightness); } - else { - + else + { float h=6.0f*(hue-floor(hue)); float f=h-floor(h); float p=brightness*(1.0f-saturation); float q=brightness*(1.0f-saturation*f); float t=brightness*(1.0f-(saturation*(1.0f-f))); - - float clampedBrightness = ClampToQuantum(QuantumRange*brightness); - float clamped_t = ClampToQuantum(QuantumRange*t); - float clamped_p = ClampToQuantum(QuantumRange*p); - float clamped_q = ClampToQuantum(QuantumRange*q); int ih = (int)h; - setRed(&rgb, (ih == 1)?clamped_q: - (ih == 2 || ih == 3)?clamped_p: - (ih == 4)?clamped_t: - clampedBrightness); - - setGreen(&rgb, (ih == 1 || ih == 2)?clampedBrightness: - (ih == 3)?clamped_q: - (ih == 4 || ih == 5)?clamped_p: - clamped_t); - - setBlue(&rgb, (ih == 2)?clamped_t: - (ih == 3 || ih == 4)?clampedBrightness: - (ih == 5)?clamped_q: - clamped_p); + + if (ih == 1) + { + result.x=ClampToQuantum(QuantumRange*q); + result.y=ClampToQuantum(QuantumRange*brightness); + result.z=ClampToQuantum(QuantumRange*p); + } + else if (ih == 2) + { + result.x=ClampToQuantum(QuantumRange*p); + result.y=ClampToQuantum(QuantumRange*brightness); + result.z=ClampToQuantum(QuantumRange*t); + } + else if (ih == 3) + { + result.x=ClampToQuantum(QuantumRange*p); + result.y=ClampToQuantum(QuantumRange*q); + result.z=ClampToQuantum(QuantumRange*brightness); + } + else if (ih == 4) + { + result.x=ClampToQuantum(QuantumRange*t); + result.y=ClampToQuantum(QuantumRange*p); + result.z=ClampToQuantum(QuantumRange*brightness); + } + else if (ih == 5) + { + result.x=ClampToQuantum(QuantumRange*brightness); + result.y=ClampToQuantum(QuantumRange*p); + result.z=ClampToQuantum(QuantumRange*q); + } + else + { + result.x=ClampToQuantum(QuantumRange*q); + result.y=ClampToQuantum(QuantumRange*brightness); + result.z=ClampToQuantum(QuantumRange*p); + } } - return rgb; + return(result); } - __kernel void Contrast(__global CLPixelType *im, const unsigned int sharpen) + __kernel void Contrast(__global CLQuantum *image,const int sign, + const unsigned int number_channels) { + const int x=get_global_id(0); + const int y=get_global_id(1); + const unsigned int columns=get_global_size(0); - const int sign = sharpen!=0?1:-1; - 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; + float4 pixel=ReadAllChannels(image,number_channels,columns,x,y); + if (number_channels < 3) + pixel.y=pixel.z=pixel.x; - CLPixelType pixel = im[c]; - float3 HueSaturationBrightness = ConvertRGBToHSB(pixel); - float brightness = HueSaturationBrightness.z; + pixel=ConvertRGBToHSB(pixel); + float brightness=pixel.z; brightness+=0.5f*sign*(0.5f*(sinpi(brightness-0.5f)+1.0f)-brightness); - brightness = clamp(brightness,0.0f,1.0f); - HueSaturationBrightness.z = brightness; + brightness=clamp(brightness,0.0f,1.0f); + pixel.z=brightness; + pixel=ConvertHSBToRGB(pixel); - CLPixelType filteredPixel = ConvertHSBToRGB(HueSaturationBrightness); - filteredPixel.w = pixel.w; - im[c] = filteredPixel; + WriteAllChannels(image,number_channels,columns,x,y,pixel); } ) diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c index 192633ee4..23c7d8f81 100644 --- a/MagickCore/accelerate.c +++ b/MagickCore/accelerate.c @@ -793,17 +793,15 @@ MagickPrivate Image* AccelerateBlurImage(const Image *image, static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv, const MagickBooleanType sharpen,ExceptionInfo *exception) { - CacheView - *image_view; - cl_command_queue queue; cl_int - clStatus; + status, + sign; cl_kernel - filterKernel; + contrastKernel; cl_event event; @@ -811,8 +809,8 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv, cl_mem imageBuffer; - cl_mem_flags - mem_flags; + cl_uint + number_channels; MagickBooleanType outputReady; @@ -820,114 +818,54 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv, MagickCLDevice device; - MagickSizeType - length; - size_t - global_work_size[2]; - - unsigned int - i, - uSharpen; - - void - *inputPixels; - - outputReady = MagickFalse; - inputPixels = NULL; - imageBuffer = NULL; - filterKernel = NULL; - queue = NULL; + gsize[2], + i; - device = RequestOpenCLDevice(clEnv); + contrastKernel=NULL; + outputReady=MagickFalse; - /* Create and initialize OpenCL buffers. */ - image_view=AcquireAuthenticCacheView(image,exception); - inputPixels=GetCacheViewAuthenticPixels(image_view,0,0,image->columns,image->rows,exception); - if (inputPixels == (void *) NULL) - { - (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename); + device=RequestOpenCLDevice(clEnv); + queue=AcquireOpenCLCommandQueue(device); + imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception); + if (imageBuffer == (cl_mem) NULL) 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(device->context, mem_flags, length * sizeof(CLPixelPacket), (void*)inputPixels, &clStatus); - if (clStatus != CL_SUCCESS) + contrastKernel=AcquireOpenCLKernel(device,"Contrast"); + if (contrastKernel == (cl_kernel) NULL) { - (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.","."); - goto cleanup; - } - - filterKernel = AcquireOpenCLKernel(device,"Contrast"); - if (filterKernel == NULL) - { - (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "."); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), + ResourceLimitWarning,"AcquireOpenCLKernel failed.","."); goto cleanup; } - i = 0; - clStatus=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); + number_channels=(cl_uint) image->number_channels; + sign=sharpen != MagickFalse ? 1 : -1; - uSharpen = (sharpen == MagickFalse)?0:1; - clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen); - if (clStatus != CL_SUCCESS) + i=0; + status =SetOpenCLKernelArg(contrastKernel,i++,sizeof(cl_mem),(void *)&imageBuffer); + status|=SetOpenCLKernelArg(contrastKernel,i++,sizeof(cl_uint),&number_channels); + status|=SetOpenCLKernelArg(contrastKernel,i++,sizeof(cl_int),&sign); + if (status != CL_SUCCESS) { - (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "."); + (void) OpenCLThrowMagickException(device,exception,GetMagickModule(), + ResourceLimitWarning,"SetOpenCLKernelArg failed.","."); goto cleanup; } - global_work_size[0] = image->columns; - global_work_size[1] = image->rows; - /* launch the kernel */ - queue = AcquireOpenCLCommandQueue(device); - clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, filterKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event); - if (clStatus != CL_SUCCESS) - { - (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "."); - goto cleanup; - } - RecordProfileData(device,filterKernel,event); + gsize[0]=image->columns; + gsize[1]=image->rows; - 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(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "."); - goto cleanup; - } - outputReady=SyncCacheViewAuthenticPixels(image_view,exception); + outputReady=EnqueueOpenCLKernel(queue,contrastKernel,2,(const size_t *) NULL, + gsize,(const size_t *) NULL,image,(Image *) NULL,exception); cleanup: - image_view=DestroyCacheView(image_view); - - if (imageBuffer!=NULL) - clEnv->library->clReleaseMemObject(imageBuffer); - if (filterKernel!=NULL) - ReleaseOpenCLKernel(filterKernel); - if (queue != NULL) + if (contrastKernel != (cl_kernel) NULL) + ReleaseOpenCLKernel(contrastKernel); + if (queue != (cl_command_queue) NULL) ReleaseOpenCLCommandQueue(device,queue); - if (device != NULL) + if (device != (MagickCLDevice) NULL) ReleaseOpenCLDevice(device); return(outputReady); @@ -945,7 +883,7 @@ MagickPrivate MagickBooleanType AccelerateContrastImage(Image *image, assert(image != NULL); assert(exception != (ExceptionInfo *) NULL); - if (checkAccelerateConditionRGBA(image) == MagickFalse) + if (checkAccelerateCondition(image) == MagickFalse) return(MagickFalse); clEnv=getOpenCLEnvironment(exception);