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);
}
)
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;
cl_mem
imageBuffer;
- cl_mem_flags
- mem_flags;
+ cl_uint
+ number_channels;
MagickBooleanType
outputReady;
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);
assert(image != NULL);
assert(exception != (ExceptionInfo *) NULL);
- if (checkAccelerateConditionRGBA(image) == MagickFalse)
+ if (checkAccelerateCondition(image) == MagickFalse)
return(MagickFalse);
clEnv=getOpenCLEnvironment(exception);