AccelerateContrastImage now supports R/RA/RGB images.

diff --git a/MagickCore/accelerate-kernels-private.h b/MagickCore/accelerate-kernels-private.h
index a97c75a707a0bae548f403f79a4e855e9cc0bb2b..5ed4da4c39f99d6c7fd86ee2814c82c5122626f5 100644 (file)
--- 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 192633ee45e1d3e33c00f71e354cb8151b173bd9..23c7d8f81c5af8f1da205f734ac0c0addd67d82f 100644 (file)
--- 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);