]> granicus.if.org Git - imagemagick/commitdiff
Some OpenCL methods are now executed asynchronous.
authordirk <dirk@git.imagemagick.org>
Mon, 13 Jun 2016 20:30:26 +0000 (22:30 +0200)
committerdirk <dirk@git.imagemagick.org>
Mon, 13 Jun 2016 20:30:26 +0000 (22:30 +0200)
Removed empty Accelerate methods that are used when OpenCL is disabled.

13 files changed:
MagickCore/accelerate-private.h
MagickCore/accelerate.c
MagickCore/cache-private.h
MagickCore/cache-view.c
MagickCore/cache.c
MagickCore/effect.c
MagickCore/enhance.c
MagickCore/fx.c
MagickCore/opencl-private.h
MagickCore/opencl.c
MagickCore/opencl.h
MagickCore/resize.c
MagickCore/statistic.c

index 46c1c975d56e3f254a23c01ea3a7c61670b8c58e..dd4303b55c02cdb5df43fc1a2d44c046073e47e8 100644 (file)
@@ -29,6 +29,8 @@
 extern "C" {
 #endif
 
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+
 extern MagickPrivate Image
   *AccelerateAddNoiseImage(const Image*,const NoiseType,ExceptionInfo *),
   *AccelerateBlurImage(const Image *,const double,const double,ExceptionInfo *),
@@ -59,8 +61,10 @@ extern MagickPrivate MagickBooleanType
   AccelerateModulateImage(Image *,const double,const double,const double,
     const ColorspaceType, ExceptionInfo*);
 
+#endif /* MAGICKCORE_OPENCL_SUPPORT */
+
 #if defined(__cplusplus) || defined(c_plusplus)
 }
 #endif
 
-#endif // MAGICKCORE_ACCELERATE_PRIVATE_H
+#endif /* MAGICKCORE_ACCELERATE_PRIVATE_H */
index 92daf1e8b89f06eef75116280e2a542908c18875..631e1dafdadb6f32e655f82c61ede9a2b135fcfe 100644 (file)
 %                              Software Design                                %
 %                                  Cristy                                     %
 %                               SiuChi Chan                                   %
-%                               Guansong Zhang                                %
+%                              Guansong Zhang                                 %
 %                               January 2010                                  %
+%                               Dirk Lemstra                                  %
+%                                April 2016                                   %
 %                                                                             %
 %                                                                             %
 %  Copyright 1999-2016 ImageMagick Studio LLC, a non-profit organization      %
@@ -232,111 +234,8 @@ inline static unsigned int padGlobalWorkgroupSizeToLocalWorkgroupSize(
   return ((orgGlobalSize+(localGroupSize-1))/localGroupSize*localGroupSize);
 }
 
-static cl_mem createBuffer(const Image *image,CacheView *image_view,
-  MagickCLEnv clEnv,MagickCLDevice device,cl_mem_flags flags,void *pixels,
-  ExceptionInfo *exception)
-{
-  cl_mem
-    buffer;
-
-  cl_mem_flags
-    mem_flags;
-
-  cl_int
-    status;
-
-  size_t
-    length;
-
-  void
-    *hostPtr;
-
-  pixels=(void *) GetCacheViewAuthenticPixels(image_view,0,0,image->columns,
-    image->rows,exception);
-  if (pixels == (void *) NULL)
-    {
-      (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
-        CacheWarning,"UnableToReadPixelCache.","`%s'",image->filename);
-      return (cl_mem) NULL;
-    }
-
-  mem_flags=flags;
-  hostPtr=pixels;
-  if (ALIGNED(pixels,CLQuantum))
-    mem_flags=mem_flags | CL_MEM_USE_HOST_PTR;
-  else if ((mem_flags == CL_MEM_READ_ONLY) || (mem_flags == CL_MEM_READ_WRITE))
-    mem_flags=mem_flags | CL_MEM_COPY_HOST_PTR;
-  else if (mem_flags == CL_MEM_WRITE_ONLY)
-    hostPtr=NULL;
-
-  length=image->columns*image->rows*image->number_channels;
-  buffer=clEnv->library->clCreateBuffer(device->context,mem_flags,length*
-    sizeof(CLQuantum),hostPtr,&status);
-  if (status != CL_SUCCESS)
-    {
-      (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
-        ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.",".");
-    }
-
-  return(buffer);
-}
-
-static inline cl_mem createReadBuffer(const Image *image,CacheView *image_view,
-  MagickCLEnv clEnv,MagickCLDevice device,ExceptionInfo *exception)
-{
-  void
-    *pixels;
-
-  pixels=(void *) NULL;
-  return(createBuffer(image,image_view,clEnv,device,CL_MEM_READ_ONLY,
-    pixels,exception));
-}
-
-static inline cl_mem createReadWriteBuffer(const Image *image,
-  CacheView *image_view,MagickCLEnv clEnv,MagickCLDevice device,void *pixels,
-  ExceptionInfo *exception)
-{
-  return(createBuffer(image,image_view,clEnv,device,CL_MEM_READ_WRITE,pixels,
-    exception));
-}
-
-static inline cl_mem createWriteBuffer(Image *image,CacheView *image_view,
-  MagickCLEnv clEnv,MagickCLDevice device,void *pixels,ExceptionInfo *exception)
-{
-  return(createBuffer(image,image_view,clEnv,device,CL_MEM_WRITE_ONLY,pixels,
-    exception));
-}
-
-static inline MagickBooleanType copyWriteBuffer(const Image *image,
-  MagickCLEnv clEnv,MagickCLDevice device,cl_command_queue queue,cl_mem buffer,
-  void *pixels,ExceptionInfo *exception)
-{
-  cl_int
-    status;
-
-  size_t
-    length;
-
-  length=image->columns*image->rows*image->number_channels*sizeof(CLQuantum);
-  if (ALIGNED(pixels,CLQuantum))
-    clEnv->library->clEnqueueMapBuffer(queue,buffer,CL_TRUE,CL_MAP_READ |
-      CL_MAP_WRITE,0,length,0,NULL,NULL,&status);
-  else 
-    status=clEnv->library->clEnqueueReadBuffer(queue,buffer,CL_TRUE,0,length,
-      pixels,0,NULL,NULL);
-  if (status != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
-      ResourceLimitWarning,"Reading output image from CL buffer failed.",
-      "'%s'",".");
-    return(MagickFalse);
-  }
-  return(MagickTrue);
-}
-
-static cl_mem createKernelInfo(MagickCLEnv clEnv,MagickCLDevice device,
-  cl_command_queue queue,const double radius,const double sigma,cl_uint *width,
-  ExceptionInfo *exception)
+static cl_mem createKernelInfo(MagickCLDevice device,const double radius,
+  const double sigma,cl_uint *width,ExceptionInfo *exception)
 {
   char
     geometry[MagickPathExtent];
@@ -353,7 +252,7 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,MagickCLDevice device,
   KernelInfo
     *kernel;
 
-  size_t
+  ssize_t
     i;
 
   (void) FormatLocaleString(geometry,MagickPathExtent,
@@ -365,44 +264,17 @@ static cl_mem createKernelInfo(MagickCLEnv clEnv,MagickCLDevice device,
       ResourceLimitWarning,"AcquireKernelInfo failed.",".");
     return((cl_mem) NULL);
   }
-
-  imageKernelBuffer=clEnv->library->clCreateBuffer(device->context,
-    CL_MEM_READ_ONLY,kernel->width*sizeof(float),NULL,&status);
-  if (status != CL_SUCCESS)
-  {
-    kernel=DestroyKernelInfo(kernel);
-    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
-      ResourceLimitWarning,"clEnv->library->clCreateBuffer failed.",".");
-    return((cl_mem) NULL);
-  }
-
-  kernelBufferPtr=(float*)clEnv->library->clEnqueueMapBuffer(queue,
-    imageKernelBuffer,CL_TRUE,CL_MAP_WRITE,0,kernel->width*sizeof(float),0,
-      NULL,NULL,&status);
-  if (status != CL_SUCCESS)
-  {
-    kernel=DestroyKernelInfo(kernel);
-    clEnv->library->clReleaseMemObject(imageKernelBuffer);
-    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
-      ResourceLimitWarning,"clEnv->library->clEnqueueMapBuffer failed.",".");
-    return((cl_mem) NULL);
-  }
-  for (i = 0; i < kernel->width; i++)
-    kernelBufferPtr[i]=(float)kernel->values[i];
-
-  *width=(cl_uint) kernel->width;
+  kernelBufferPtr=(float *)AcquireMagickMemory(kernel->width*sizeof(float));
+  for (i = 0; i < (ssize_t) kernel->width; i++)
+    kernelBufferPtr[i] = (float)kernel->values[i];
+  imageKernelBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
+    CL_MEM_COPY_HOST_PTR,kernel->width*sizeof(float),kernelBufferPtr);
+  *width=kernel->width;
+  kernelBufferPtr=(float *) RelinquishMagickMemory(kernelBufferPtr);
   kernel=DestroyKernelInfo(kernel);
-
-  status=clEnv->library->clEnqueueUnmapMemObject(queue,imageKernelBuffer,
-    kernelBufferPtr,0,NULL,NULL);
-  if (status != CL_SUCCESS)
-  {
-    clEnv->library->clReleaseMemObject(imageKernelBuffer);
+  if (imageKernelBuffer == (cl_mem) NULL)
     (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
-      ResourceLimitWarning,"clEnv->library->clEnqueueUnmapMemObject failed.",
-      "'%s'",".");
-    return((cl_mem) NULL);
-  }
+      ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
   return(imageKernelBuffer);
 }
 
@@ -443,7 +315,7 @@ static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
   histogramKernel = AcquireOpenCLKernel(device,"Histogram");
   if (histogramKernel == NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
     goto cleanup;
   }
 
@@ -456,7 +328,7 @@ static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
   clStatus|=clEnv->library->clSetKernelArg(histogramKernel,i++,sizeof(cl_mem),(void *)&histogramBuffer);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -468,7 +340,7 @@ static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
 
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
     goto cleanup;
   }
   RecordProfileData(device,histogramKernel,event);
@@ -478,7 +350,7 @@ static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
 cleanup:
  
   if (histogramKernel!=NULL)
-    RelinquishOpenCLKernel(histogramKernel);
+    ReleaseOpenCLKernel(histogramKernel);
 
   return(outputReady);
 }
@@ -498,25 +370,15 @@ cleanup:
 static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv,
   const NoiseType noise_type,ExceptionInfo *exception)
 {
-  CacheView
-    *filteredImage_view,
-    *image_view;
-
-  cl_command_queue
-    queue;
-
   cl_float
     attenuate;
 
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
     addNoiseKernel;
 
-  cl_event
-    event;
-
   cl_mem
     filteredImageBuffer,
     imageBuffer;
@@ -528,11 +390,15 @@ static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv,
     numRandomNumberPerPixel,
     pixelsPerWorkitem,
     seed0,
-    seed1;
+    seed1,
+    workItemCount;
 
   const char
     *option;
 
+  const unsigned long
+    *s;
+
   MagickBooleanType
     outputReady;
 
@@ -542,157 +408,116 @@ static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv,
   Image
     *filteredImage;
 
-  size_t
-    global_work_size[1],
-    local_work_size[1];
-
-  unsigned int
-    k;
-
-  void
-    *filteredPixels;
+  RandomInfo
+    *randomInfo;
 
-  outputReady = MagickFalse;
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  filteredPixels = NULL;
-  filteredImageBuffer = NULL;
-  addNoiseKernel = NULL;
+  size_t
+    gsize[1],
+    i,
+    lsize[1],
+    numRandPerChannel;
 
-  device = RequestOpenCLDevice(clEnv);
-  queue = AcquireOpenCLCommandQueue(device);
+  filteredImage=NULL;
+  addNoiseKernel=NULL;
+  outputReady=MagickFalse;
 
-  image_view=AcquireAuthenticCacheView(image,exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
-  filteredImage=CloneImage(image,image->columns,image->rows,MagickTrue,exception);
+  filteredImage=CloneImage(image,image->columns,image->rows,MagickTrue,
+    exception);
   if (filteredImage == (Image *) NULL)
     goto cleanup;
-  if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
-  if (filteredImageBuffer == (void *) NULL)
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
+  if (filteredImageBuffer == (cl_mem) NULL)
     goto cleanup;
 
   /* find out how many random numbers needed by pixel */
-  numRandomNumberPerPixel = 0;
+  numRandPerChannel=0;
+  numRandomNumberPerPixel=0;
+  switch (noise_type)
   {
-    unsigned int numRandPerChannel = 0;
-    switch (noise_type)
-    {
     case UniformNoise:
     case ImpulseNoise:
     case LaplacianNoise:
     case RandomNoise:
     default:
-      numRandPerChannel = 1;
+      numRandPerChannel=1;
       break;
     case GaussianNoise:
     case MultiplicativeGaussianNoise:
     case PoissonNoise:
-      numRandPerChannel = 2;
+      numRandPerChannel=2;
       break;
-    };
-
-    if (GetPixelRedTraits(image) != UndefinedPixelTrait)
-      numRandomNumberPerPixel+=numRandPerChannel;
-    if (GetPixelGreenTraits(image) != UndefinedPixelTrait)
-      numRandomNumberPerPixel+=numRandPerChannel;
-    if (GetPixelBlueTraits(image) != UndefinedPixelTrait)
-      numRandomNumberPerPixel+=numRandPerChannel;
-    if (GetPixelAlphaTraits(image) != UndefinedPixelTrait)
-      numRandomNumberPerPixel+=numRandPerChannel;
-  }
-
-  addNoiseKernel = AcquireOpenCLKernel(device,"AddNoise");
-  if (addNoiseKernel == NULL)
-  {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  {
-    cl_uint workItemCount;
-    workItemCount = device->max_compute_units * 2 * 256; // 256 work items per group, 2 groups per CU
-    inputPixelCount = (cl_int) (image->columns * image->rows);
-    pixelsPerWorkitem = (inputPixelCount + workItemCount - 1) / workItemCount;
-    pixelsPerWorkitem = ((pixelsPerWorkitem + 3) / 4) * 4;
-
-    local_work_size[0] = 256;
-    global_work_size[0] = workItemCount;
-  }
-  {
-    RandomInfo* randomInfo = AcquireRandomInfo();
-    const unsigned long* s = GetRandomInfoSeed(randomInfo);
-    seed0 = s[0];
-    (void) GetPseudoRandomValue(randomInfo);
-    seed1 = s[0];
-    randomInfo = DestroyRandomInfo(randomInfo);
-  }
-
-  number_channels = (cl_uint) image->number_channels;
-  bufferLength = (cl_uint)(image->columns * image->rows * image->number_channels);
+  };
+  if (GetPixelRedTraits(image) != UndefinedPixelTrait)
+    numRandomNumberPerPixel+=numRandPerChannel;
+  if (GetPixelGreenTraits(image) != UndefinedPixelTrait)
+    numRandomNumberPerPixel+=numRandPerChannel;
+  if (GetPixelBlueTraits(image) != UndefinedPixelTrait)
+    numRandomNumberPerPixel+=numRandPerChannel;
+  if (GetPixelAlphaTraits(image) != UndefinedPixelTrait)
+    numRandomNumberPerPixel+=numRandPerChannel;
+
+  addNoiseKernel=AcquireOpenCLKernel(device,"AddNoise");
+  if (addNoiseKernel == (cl_kernel) NULL)
+  {
+    (void)OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
+
+  /* 256 work items per group, 2 groups per CU */
+  workItemCount=device->max_compute_units*2*256;
+  inputPixelCount=(cl_int) (image->columns*image->rows);
+  pixelsPerWorkitem=(inputPixelCount+workItemCount-1)/workItemCount;
+  pixelsPerWorkitem=((pixelsPerWorkitem+3)/4)*4;
+  lsize[0]=256;
+  gsize[0]=workItemCount;
+
+  randomInfo=AcquireRandomInfo();
+  s=GetRandomInfoSeed(randomInfo);
+  seed0=s[0];
+  (void) GetPseudoRandomValue(randomInfo);
+  seed1=s[0];
+  randomInfo=DestroyRandomInfo(randomInfo);
+
+  number_channels=(cl_uint) image->number_channels;
+  bufferLength=(cl_uint) (image->columns*image->rows*image->number_channels);
   attenuate=1.0f;
   option=GetImageArtifact(image,"attenuate");
   if (option != (char *) NULL)
     attenuate=(float)StringToDouble(option,(char **) NULL);
 
-  k = 0;
-  clStatus=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&imageBuffer);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&number_channels);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(ChannelType),(void *)&image->channel_mask);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&bufferLength);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&pixelsPerWorkitem);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(NoiseType),(void *)&noise_type);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_float),(void *)&attenuate);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed0);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&seed1);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_uint),(void *)&numRandomNumberPerPixel);
-  clStatus|=clEnv->library->clSetKernelArg(addNoiseKernel,k++,sizeof(cl_mem),(void *)&filteredImageBuffer);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  clStatus=clEnv->library->clEnqueueNDRangeKernel(queue,addNoiseKernel,1,NULL,global_work_size,local_work_size,0,NULL,&event);
-  if (clStatus != CL_SUCCESS)
+  i=0;
+  status =SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_uint),(void *)&number_channels);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_uint),(void *)&bufferLength);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_uint),(void *)&pixelsPerWorkitem);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(NoiseType),(void *)&noise_type);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_float),(void *)&attenuate);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_uint),(void *)&seed0);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_uint),(void *)&seed1);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_uint),(void *)&numRandomNumberPerPixel);
+  status|=SetOpenCLKernelArg(addNoiseKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+  if (status != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"clSetKernelArg failed.",".");
     goto cleanup;
   }
 
-  RecordProfileData(device,addNoiseKernel,event);
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
-    goto cleanup;
-
-  outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
-
+  outputReady=EnqueueOpenCLKernel(addNoiseKernel,1,(const size_t *) NULL,gsize,
+    lsize,image,filteredImage,exception);
 cleanup:
 
-  image_view=DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view=DestroyCacheView(filteredImage_view);
-
-  if (queue!=NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (addNoiseKernel!=NULL)
-    RelinquishOpenCLKernel(addNoiseKernel);
-  if (imageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (filteredImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (outputReady == MagickFalse && filteredImage != NULL) 
+  if (addNoiseKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(addNoiseKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
     filteredImage=DestroyImage(filteredImage);
 
   return(filteredImage);
@@ -736,23 +561,13 @@ MagickPrivate Image *AccelerateAddNoiseImage(const Image *image,
 static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv,
   const double radius,const double sigma,ExceptionInfo *exception)
 {
-  CacheView
-    *filteredImage_view,
-    *image_view;
-
-  cl_command_queue
-    queue;
-
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
     blurColumnKernel,
     blurRowKernel;
 
-  cl_event
-    event;
-
   cl_mem
     filteredImageBuffer,
     imageBuffer,
@@ -777,202 +592,126 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv,
   MagickSizeType
     length;
 
-  unsigned int
-    i;
-
-  void
-    *filteredPixels;
-
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  imageBuffer = NULL;
-  tempImageBuffer = NULL;
-  filteredImageBuffer = NULL;
-  filteredPixels = NULL;
-  imageKernelBuffer = NULL;
-  blurRowKernel = NULL;
-  blurColumnKernel = NULL;
-  queue = NULL;
-
-  outputReady = MagickFalse;
+  size_t
+    chunkSize=256,
+    gsize[2],
+    i,
+    lsize[2];
 
-  device = RequestOpenCLDevice(clEnv);
-  queue = AcquireOpenCLCommandQueue(device);
+  filteredImage=NULL;
+  tempImageBuffer=NULL;
+  imageKernelBuffer=NULL;
+  blurRowKernel=NULL;
+  blurColumnKernel=NULL;
+  outputReady=MagickFalse;
 
-  image_view=AcquireAuthenticCacheView(image,exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
   filteredImage=CloneImage(image,image->columns,image->rows,MagickTrue,exception);
   if (filteredImage == (Image *) NULL)
     goto cleanup;
-  if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
+  if (filteredImageBuffer == (cl_mem) NULL)
     goto cleanup;
-  }
 
-  filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
-  if (filteredImageBuffer == (void *) NULL)
+  imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
+    exception);
+  if (imageKernelBuffer == (cl_mem) NULL)
     goto cleanup;
 
-  imageKernelBuffer=createKernelInfo(clEnv,device,queue,radius,sigma,
-    &kernelWidth,exception);
+  length=image->columns*image->rows;
+  tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
+    sizeof(cl_float4),NULL);
+  if (tempImageBuffer == (cl_mem) NULL)
+    goto cleanup;
 
+  blurRowKernel=AcquireOpenCLKernel(device,"BlurRow");
+  if (blurRowKernel == (cl_kernel) NULL)
   {
-    /* create temp buffer */
-    {
-      length = image->columns * image->rows;
-      tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
-        goto cleanup;
-      }
-    }
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
+  blurColumnKernel=AcquireOpenCLKernel(device,"BlurColumn");
+  if (blurColumnKernel == (cl_kernel) NULL)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
 
-    /* get the OpenCL kernels */
-    {
-      blurRowKernel = AcquireOpenCLKernel(device,"BlurRow");
-      if (blurRowKernel == NULL)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-        goto cleanup;
-      };
+  number_channels=(cl_uint) image->number_channels;
+  imageColumns=(cl_uint) image->columns;
+  imageRows=(cl_uint) image->rows;
 
-      blurColumnKernel = AcquireOpenCLKernel(device,"BlurColumn");
-      if (blurColumnKernel == NULL)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-        goto cleanup;
-      };
-    }
+  i=0;
+  status =SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
+  if (status != CL_SUCCESS)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
+    goto cleanup;
+  }
 
-    number_channels = (cl_uint) image->number_channels;
-    imageColumns = (cl_uint) image->columns;
-    imageRows = (cl_uint) image->rows;
+  gsize[0]=chunkSize*((image->columns+chunkSize-1)/chunkSize);
+  gsize[1]=image->rows;
+  lsize[0]=chunkSize;
+  lsize[1]=1;
 
-    {
-      /* need logic to decide this value */
-      int chunkSize = 256;
+  outputReady=EnqueueOpenCLKernel(blurRowKernel,2,NULL,gsize,lsize,image,
+    filteredImage,exception);
+  if (outputReady == MagickFalse)
+    goto cleanup;
 
-      {
-        /* set the kernel arguments */
-        i = 0;
-        clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
-        clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
-        if (clStatus != CL_SUCCESS)
-        {
-          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-          goto cleanup;
-        }
-      }
+  i=0;
+  status =SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_uint),&number_channels);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
+  status|=SetOpenCLKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+  if (status != CL_SUCCESS)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
+    goto cleanup;
+  }
 
-      /* launch the kernel */
-      {
-        size_t gsize[2];
-        size_t wsize[2];
+  gsize[0]=image->columns;
+  gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize);
+  lsize[0]=1;
+  lsize[1]=chunkSize;
 
-        gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
-        gsize[1] = image->rows;
-        wsize[0] = chunkSize;
-        wsize[1] = 1;
+  outputReady=EnqueueOpenCLKernel(blurRowKernel,2,NULL,gsize,lsize,image,
+    filteredImage,exception);
 
-        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
-        if (clStatus != CL_SUCCESS)
-        {
-          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-          goto cleanup;
-        }
-        RecordProfileData(device,blurRowKernel,event);
-      }
-    }
+cleanup:
 
-    {
-      /* need logic to decide this value */
-      int chunkSize = 256;
-
-      {
-        /* set the kernel arguments */
-        i = 0;
-        clStatus=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),&number_channels);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
-        clStatus|=clEnv->library->clSetKernelArg(blurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
-        if (clStatus != CL_SUCCESS)
-        {
-          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-          goto cleanup;
-        }
-      }
-
-      /* launch the kernel */
-      {
-        size_t gsize[2];
-        size_t wsize[2];
-
-        gsize[0] = image->columns;
-        gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
-        wsize[0] = 1;
-        wsize[1] = chunkSize;
-
-        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
-        if (clStatus != CL_SUCCESS)
-        {
-          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-          goto cleanup;
-        }
-        RecordProfileData(device,blurColumnKernel,event);
-      }
-    }
-
-  }
-
-  /* get result */
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
-    goto cleanup;
-
-  outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
-
-cleanup:
-
-  image_view=DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view=DestroyCacheView(filteredImage_view);
-
-  if (imageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (tempImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(tempImageBuffer);
-  if (filteredImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (imageKernelBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageKernelBuffer);
-  if (blurRowKernel!=NULL)
-    RelinquishOpenCLKernel(blurRowKernel);
-  if (blurColumnKernel!=NULL)
-    RelinquishOpenCLKernel(blurColumnKernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (outputReady == MagickFalse && filteredImage != NULL)
-    filteredImage=DestroyImage(filteredImage);
+  if (tempImageBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(tempImageBuffer);
+  if (imageKernelBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(imageKernelBuffer);
+  if (blurRowKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(blurRowKernel);
+  if (blurColumnKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(blurColumnKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
+    filteredImage=DestroyImage(filteredImage);
 
   return(filteredImage);
 }
@@ -1069,7 +808,7 @@ static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv,
 
   RecordProfileData(device,compositeKernel,event);
 
-  RelinquishOpenCLKernel(compositeKernel);
+  ReleaseOpenCLKernel(compositeKernel);
 
   return((clStatus==CL_SUCCESS) ? MagickTrue : MagickFalse);
 }
@@ -1220,7 +959,7 @@ cleanup:
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
 
   return(outputReady);
 }
@@ -1361,7 +1100,7 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv,
   filterKernel = AcquireOpenCLKernel(device,"Contrast");
   if (filterKernel == NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
     goto cleanup;
   }
 
@@ -1372,7 +1111,7 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv,
   clStatus|=clEnv->library->clSetKernelArg(filterKernel,i++,sizeof(cl_uint),&uSharpen);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -1383,7 +1122,7 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv,
   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.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
     goto cleanup;
   }
   RecordProfileData(device,filterKernel,event);
@@ -1400,7 +1139,7 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
   outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
@@ -1412,11 +1151,11 @@ cleanup:
   if (imageBuffer!=NULL)
     clEnv->library->clReleaseMemObject(imageBuffer);
   if (filterKernel!=NULL)
-    RelinquishOpenCLKernel(filterKernel);
+    ReleaseOpenCLKernel(filterKernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
 
   return(outputReady);
 }
@@ -1637,7 +1376,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -1647,7 +1386,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
     clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", ".");
       goto cleanup;
     }
   }
@@ -1945,7 +1684,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
   stretchKernel = AcquireOpenCLKernel(device,"ContrastStretch");
   if (stretchKernel == NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
     goto cleanup;
   }
 
@@ -1958,7 +1697,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
   clStatus|=clEnv->library->clSetKernelArg(stretchKernel,i++,sizeof(FloatPixelPacket),&black);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -1970,7 +1709,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
 
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
     goto cleanup;
   }
   RecordProfileData(device,stretchKernel,event);
@@ -1988,7 +1727,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -2010,13 +1749,13 @@ cleanup:
   if (histogram!=NULL)
     histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
   if (histogramKernel!=NULL)
-    RelinquishOpenCLKernel(histogramKernel);
+    ReleaseOpenCLKernel(histogramKernel);
   if (stretchKernel!=NULL)
-    RelinquishOpenCLKernel(stretchKernel);
+    ReleaseOpenCLKernel(stretchKernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
 
   return(outputReady);
 }
@@ -2171,7 +1910,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
   assert(filteredImage != NULL);
   if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", ".");
     goto cleanup;
   }
   filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
@@ -2225,7 +1964,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
   clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, convolutionKernel, kernelBufferPtr, 0, NULL, NULL);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", ".");
     goto cleanup;
   }
 
@@ -2249,7 +1988,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
     clkernel = AcquireOpenCLKernel(device,"ConvolveOptimized");
     if (clkernel == NULL)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
       goto cleanup;
     }
 
@@ -2273,7 +2012,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
     clStatus|=clEnv->library->clSetKernelArg(clkernel,i++, kernel->width*kernel->height*sizeof(float),NULL);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
       goto cleanup;
     }
 
@@ -2285,7 +2024,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
     clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, global_work_size, localGroupSize, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,clkernel,event);
@@ -2296,7 +2035,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
     clkernel = AcquireOpenCLKernel(device,"Convolve");
     if (clkernel == NULL)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
       goto cleanup;
     }
 
@@ -2318,7 +2057,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
     clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
       goto cleanup;
     }
 
@@ -2330,7 +2069,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
     
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
   }
@@ -2348,7 +2087,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -2366,11 +2105,11 @@ cleanup:
   if (convolutionKernel != NULL)
     clEnv->library->clReleaseMemObject(convolutionKernel);
   if (clkernel != NULL)
-    RelinquishOpenCLKernel(clkernel);
+    ReleaseOpenCLKernel(clkernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
   if (outputReady == MagickFalse)
   {
     if (filteredImage != NULL)
@@ -2537,7 +2276,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
   assert(filteredImage != NULL);
   if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", ".");
     goto cleanup;
   }
   filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
@@ -2580,7 +2319,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
   clStatus |=clEnv->library->clSetKernelArg(hullPass1,6,sizeof(int),(void *)&matte);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -2594,7 +2333,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
   clStatus |=clEnv->library->clSetKernelArg(hullPass2,6,sizeof(int),(void *)&matte);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -2618,14 +2357,14 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
     clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
       goto cleanup;
     }
     /* launch the kernel */
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass1,event);
@@ -2634,7 +2373,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass2,event);
@@ -2650,14 +2389,14 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
     clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
       goto cleanup;
     }
     /* launch the kernel */
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass1,event);
@@ -2666,7 +2405,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass2,event);
@@ -2680,14 +2419,14 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
     clStatus|=clEnv->library->clSetKernelArg(hullPass2,5,sizeof(int),(void *)&polarity);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
       goto cleanup;
     }
     /* launch the kernel */
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass1,event);
@@ -2696,7 +2435,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass2,event);
@@ -2714,14 +2453,14 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
 
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
       goto cleanup;
     }
     /* launch the kernel */
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass1, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass1,event);
@@ -2730,7 +2469,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,hullPass2,event);
@@ -2748,7 +2487,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -2763,7 +2502,7 @@ cleanup:
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
   if (imageBuffer!=NULL)
     clEnv->library->clReleaseMemObject(imageBuffer);
   for (k = 0; k < 2; k++)
@@ -2774,9 +2513,9 @@ cleanup:
   if (filteredImageBuffer!=NULL)
     clEnv->library->clReleaseMemObject(filteredImageBuffer);
   if (hullPass1!=NULL)
-    RelinquishOpenCLKernel(hullPass1);
+    ReleaseOpenCLKernel(hullPass1);
   if (hullPass2!=NULL)
-    RelinquishOpenCLKernel(hullPass2);
+    ReleaseOpenCLKernel(hullPass2);
   if (outputReady == MagickFalse && filteredImage != NULL)
     filteredImage=DestroyImage(filteredImage);
 
@@ -2982,7 +2721,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -2992,7 +2731,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
     clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, histogramBuffer, histogram, 0, NULL, NULL);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", ".");
       goto cleanup;
     }
   }
@@ -3167,7 +2906,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
   equalizeKernel = AcquireOpenCLKernel(device,"Equalize");
   if (equalizeKernel == NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
     goto cleanup;
   }
 
@@ -3180,7 +2919,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
   clStatus|=clEnv->library->clSetKernelArg(equalizeKernel,i++,sizeof(FloatPixelPacket),&black);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -3192,7 +2931,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
 
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
     goto cleanup;
   }
   RecordProfileData(device,equalizeKernel,event);
@@ -3210,7 +2949,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -3233,13 +2972,13 @@ cleanup:
   if (histogram!=NULL)
     histogram=(cl_uint4 *) RelinquishMagickMemory(histogram);
   if (histogramKernel!=NULL)
-    RelinquishOpenCLKernel(histogramKernel);
+    ReleaseOpenCLKernel(histogramKernel);
   if (equalizeKernel!=NULL)
-    RelinquishOpenCLKernel(equalizeKernel);
+    ReleaseOpenCLKernel(equalizeKernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device, queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
 
   return(outputReady);
 }
@@ -3284,144 +3023,99 @@ static MagickBooleanType ComputeFunctionImage(Image *image,MagickCLEnv clEnv,
   const MagickFunction function,const size_t number_parameters,
   const double *parameters,ExceptionInfo *exception)
 {
-  CacheView
-    *image_view;
-
-  cl_command_queue
-    queue;
-
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
-    clkernel;
-
-  cl_event
-    event;
+    functionKernel;
 
   cl_mem
     imageBuffer,
     parametersBuffer;
 
   cl_uint
+    number_params,
     number_channels;
 
   float
     *parametersBufferPtr;
 
   MagickBooleanType
-    status;
+    outputReady;
 
   MagickCLDevice
     device;
 
   size_t
-    globalWorkSize[2];
-
-  unsigned int
+    gsize[2],
     i;
 
-  void
-    *pixels;
-
-  status = MagickFalse;
-
-  clkernel = NULL;
-  queue = NULL;
-  imageBuffer = NULL;
-  parametersBuffer = NULL;
-  pixels = NULL;
+  outputReady=MagickFalse;
 
-  device = RequestOpenCLDevice(clEnv);
+  functionKernel=NULL;
+  parametersBuffer=NULL;
 
-  image_view=AcquireAuthenticCacheView(image,exception);
-  imageBuffer=createReadWriteBuffer(image,image_view,clEnv,device,pixels,
-    exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
 
-  parametersBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, number_parameters * sizeof(float), NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
-    goto cleanup;
-  }
-
-  queue = AcquireOpenCLCommandQueue(device);
-
-  parametersBufferPtr = (float*)clEnv->library->clEnqueueMapBuffer(queue, parametersBuffer, CL_TRUE, CL_MAP_WRITE, 0, number_parameters * sizeof(float)
-                , 0, NULL, NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
+  parametersBufferPtr=(float *) AcquireQuantumMemory(number_parameters,
+    sizeof(float));
+  if (parametersBufferPtr == (float *) NULL)
     goto cleanup;
-  }
-  for (i = 0; i < number_parameters; i++)
-  {
-    parametersBufferPtr[i] = (float)parameters[i];
-  }
-  clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, parametersBuffer, parametersBufferPtr, 0, NULL, NULL);
-  if (clStatus != CL_SUCCESS)
+  for (i=0; i<number_parameters; i++)
+    parametersBufferPtr[i]=(float) parameters[i];
+  parametersBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
+    CL_MEM_COPY_HOST_PTR,number_parameters*sizeof(float),
+    parametersBufferPtr);
+  parametersBufferPtr=RelinquishMagickMemory(parametersBufferPtr);
+  if (parametersBuffer == (cl_mem) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
     goto cleanup;
   }
 
-  clkernel = AcquireOpenCLKernel(device,"ComputeFunction");
-  if (clkernel == NULL)
+  functionKernel=AcquireOpenCLKernel(device,"ComputeFunction");
+  if (functionKernel == (cl_kernel) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
     goto cleanup;
   }
 
-  number_channels = (cl_uint) image->number_channels;
-
-  /* set the kernel arguments */
-  i = 0;
-  clStatus =clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
-  clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_uint),(void *)&number_channels);
-  clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
-  clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(MagickFunction),(void *)&function);
-  clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(unsigned int),(void *)&number_parameters);
-  clStatus|=clEnv->library->clSetKernelArg(clkernel,i++,sizeof(cl_mem),(void *)&parametersBuffer);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-    goto cleanup;
-  }
+  number_channels=(cl_uint) image->number_channels;
+  number_params=(cl_uint) number_parameters;
 
-  globalWorkSize[0] = image->columns;
-  globalWorkSize[1] = image->rows;
-  /* launch the kernel */
-  clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, clkernel, 2, NULL, globalWorkSize, NULL, 0, NULL, &event);
-  if (clStatus != CL_SUCCESS)
+  i=0;
+  status =SetOpenCLKernelArg(functionKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(functionKernel,i++,sizeof(cl_uint),(void *)&number_channels);
+  status|=SetOpenCLKernelArg(functionKernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
+  status|=SetOpenCLKernelArg(functionKernel,i++,sizeof(MagickFunction),(void *)&function);
+  status|=SetOpenCLKernelArg(functionKernel,i++,sizeof(cl_uint),(void *)&number_params);
+  status|=SetOpenCLKernelArg(functionKernel,i++,sizeof(cl_mem),(void *)&parametersBuffer);
+  if (status != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
     goto cleanup;
   }
-  RecordProfileData(device,clkernel,event);
-
-  if (copyWriteBuffer(image,clEnv,device,queue,imageBuffer,pixels,exception) == MagickFalse)
-      goto cleanup;
 
-  status=SyncCacheViewAuthenticPixels(image_view,exception);
+  gsize[0]=image->columns;
+  gsize[1]=image->rows;
+  outputReady=EnqueueOpenCLKernel(functionKernel,2,(const size_t *) NULL,
+    gsize,(const size_t *) NULL,image,(const Image *) NULL,exception);
 
 cleanup:
 
-  image_view=DestroyCacheView(image_view);
-  
-  if (clkernel != NULL)
-    RelinquishOpenCLKernel(clkernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (imageBuffer != NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (parametersBuffer != NULL)
-    clEnv->library->clReleaseMemObject(parametersBuffer);
-
-  return(status);
+  if (parametersBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(parametersBuffer);
+  if (functionKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(functionKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  return(outputReady);
 }
 
 MagickPrivate MagickBooleanType AccelerateFunctionImage(Image *image,
@@ -3464,21 +3158,12 @@ MagickPrivate MagickBooleanType AccelerateFunctionImage(Image *image,
 static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv,
   const PixelIntensityMethod method,ExceptionInfo *exception)
 {
-  CacheView
-    *image_view;
-
-  cl_command_queue
-    queue;
-
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
     grayscaleKernel;
 
-  cl_event
-    event;
-
   cl_mem
     imageBuffer;
 
@@ -3493,92 +3178,57 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv,
   MagickCLDevice
     device;
 
-  register ssize_t
+  size_t
+    gsize[2],
     i;
 
-  void
-    *inputPixels;
-
-  outputReady = MagickFalse;
-  inputPixels = NULL;
-  grayscaleKernel = NULL;
+  outputReady=MagickFalse;
+  grayscaleKernel=NULL;
 
   assert(image != (Image *) NULL);
   assert(image->signature == MagickCoreSignature);
-  if (image->debug != MagickFalse)
-    (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
-
-  /*
-   * initialize opencl env
-   */
-  device = RequestOpenCLDevice(clEnv);
-  queue = AcquireOpenCLCommandQueue(device);
-
-  /* Create and initialize OpenCL buffers.
-   inputPixels = AcquirePixelCachePixels(image, &length, exception);
-   assume this  will get a writable image
-   */
-  image_view=AcquireAuthenticCacheView(image,exception);
-  imageBuffer=createReadWriteBuffer(image,image_view,clEnv,device,inputPixels,
-    exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
 
-  grayscaleKernel = AcquireOpenCLKernel(device,"Grayscale");
-  if (grayscaleKernel == NULL)
+  grayscaleKernel=AcquireOpenCLKernel(device,"Grayscale");
+  if (grayscaleKernel == (cl_kernel) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
     goto cleanup;
   }
 
-  number_channels = (cl_uint) image->number_channels;
-  intensityMethod = (cl_uint) method;
-  colorspace = (cl_uint) image->colorspace;
+  number_channels=(cl_uint) image->number_channels;
+  intensityMethod=(cl_uint) method;
+  colorspace=(cl_uint) image->colorspace;
 
-  i = 0;
-  clStatus=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
-  clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&number_channels);
-  clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&colorspace);
-  clStatus|=clEnv->library->clSetKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&intensityMethod);
-  if (clStatus != CL_SUCCESS)
+  i=0;
+  status =SetOpenCLKernelArg(grayscaleKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&number_channels);
+  status|=SetOpenCLKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&colorspace);
+  status|=SetOpenCLKernelArg(grayscaleKernel,i++,sizeof(cl_uint),&intensityMethod);
+  if (status != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
     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, grayscaleKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
-    if (clStatus != CL_SUCCESS)
-    {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-      goto cleanup;
-    }
-    RecordProfileData(device,grayscaleKernel,event);
-  }
-
-  if (copyWriteBuffer(image,clEnv,device,queue,imageBuffer,inputPixels,exception) == MagickFalse)
-    goto cleanup;
-
-  outputReady=SyncCacheViewAuthenticPixels(image_view,exception);
+  gsize[0]=image->columns;
+  gsize[1]=image->rows;
+  outputReady=EnqueueOpenCLKernel(grayscaleKernel,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 (grayscaleKernel!=NULL)
-    RelinquishOpenCLKernel(grayscaleKernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+  if (grayscaleKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(grayscaleKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
 
-  return( outputReady);
+  return(outputReady);
 }
 
 MagickPrivate MagickBooleanType AccelerateGrayscaleImage(Image* image,
@@ -3731,7 +3381,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
     assert(filteredImage != NULL);
     if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", ".");
       goto cleanup;
     }
     filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
@@ -3780,14 +3430,14 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
       blurRowKernel = AcquireOpenCLKernel(device,"LocalContrastBlurRow");
       if (blurRowKernel == NULL)
       {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
         goto cleanup;
       };
 
       blurColumnKernel = AcquireOpenCLKernel(device,"LocalContrastBlurApplyColumn");
       if (blurColumnKernel == NULL)
       {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
         goto cleanup;
       };
     }
@@ -3811,7 +3461,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
       
       if (clStatus != CL_SUCCESS)
       {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
         goto cleanup;
       }
     }
@@ -3834,7 +3484,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
         clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, goffset, gsize, wsize, 0, NULL, &event);
         if (clStatus != CL_SUCCESS)
         {
-          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
           goto cleanup;
         }
         RecordProfileData(device,blurRowKernel,event);
@@ -3854,7 +3504,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
 
       if (clStatus != CL_SUCCESS)
       {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
         goto cleanup;
       }
     }
@@ -3877,7 +3527,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
         clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurColumnKernel, 2, goffset, gsize, wsize, 0, NULL, &event);
         if (clStatus != CL_SUCCESS)
         {
-          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+          (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
           goto cleanup;
         }
         RecordProfileData(device,blurColumnKernel,event);
@@ -3898,7 +3548,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -3919,13 +3569,13 @@ cleanup:
   if (imageKernelBuffer!=NULL)
     clEnv->library->clReleaseMemObject(imageKernelBuffer);
   if (blurRowKernel!=NULL)
-    RelinquishOpenCLKernel(blurRowKernel);
+    ReleaseOpenCLKernel(blurRowKernel);
   if (blurColumnKernel!=NULL)
-    RelinquishOpenCLKernel(blurColumnKernel);
+    ReleaseOpenCLKernel(blurColumnKernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device, queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
   if (outputReady == MagickFalse)
   {
     if (filteredImage != NULL)
@@ -4074,7 +3724,7 @@ static MagickBooleanType ComputeModulateImage(Image *image,MagickCLEnv clEnv,
   modulateKernel = AcquireOpenCLKernel(device, "Modulate");
   if (modulateKernel == NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", ".");
     goto cleanup;
   }
 
@@ -4091,7 +3741,7 @@ static MagickBooleanType ComputeModulateImage(Image *image,MagickCLEnv clEnv,
   clStatus|=clEnv->library->clSetKernelArg(modulateKernel,i++,sizeof(cl_float),&color);
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -4103,7 +3753,7 @@ static MagickBooleanType ComputeModulateImage(Image *image,MagickCLEnv clEnv,
        clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, modulateKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
     if (clStatus != CL_SUCCESS)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
       goto cleanup;
     }
     RecordProfileData(device,modulateKernel,event);
@@ -4121,7 +3771,7 @@ static MagickBooleanType ComputeModulateImage(Image *image,MagickCLEnv clEnv,
   }
   if (clStatus != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
 
@@ -4134,11 +3784,11 @@ cleanup:
   if (imageBuffer!=NULL)
     clEnv->library->clReleaseMemObject(imageBuffer);
   if (modulateKernel!=NULL)
-    RelinquishOpenCLKernel(modulateKernel);
+    ReleaseOpenCLKernel(modulateKernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
 
   return outputReady;
 
@@ -4306,7 +3956,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
   if (SetImageStorageClass(filteredImage,DirectClass,exception) != MagickTrue)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), 
-      ResourceLimitError, "CloneImage failed.", "'%s'", ".");
+      ResourceLimitError, "CloneImage failed.", ".");
     goto cleanup;
   }
   filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
@@ -4368,7 +4018,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
  if (clStatus != CL_SUCCESS)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, 
-      "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+      "clEnv->library->clEnqueueUnmapMemObject failed.", ".");
     goto cleanup;
   }
 
@@ -4400,7 +4050,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
  if (clStatus != CL_SUCCESS)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
-      "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+      "clEnv->library->clEnqueueUnmapMemObject failed.", ".");
     goto cleanup;
   }
 
@@ -4410,7 +4060,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
   if (motionBlurKernel == NULL)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
-      "AcquireOpenCLKernel failed.", "'%s'", ".");
+      "AcquireOpenCLKernel failed.", ".");
     goto cleanup;
   }
   
@@ -4446,7 +4096,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
   if (clStatus != CL_SUCCESS)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
-      "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+      "clEnv->library->clSetKernelArg failed.", ".");
     goto cleanup;
   }
 
@@ -4463,7 +4113,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
   if (clStatus != CL_SUCCESS)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
-      "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+      "clEnv->library->clEnqueueNDRangeKernel failed.", ".");
     goto cleanup;
   }
   RecordProfileData(device,motionBlurKernel,event);
@@ -4484,7 +4134,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
   if (clStatus != CL_SUCCESS)
   {
     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError,
-      "Reading output image from CL buffer failed.", "'%s'", ".");
+      "Reading output image from CL buffer failed.", ".");
     goto cleanup;
   }
   outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
@@ -4502,11 +4152,11 @@ cleanup:
   if (imageKernelBuffer!=NULL)
     clEnv->library->clReleaseMemObject(imageKernelBuffer);
   if (motionBlurKernel!=NULL)
-    RelinquishOpenCLKernel(motionBlurKernel);
+    ReleaseOpenCLKernel(motionBlurKernel);
   if (queue != NULL)
     RelinquishOpenCLCommandQueue(device,queue);
   if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
+    ReleaseOpenCLDevice(device);
   if (outputReady == MagickFalse && filteredImage != NULL)
     filteredImage=DestroyImage(filteredImage);
 
@@ -4552,20 +4202,18 @@ MagickPrivate Image *AccelerateMotionBlurImage(const Image *image,
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 */
 
-static MagickBooleanType resizeHorizontalFilter(MagickCLEnv clEnv,
-  MagickCLDevice device,cl_command_queue queue,cl_mem image,
-  cl_uint number_channels,cl_uint columns,cl_uint rows,cl_mem resizedImage,
-  cl_uint resizedColumns,cl_uint resizedRows,const ResizeFilter *resizeFilter,
-  cl_mem resizeFilterCubicCoefficients,const float xFactor,
-  ExceptionInfo *exception)
+static MagickBooleanType resizeHorizontalFilter(MagickCLDevice device,
+  const Image *image,Image *filteredImage,cl_mem imageBuffer,
+  cl_uint number_channels,cl_uint columns,cl_uint rows,
+  cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows,
+  const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
+  const float xFactor,ExceptionInfo *exception)
 {
   cl_kernel
     horizontalKernel;
 
-  cl_event
-    event;
-
-  cl_int clStatus;
+  cl_int
+    status;
 
   const unsigned int
     workgroupSize = 256;
@@ -4586,24 +4234,24 @@ static MagickBooleanType resizeHorizontalFilter(MagickCLEnv clEnv,
     resizeWindowType;
 
   MagickBooleanType
-    status;
+    outputReady;
 
   size_t
     gammaAccumulatorLocalMemorySize,
-    global_work_size[2],
+    gsize[2],
+    i,
     imageCacheLocalMemorySize,
     pixelAccumulatorLocalMemorySize,
-    local_work_size[2],
+    lsize[2],
     totalLocalMemorySize,
     weightAccumulatorLocalMemorySize;
 
   unsigned int
     chunkSize,
-    i,
     pixelPerWorkgroup;
 
-  horizontalKernel = NULL;
-  status = MagickFalse;
+  horizontalKernel=NULL;
+  outputReady=MagickFalse;
 
   /*
   Apply filter to resize vertically from image to resize image.
@@ -4623,13 +4271,13 @@ static MagickBooleanType resizeHorizontalFilter(MagickCLEnv clEnv,
 
   if (resizedColumns < workgroupSize) 
   {
-    chunkSize = 32;
-    pixelPerWorkgroup = 32;
+    chunkSize=32;
+    pixelPerWorkgroup=32;
   }
   else
   {
-    chunkSize = workgroupSize;
-    pixelPerWorkgroup = workgroupSize;
+    chunkSize=workgroupSize;
+    pixelPerWorkgroup=workgroupSize;
   }
 
 DisableMSCWarning(4127)
@@ -4637,35 +4285,36 @@ DisableMSCWarning(4127)
 RestoreMSCWarning
   {
     /* calculate the local memory size needed per workgroup */
-    cacheRangeStart = (int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
-    cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+MagickEpsilon)+support+0.5);
-    numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
-    imageCacheLocalMemorySize = numCachedPixels * sizeof(CLQuantum) * number_channels;
-    totalLocalMemorySize = imageCacheLocalMemorySize;
+    cacheRangeStart=(int) (((0 + 0.5)/xFactor+MagickEpsilon)-support+0.5);
+    cacheRangeEnd=(int) ((((pixelPerWorkgroup-1) + 0.5)/xFactor+
+      MagickEpsilon)+support+0.5);
+    numCachedPixels=cacheRangeEnd-cacheRangeStart+1;
+    imageCacheLocalMemorySize=numCachedPixels*sizeof(CLQuantum)*
+      number_channels;
+    totalLocalMemorySize=imageCacheLocalMemorySize;
 
     /* local size for the pixel accumulator */
-    pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
+    pixelAccumulatorLocalMemorySize=chunkSize*sizeof(cl_float4);
     totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
 
     /* local memory size for the weight accumulator */
-    weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
+    weightAccumulatorLocalMemorySize=chunkSize*sizeof(float);
     totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
 
     /* local memory size for the gamma accumulator */
     if ((number_channels == 4) || (number_channels == 2))
-      gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
+      gammaAccumulatorLocalMemorySize=chunkSize*sizeof(float);
     else
-      gammaAccumulatorLocalMemorySize = sizeof(float);
+      gammaAccumulatorLocalMemorySize=sizeof(float);
     totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
 
     if (totalLocalMemorySize <= device->local_memory_size)
       break;
     else
     {
-      pixelPerWorkgroup = pixelPerWorkgroup/2;
-      chunkSize = chunkSize/2;
-      if (pixelPerWorkgroup == 0
-          || chunkSize == 0)
+      pixelPerWorkgroup=pixelPerWorkgroup/2;
+      chunkSize=chunkSize/2;
+      if ((pixelPerWorkgroup == 0) || (chunkSize == 0))
       {
         /* quit, fallback to CPU */
         goto cleanup;
@@ -4673,95 +4322,80 @@ RestoreMSCWarning
     }
   }
 
-  resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
-  resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
-
-  horizontalKernel = AcquireOpenCLKernel(device, "ResizeHorizontalFilter");
-  if (horizontalKernel == NULL)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  i = 0;
-  clStatus = clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&image);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_uint), (void*)&number_channels);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_uint), (void*)&columns);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_uint), (void*)&rows);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_uint), (void*)&resizedColumns);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_uint), (void*)&resizedRows);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&xFactor);
-
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeFilterType);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), (void*)&resizeWindowType);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
-
-  resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
-
-  resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
-
-  resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
-
-  resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
-
+  resizeFilterType=(int)GetResizeFilterWeightingType(resizeFilter);
+  resizeWindowType=(int)GetResizeFilterWindowWeightingType(resizeFilter);
 
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, imageCacheLocalMemorySize, NULL);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(int), &numCachedPixels);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, sizeof(unsigned int), &chunkSize);
-  
-
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
-  clStatus |= clEnv->library->clSetKernelArg(horizontalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
-
-  if (clStatus != CL_SUCCESS)
+  horizontalKernel=AcquireOpenCLKernel(device,"ResizeHorizontalFilter");
+  if (horizontalKernel == (cl_kernel) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  global_work_size[0] = (resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
-  global_work_size[1] = resizedRows;
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.", ".");
+    goto cleanup;
+  }
+
+  resizeFilterScale=(float) GetResizeFilterScale(resizeFilter);
+  resizeFilterSupport=(float) GetResizeFilterSupport(resizeFilter);
+  resizeFilterWindowSupport=(float) GetResizeFilterWindowSupport(resizeFilter);
+  resizeFilterBlur=(float) GetResizeFilterBlur(resizeFilter);
+
+  i=0;
+  status =SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_mem),(void*)&imageBuffer);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_uint),(void*)&number_channels);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_uint),(void*)&columns);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_uint),(void*)&rows);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_mem),(void*)&resizedImageBuffer);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_uint),(void*)&resizedColumns);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_uint),(void*)&resizedRows);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(float),(void*)&xFactor);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(int),(void*)&resizeFilterType);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(int),(void*)&resizeWindowType);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(cl_mem),(void*)&resizeFilterCubicCoefficients);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(float),(void*)&resizeFilterSupport);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(float),(void*)&resizeFilterScale);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(float),(void*)&resizeFilterWindowSupport);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(float),(void*)&resizeFilterBlur);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,imageCacheLocalMemorySize,NULL);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(int),&numCachedPixels);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(unsigned int),&pixelPerWorkgroup);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,sizeof(unsigned int),&chunkSize);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,pixelAccumulatorLocalMemorySize,NULL);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,weightAccumulatorLocalMemorySize,NULL);
+  status|=SetOpenCLKernelArg(horizontalKernel,i++,gammaAccumulatorLocalMemorySize,NULL);
 
-  local_work_size[0] = workgroupSize;
-  local_work_size[1] = 1;
-  clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, horizontalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event);
-  if (clStatus != CL_SUCCESS)
+  if (status != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
     goto cleanup;
   }
-  RecordProfileData(device,horizontalKernel,event);
-  status = MagickTrue;
-
 
+  gsize[0]=(resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*
+    workgroupSize;
+  gsize[1]=resizedRows;
+  lsize[0]=workgroupSize;
+  lsize[1]=1;
+  outputReady=EnqueueOpenCLKernel(horizontalKernel,2,(const size_t *) NULL,
+    gsize,lsize,image,filteredImage,exception);
 cleanup:
 
-  if (horizontalKernel != NULL) RelinquishOpenCLKernel(horizontalKernel);
+  if (horizontalKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(horizontalKernel);
 
-  return(status);
+  return(outputReady);
 }
 
-static MagickBooleanType resizeVerticalFilter(MagickCLEnv clEnv,
-  MagickCLDevice device,cl_command_queue queue,cl_mem image,
-  cl_uint number_channels,cl_uint columns,cl_uint rows,cl_mem resizedImage,
-  cl_uint resizedColumns,cl_uint resizedRows,const ResizeFilter *resizeFilter,
-  cl_mem resizeFilterCubicCoefficients,const float yFactor,
-  ExceptionInfo *exception)
+static MagickBooleanType resizeVerticalFilter(MagickCLDevice device,
+  const Image *image,Image * filteredImage,cl_mem imageBuffer,
+  cl_uint number_channels,cl_uint columns,cl_uint rows,
+  cl_mem resizedImageBuffer,cl_uint resizedColumns,cl_uint resizedRows,
+  const ResizeFilter *resizeFilter,cl_mem resizeFilterCubicCoefficients,
+  const float yFactor,ExceptionInfo *exception)
 {
   cl_kernel
     verticalKernel;
 
-  cl_event
-    event;
-
-  cl_int clStatus;
+  cl_int
+    status;
 
   const unsigned int
     workgroupSize = 256;
@@ -4782,24 +4416,24 @@ static MagickBooleanType resizeVerticalFilter(MagickCLEnv clEnv,
     resizeWindowType;
 
   MagickBooleanType
-    status;
+    outputReady;
 
   size_t
     gammaAccumulatorLocalMemorySize,
-    global_work_size[2],
+    gsize[2],
+    i,
     imageCacheLocalMemorySize,
     pixelAccumulatorLocalMemorySize,
-    local_work_size[2],
+    lsize[2],
     totalLocalMemorySize,
     weightAccumulatorLocalMemorySize;
 
   unsigned int
     chunkSize,
-    i,
     pixelPerWorkgroup;
 
-  verticalKernel = NULL;
-  status = MagickFalse;
+  verticalKernel=NULL;
+  outputReady=MagickFalse;
 
   /*
   Apply filter to resize vertically from image to resize image.
@@ -4819,13 +4453,13 @@ static MagickBooleanType resizeVerticalFilter(MagickCLEnv clEnv,
 
   if (resizedRows < workgroupSize) 
   {
-    chunkSize = 32;
-    pixelPerWorkgroup = 32;
+    chunkSize=32;
+    pixelPerWorkgroup=32;
   }
   else
   {
-    chunkSize = workgroupSize;
-    pixelPerWorkgroup = workgroupSize;
+    chunkSize=workgroupSize;
+    pixelPerWorkgroup=workgroupSize;
   }
 
 DisableMSCWarning(4127)
@@ -4833,35 +4467,36 @@ DisableMSCWarning(4127)
 RestoreMSCWarning
   {
     /* calculate the local memory size needed per workgroup */
-    cacheRangeStart = (int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
-    cacheRangeEnd = (int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+MagickEpsilon)+support+0.5);
-    numCachedPixels = cacheRangeEnd - cacheRangeStart + 1;
-    imageCacheLocalMemorySize = numCachedPixels * sizeof(CLQuantum) * number_channels;
-    totalLocalMemorySize = imageCacheLocalMemorySize;
+    cacheRangeStart=(int) (((0 + 0.5)/yFactor+MagickEpsilon)-support+0.5);
+    cacheRangeEnd=(int) ((((pixelPerWorkgroup-1) + 0.5)/yFactor+
+      MagickEpsilon)+support+0.5);
+    numCachedPixels=cacheRangeEnd-cacheRangeStart+1;
+    imageCacheLocalMemorySize=numCachedPixels*sizeof(CLQuantum)*
+      number_channels;
+    totalLocalMemorySize=imageCacheLocalMemorySize;
 
     /* local size for the pixel accumulator */
-    pixelAccumulatorLocalMemorySize = chunkSize * sizeof(cl_float4);
+    pixelAccumulatorLocalMemorySize=chunkSize*sizeof(cl_float4);
     totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;
 
     /* local memory size for the weight accumulator */
-    weightAccumulatorLocalMemorySize = chunkSize * sizeof(float);
+    weightAccumulatorLocalMemorySize=chunkSize*sizeof(float);
     totalLocalMemorySize+=weightAccumulatorLocalMemorySize;
 
     /* local memory size for the gamma accumulator */
     if ((number_channels == 4) || (number_channels == 2))
-      gammaAccumulatorLocalMemorySize = chunkSize * sizeof(float);
+      gammaAccumulatorLocalMemorySize=chunkSize*sizeof(float);
     else
-      gammaAccumulatorLocalMemorySize = sizeof(float);
+      gammaAccumulatorLocalMemorySize=sizeof(float);
     totalLocalMemorySize+=gammaAccumulatorLocalMemorySize;
 
     if (totalLocalMemorySize <= device->local_memory_size)
       break;
     else
     {
-      pixelPerWorkgroup = pixelPerWorkgroup/2;
-      chunkSize = chunkSize/2;
-      if (pixelPerWorkgroup == 0
-          || chunkSize == 0)
+      pixelPerWorkgroup=pixelPerWorkgroup/2;
+      chunkSize=chunkSize/2;
+      if ((pixelPerWorkgroup == 0) || (chunkSize == 0))
       {
         /* quit, fallback to CPU */
         goto cleanup;
@@ -4869,95 +4504,73 @@ RestoreMSCWarning
     }
   }
 
-  resizeFilterType = (int)GetResizeFilterWeightingType(resizeFilter);
-  resizeWindowType = (int)GetResizeFilterWindowWeightingType(resizeFilter);
-
-  verticalKernel = AcquireOpenCLKernel(device,"ResizeVerticalFilter");
-  if (verticalKernel == NULL)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  i = 0;
-  clStatus = clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_mem), (void*)&image);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_uint), (void*)&number_channels);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_uint), (void*)&columns);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_uint), (void*)&rows);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_mem), (void*)&resizedImage);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_uint), (void*)&resizedColumns);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_uint), (void*)&resizedRows);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&yFactor);
-
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(int), (void*)&resizeFilterType);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(int), (void*)&resizeWindowType);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(cl_mem), (void*)&resizeFilterCubicCoefficients);
-
-  resizeFilterScale = (float) GetResizeFilterScale(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterScale);
-
-  resizeFilterSupport = (float) GetResizeFilterSupport(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterSupport);
-
-  resizeFilterWindowSupport = (float) GetResizeFilterWindowSupport(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterWindowSupport);
-
-  resizeFilterBlur = (float) GetResizeFilterBlur(resizeFilter);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(float), (void*)&resizeFilterBlur);
-
-
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, imageCacheLocalMemorySize, NULL);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(int), &numCachedPixels);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), &pixelPerWorkgroup);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, sizeof(unsigned int), &chunkSize);
-  
-
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, pixelAccumulatorLocalMemorySize, NULL);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, weightAccumulatorLocalMemorySize, NULL);
-  clStatus |= clEnv->library->clSetKernelArg(verticalKernel, i++, gammaAccumulatorLocalMemorySize, NULL);
+  resizeFilterType=(int)GetResizeFilterWeightingType(resizeFilter);
+  resizeWindowType=(int)GetResizeFilterWindowWeightingType(resizeFilter);
 
-  if (clStatus != CL_SUCCESS)
+  verticalKernel=AcquireOpenCLKernel(device,"ResizeVerticalFilter");
+  if (verticalKernel == (cl_kernel) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  global_work_size[0] = resizedColumns;
-  global_work_size[1] = (resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*workgroupSize;
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
+
+  resizeFilterScale=(float) GetResizeFilterScale(resizeFilter);
+  resizeFilterSupport=(float) GetResizeFilterSupport(resizeFilter);
+  resizeFilterBlur=(float) GetResizeFilterBlur(resizeFilter);
+  resizeFilterWindowSupport=(float) GetResizeFilterWindowSupport(resizeFilter);
+
+  i=0;
+  status =SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_mem),(void*)&imageBuffer);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_uint),(void*)&number_channels);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_uint),(void*)&columns);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_uint),(void*)&rows);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_mem),(void*)&resizedImageBuffer);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_uint),(void*)&resizedColumns);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_uint),(void*)&resizedRows);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(float),(void*)&yFactor);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(int),(void*)&resizeFilterType);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(int),(void*)&resizeWindowType);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(cl_mem),(void*)&resizeFilterCubicCoefficients);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(float),(void*)&resizeFilterScale);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(float),(void*)&resizeFilterSupport);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(float),(void*)&resizeFilterWindowSupport);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(float),(void*)&resizeFilterBlur);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,imageCacheLocalMemorySize, NULL);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(int), &numCachedPixels);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(unsigned int), &pixelPerWorkgroup);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,sizeof(unsigned int), &chunkSize);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,pixelAccumulatorLocalMemorySize, NULL);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,weightAccumulatorLocalMemorySize, NULL);
+  status|=SetOpenCLKernelArg(verticalKernel,i++,gammaAccumulatorLocalMemorySize, NULL);
 
-  local_work_size[0] = 1;
-  local_work_size[1] = workgroupSize;
-  clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, verticalKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event);
-  if (clStatus != CL_SUCCESS)
+  if (status != CL_SUCCESS)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
     goto cleanup;
   }
-  RecordProfileData(device,verticalKernel,event);
-  status = MagickTrue;
 
+  gsize[0]=resizedColumns;
+  gsize[1]=(resizedRows+pixelPerWorkgroup-1)/pixelPerWorkgroup*
+    workgroupSize;
+  lsize[0]=1;
+  lsize[1]=workgroupSize;
+  outputReady=EnqueueOpenCLKernel(verticalKernel,2,(const size_t *) NULL,
+    gsize,lsize,image,filteredImage,exception);
 
 cleanup:
 
-  if (verticalKernel != NULL) RelinquishOpenCLKernel(verticalKernel);
+  if (verticalKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(verticalKernel);
 
-  return(status);
+  return(outputReady);
 }
 
 static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv,
   const size_t resizedColumns,const size_t resizedRows,
   const ResizeFilter *resizeFilter,ExceptionInfo *exception)
 {
-  CacheView
-    *filteredImage_view,
-    *image_view;
-
-  cl_command_queue
-    queue;
-
-  cl_int
-    clStatus;
-
   cl_mem
     cubicCoefficientsBuffer,
     filteredImageBuffer,
@@ -4971,13 +4584,12 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv,
     *resizeFilterCoefficient;
 
   float
-    *mappedCoefficientBuffer,
+    coefficientBuffer[7],
     xFactor,
     yFactor;
 
   MagickBooleanType
-    outputReady,
-    status;
+    outputReady;
 
   MagickCLDevice
     device;
@@ -4988,145 +4600,103 @@ static Image *ComputeResizeImage(const Image* image,MagickCLEnv clEnv,
   Image
     *filteredImage;
 
-  unsigned int
+  size_t
     i;
 
-  void
-    *filteredPixels;
-
-  outputReady = MagickFalse;
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  imageBuffer = NULL;
-  tempImageBuffer = NULL;
-  filteredImageBuffer = NULL;
-  filteredPixels = NULL;
-  cubicCoefficientsBuffer = NULL;
-  queue = NULL;
-
-  device = RequestOpenCLDevice(clEnv);
+  filteredImage=NULL;
+  tempImageBuffer=NULL;
+  cubicCoefficientsBuffer=NULL;
+  outputReady=MagickFalse;
 
-  image_view = AcquireAuthenticCacheView(image, exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
-  cubicCoefficientsBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY, 7 * sizeof(float), NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
-    goto cleanup;
-  }
-  queue = AcquireOpenCLCommandQueue(device);
-  mappedCoefficientBuffer = (float*)clEnv->library->clEnqueueMapBuffer(queue, cubicCoefficientsBuffer, CL_TRUE, CL_MAP_WRITE, 0, 7 * sizeof(float)
-          , 0, NULL, NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueMapBuffer failed.",".");
+  filteredImage=CloneImage(image,resizedColumns,resizedRows,MagickTrue,
+    exception);
+  if (filteredImage == (Image *) NULL)
     goto cleanup;
-  }
-  resizeFilterCoefficient = GetResizeFilterCoefficient(resizeFilter);
-  for (i = 0; i < 7; i++)
-  {
-    mappedCoefficientBuffer[i] = (float) resizeFilterCoefficient[i];
-  }
-  clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, cubicCoefficientsBuffer, mappedCoefficientBuffer, 0, NULL, NULL);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
+  if (filteredImageBuffer == (cl_mem) NULL)
     goto cleanup;
-  }
 
-  filteredImage = CloneImage(image,resizedColumns,resizedRows,MagickTrue,exception);
-  if (filteredImage == (Image *) NULL)
-    goto cleanup;
-  if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue)
+  resizeFilterCoefficient=GetResizeFilterCoefficient(resizeFilter);
+  for (i = 0; i < 7; i++)
+    coefficientBuffer[i]=(float) resizeFilterCoefficient[i];
+  cubicCoefficientsBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
+    CL_MEM_COPY_HOST_PTR,7*sizeof(float),&coefficientBuffer);
+  if (cubicCoefficientsBuffer == (cl_mem) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
     goto cleanup;
   }
-  filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
-  if (filteredImageBuffer == (cl_mem) NULL)
-    goto cleanup;
 
-  number_channels = image->number_channels;
+  number_channels=(cl_uint) image->number_channels;
   xFactor=(float) resizedColumns/(float) image->columns;
   yFactor=(float) resizedRows/(float) image->rows;
   if (xFactor > yFactor)
   {
-    length = resizedColumns*image->rows*number_channels;
-    tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length*sizeof(CLQuantum), NULL, &clStatus);
-    if (clStatus != CL_SUCCESS)
+    length=resizedColumns*image->rows*number_channels;
+    tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
+      sizeof(CLQuantum),(void *) NULL);
+    if (tempImageBuffer == (cl_mem) NULL)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
+      (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+        ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
       goto cleanup;
     }
 
-    status = resizeHorizontalFilter(clEnv,device,queue,imageBuffer,number_channels,
-      (cl_uint) image->columns,(cl_uint) image->rows,tempImageBuffer,
-      (cl_uint) resizedColumns,(cl_uint) image->rows,resizeFilter,
-      cubicCoefficientsBuffer,xFactor,exception);
-    if (status != MagickTrue)
+    outputReady=resizeHorizontalFilter(device,image,filteredImage,imageBuffer,
+      number_channels,(cl_uint) image->columns,(cl_uint) image->rows,
+      tempImageBuffer,(cl_uint) resizedColumns,(cl_uint) image->rows,
+      resizeFilter,cubicCoefficientsBuffer,xFactor,exception);
+    if (outputReady == MagickFalse)
       goto cleanup;
 
-    status = resizeVerticalFilter(clEnv,device,queue,tempImageBuffer,number_channels,
-      (cl_uint) resizedColumns,(cl_uint) image->rows,filteredImageBuffer,
-      (cl_uint) resizedColumns,(cl_uint) resizedRows,resizeFilter,
-      cubicCoefficientsBuffer,yFactor,exception);
-    if (status != MagickTrue)
+    outputReady=resizeVerticalFilter(device,image,filteredImage,tempImageBuffer,
+      number_channels,(cl_uint) resizedColumns,(cl_uint) image->rows,
+      filteredImageBuffer,(cl_uint) resizedColumns,(cl_uint) resizedRows,
+      resizeFilter,cubicCoefficientsBuffer,yFactor,exception);
+    if (outputReady == MagickFalse)
       goto cleanup;
   }
   else
   {
-    length = image->columns*resizedRows*number_channels;
-    tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length*sizeof(CLQuantum), NULL, &clStatus);
-    if (clStatus != CL_SUCCESS)
+    length=image->columns*resizedRows*number_channels;
+    tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,
+      length*sizeof(CLQuantum),(void *) NULL);
+    if (tempImageBuffer == (cl_mem) NULL)
     {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
+      (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+        ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
       goto cleanup;
     }
 
-    status = resizeVerticalFilter(clEnv,device,queue,imageBuffer,number_channels,
-      (cl_uint) image->columns,(cl_int) image->rows,tempImageBuffer,
-      (cl_uint) image->columns,(cl_uint) resizedRows,resizeFilter,
-      cubicCoefficientsBuffer,yFactor,exception);
-    if (status != MagickTrue)
+    outputReady=resizeVerticalFilter(device,image,filteredImage,imageBuffer,
+      number_channels,(cl_uint) image->columns,(cl_int) image->rows,
+      tempImageBuffer,(cl_uint) image->columns,(cl_uint) resizedRows,
+      resizeFilter,cubicCoefficientsBuffer,yFactor,exception);
+    if (outputReady == MagickFalse)
       goto cleanup;
 
-    status = resizeHorizontalFilter(clEnv,device,queue,tempImageBuffer,number_channels,
-      (cl_uint) image->columns, (cl_uint) resizedRows,filteredImageBuffer,
-      (cl_uint) resizedColumns, (cl_uint) resizedRows,resizeFilter,
-      cubicCoefficientsBuffer,xFactor,exception);
-    if (status != MagickTrue)
+    outputReady=resizeHorizontalFilter(device,image,filteredImage,tempImageBuffer,
+      number_channels,(cl_uint) image->columns, (cl_uint) resizedRows,
+      filteredImageBuffer,(cl_uint) resizedColumns, (cl_uint) resizedRows,
+      resizeFilter,cubicCoefficientsBuffer,xFactor,exception);
+    if (outputReady == MagickFalse)
       goto cleanup;
   }
 
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
-    goto cleanup;
-
-  outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
-
 cleanup:
 
-  image_view=DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view=DestroyCacheView(filteredImage_view);
-
-  if (imageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (tempImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(tempImageBuffer);
-  if (filteredImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (cubicCoefficientsBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(cubicCoefficientsBuffer);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (outputReady == MagickFalse && filteredImage != NULL)
+  if (tempImageBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(tempImageBuffer);
+  if (cubicCoefficientsBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(cubicCoefficientsBuffer);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
     filteredImage=DestroyImage(filteredImage);
 
   return(filteredImage);
@@ -5194,13 +4764,6 @@ MagickPrivate Image *AccelerateResizeImage(const Image *image,
 static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv,
   const double angle,ExceptionInfo *exception)
 {
-  CacheView
-    *image_view,
-    *filteredImage_view;
-
-  cl_command_queue
-    queue;
-
   cl_float2
     blurCenter;
 
@@ -5208,7 +4771,7 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv,
     biasPixel;
 
   cl_int
-    clStatus;
+    status;
 
   cl_mem
     cosThetaBuffer,
@@ -5219,9 +4782,6 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv,
   cl_kernel
     rotationalBlurKernel;
 
-  cl_event
-    event;
-
   cl_uint
     cossin_theta_size,
     number_channels;
@@ -5246,78 +4806,39 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv,
     bias;
 
   size_t
-    global_work_size[2];
-
-  unsigned int
+    gsize[2],
     i;
 
-  void
-    *filteredPixels;
-
-  outputReady = MagickFalse;
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  filteredPixels = NULL;
-  imageBuffer = NULL;
-  filteredImageBuffer = NULL;
-  sinThetaBuffer = NULL;
-  cosThetaBuffer = NULL;
-  queue = NULL;
-  rotationalBlurKernel = NULL;
-
-  device = RequestOpenCLDevice(clEnv);
+  filteredImage=NULL;
+  sinThetaBuffer=NULL;
+  cosThetaBuffer=NULL;
+  rotationalBlurKernel=NULL;
+  outputReady=MagickFalse;
 
-  image_view=AcquireAuthenticCacheView(image, exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
-  filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
+  filteredImage=CloneImage(image,image->columns,image->rows,MagickTrue,
+    exception);
   if (filteredImage == (Image *) NULL)
     goto cleanup;
-  if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue)
-  {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
-    goto cleanup;
-  }
-  filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,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;
+  blurCenter.s[0]=(float) (image->columns-1)/2.0;
+  blurCenter.s[1]=(float) (image->rows-1)/2.0;
   blurRadius=hypot(blurCenter.s[0],blurCenter.s[1]);
   cossin_theta_size=(unsigned int) fabs(4.0*DegreesToRadians(angle)*sqrt((double)blurRadius)+2UL);
 
-  /* create a buffer for sin_theta and cos_theta */
-  sinThetaBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
-    goto cleanup;
-  }
-  cosThetaBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, cossin_theta_size * sizeof(float), NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
-    goto cleanup;
-  }
-
-
-  queue = AcquireOpenCLCommandQueue(device);
-  sinThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, sinThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
+  cosThetaPtr=AcquireQuantumMemory(cossin_theta_size,sizeof(float));
+  if (cosThetaPtr == (float *) NULL)
     goto cleanup;
-  }
-
-  cosThetaPtr = (float*) clEnv->library->clEnqueueMapBuffer(queue, cosThetaBuffer, CL_TRUE, CL_MAP_WRITE, 0, cossin_theta_size*sizeof(float), 0, NULL, NULL, &clStatus);
-  if (clStatus != CL_SUCCESS)
+  sinThetaPtr=AcquireQuantumMemory(cossin_theta_size,sizeof(float));
+  if (sinThetaPtr == (float *) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnqueuemapBuffer failed.",".");
+    cosThetaPtr=RelinquishMagickMemory(cosThetaPtr);
     goto cleanup;
   }
 
@@ -5328,97 +4849,72 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv,
     cosThetaPtr[i]=(float)cos((double) (theta*i-offset));
     sinThetaPtr[i]=(float)sin((double) (theta*i-offset));
   }
-  clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, sinThetaBuffer, sinThetaPtr, 0, NULL, NULL);
-  clStatus |= clEnv->library->clEnqueueUnmapMemObject(queue, cosThetaBuffer, cosThetaPtr, 0, NULL, NULL);
-  if (clStatus != CL_SUCCESS)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueUnmapMemObject failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  /* get the OpenCL kernel */
-  rotationalBlurKernel = AcquireOpenCLKernel(device,"RotationalBlur");
-  if (rotationalBlurKernel == NULL)
-  {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  GetPixelInfo(image,&bias);
-  biasPixel.s[0] = bias.red;
-  biasPixel.s[1] = bias.green;
-  biasPixel.s[2] = bias.blue;
-  biasPixel.s[3] = bias.alpha;
 
-  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(cl_uint), &cossin_theta_size);
-  clStatus|=clEnv->library->clSetKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
-  if (clStatus != CL_SUCCESS)
+  sinThetaBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
+    CL_MEM_COPY_HOST_PTR,cossin_theta_size*sizeof(float),sinThetaPtr);
+  sinThetaPtr=RelinquishMagickMemory(sinThetaPtr);
+  cosThetaBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_ONLY |
+    CL_MEM_COPY_HOST_PTR,cossin_theta_size*sizeof(float),cosThetaPtr);
+  cosThetaPtr=RelinquishMagickMemory(cosThetaPtr);
+  if ((sinThetaBuffer == (cl_mem) NULL) || (cosThetaBuffer == (cl_mem)NULL))
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
     goto cleanup;
   }
 
-  global_work_size[0] = image->columns;
-  global_work_size[1] = image->rows;
-  /* launch the kernel */
-  clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, rotationalBlurKernel, 2, NULL, global_work_size, NULL, 0, NULL, &event);
-  if (clStatus != CL_SUCCESS)
+  rotationalBlurKernel=AcquireOpenCLKernel(device,"RotationalBlur");
+  if (rotationalBlurKernel == (cl_kernel) NULL)
   {
-    (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
     goto cleanup;
   }
-  RecordProfileData(device,rotationalBlurKernel,event);
 
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
+  GetPixelInfo(image,&bias);
+  biasPixel.s[0]=bias.red;
+  biasPixel.s[1]=bias.green;
+  biasPixel.s[2]=bias.blue;
+  biasPixel.s[3]=bias.alpha;
+
+  number_channels=(cl_uint) image->number_channels;
+
+  i=0;
+  status =SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_uint),&number_channels);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(ChannelType), &image->channel_mask);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_float4), &biasPixel);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_float2), &blurCenter);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&cosThetaBuffer);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&sinThetaBuffer);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_uint), &cossin_theta_size);
+  status|=SetOpenCLKernelArg(rotationalBlurKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+  if (status != CL_SUCCESS)
   {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"clEnv->library->clSetKernelArg failed.",".");
     goto cleanup;
   }
 
-  outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
+  gsize[0]=image->columns;
+  gsize[1]=image->rows;
+  outputReady=EnqueueOpenCLKernel(rotationalBlurKernel,2,(const size_t *) NULL,
+    gsize,(const size_t *) NULL,image,filteredImage,exception);
 
 cleanup:
 
-  image_view=DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view=DestroyCacheView(filteredImage_view);
-
-  if (filteredImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (imageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (sinThetaBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(sinThetaBuffer);
-  if (cosThetaBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(cosThetaBuffer);
-  if (rotationalBlurKernel!=NULL)
-    RelinquishOpenCLKernel(rotationalBlurKernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (outputReady == MagickFalse)
-  {
-    if (filteredImage != NULL)
-    {
-      DestroyImage(filteredImage);
-      filteredImage = NULL;
-    }
-  }
+  if (sinThetaBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(sinThetaBuffer);
+  if (cosThetaBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(cosThetaBuffer);
+  if (rotationalBlurKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(rotationalBlurKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
+    filteredImage=DestroyImage(filteredImage);
 
-  return filteredImage;
+  return(filteredImage);
 }
 
 MagickPrivate Image* AccelerateRotationalBlurImage(const Image *image,
@@ -5460,23 +4956,13 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv,
   const double radius,const double sigma,const double gain,
   const double threshold,ExceptionInfo *exception)
 {
-  CacheView
-    *filteredImage_view,
-    *image_view;
-
-  cl_command_queue
-    queue;
-
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
     blurRowKernel,
     unsharpMaskBlurColumnKernel;
 
-  cl_event
-    event;
-
   cl_mem
     filteredImageBuffer,
     imageBuffer,
@@ -5508,210 +4994,135 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv,
   MagickSizeType
     length;
 
-  void
-    *filteredPixels;
-
-  unsigned int
-    i;
-
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  imageBuffer = NULL;
-  filteredImageBuffer = NULL;
-  filteredPixels = NULL;
-  tempImageBuffer = NULL;
-  imageKernelBuffer = NULL;
-  blurRowKernel = NULL;
-  unsharpMaskBlurColumnKernel = NULL;
-  queue = NULL;
-  outputReady = MagickFalse;
+  size_t
+    gsize[2],
+    i,
+    lsize[2];
 
-  device = RequestOpenCLDevice(clEnv);
-  queue = AcquireOpenCLCommandQueue(device);
+  filteredImage=NULL;
+  tempImageBuffer=NULL;
+  imageKernelBuffer=NULL;
+  blurRowKernel=NULL;
+  unsharpMaskBlurColumnKernel=NULL;
+  outputReady=MagickFalse;
 
-  image_view = AcquireAuthenticCacheView(image, exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
   filteredImage=CloneImage(image,0,0,MagickTrue,exception);
   if (filteredImage == (Image *) NULL)
     goto cleanup;
-  if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue)
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
+  if (filteredImageBuffer == (cl_mem) NULL)
+    goto cleanup;
+
+  imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
+    exception);
+
+  length=image->columns*image->rows;
+  tempImageBuffer=CreateOpenCLBuffer(device,CL_MEM_READ_WRITE,length*
+    sizeof(cl_float4),NULL);
+  if (tempImageBuffer == (cl_mem) NULL)
   {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"CreateOpenCLBuffer failed.",".");
     goto cleanup;
   }
-  filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
-  if (filteredImageBuffer == (cl_mem) NULL)
+
+  blurRowKernel=AcquireOpenCLKernel(device,"BlurRow");
+  if (blurRowKernel == (cl_kernel) NULL)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
     goto cleanup;
+  }
+
+  unsharpMaskBlurColumnKernel=AcquireOpenCLKernel(device,
+    "UnsharpMaskBlurColumn");
+  if (unsharpMaskBlurColumnKernel == (cl_kernel) NULL)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
 
-  imageKernelBuffer=createKernelInfo(clEnv,device,queue,radius,sigma,
-    &kernelWidth,exception);
+  number_channels=(cl_uint) image->number_channels;
+  imageColumns=(cl_uint) image->columns;
+  imageRows=(cl_uint) image->rows;
 
+  chunkSize = 256;
+
+  i=0;
+  status =SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
+  status|=SetOpenCLKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
+  if (status != CL_SUCCESS)
   {
-    /* create temp buffer */
-    {
-      length = image->columns * image->rows;
-      tempImageBuffer = clEnv->library->clCreateBuffer(device->context, CL_MEM_READ_WRITE, length * sizeof(cl_float4), NULL, &clStatus);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clCreateBuffer failed.",".");
-        goto cleanup;
-      }
-    }
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"clEnv->library->clSetKernelArg failed.",".");
+    goto cleanup;
+  }
 
-    /* get the opencl kernel */
-    {
-      blurRowKernel = AcquireOpenCLKernel(device,"BlurRow");
-      if (blurRowKernel == NULL)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-        goto cleanup;
-      };
+  gsize[0]=chunkSize*((image->columns+chunkSize-1)/chunkSize);
+  gsize[1]=image->rows;
+  lsize[0]=chunkSize;
+  lsize[1]=1;
+  outputReady=EnqueueOpenCLKernel(blurRowKernel,2,(const size_t *) NULL,gsize,
+    lsize,image,filteredImage,exception);
+  
+  chunkSize=256;
+  fGain=(float) gain;
+  fThreshold=(float) threshold;
+
+  i=0;
+  status =SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),&number_channels);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,(chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,kernelWidth*sizeof(float),NULL);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
+  status|=SetOpenCLKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+  if (status != CL_SUCCESS)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"clEnv->library->clSetKernelArg failed.",".");
+    goto cleanup;
+  }
 
-      unsharpMaskBlurColumnKernel = AcquireOpenCLKernel(device,"UnsharpMaskBlurColumn");
-      if (unsharpMaskBlurColumnKernel == NULL)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-        goto cleanup;
-      };
-    }
+  gsize[0]=image->columns;
+  gsize[1]=chunkSize*((image->rows+chunkSize-1)/chunkSize);
+  lsize[0]=1;
+  lsize[1]=chunkSize;
+  outputReady=EnqueueOpenCLKernel(unsharpMaskBlurColumnKernel,2,
+    (const size_t *) NULL,gsize,lsize,image,filteredImage,exception);
 
-    number_channels = (cl_uint) image->number_channels;
-    imageColumns = (cl_uint) image->columns;
-    imageRows = (cl_uint) image->rows;
+cleanup:
 
-    {
-      chunkSize = 256;
-
-      /* set the kernel arguments */
-      i = 0;
-      clStatus=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),&number_channels);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(ChannelType),&image->channel_mask);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_uint),(void *)&imageRows);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_float4)*(chunkSize+kernelWidth),(void *) NULL);
-      clStatus|=clEnv->library->clSetKernelArg(blurRowKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-        goto cleanup;
-      }
-    }
-
-    /* launch the kernel */
-    {
-      size_t gsize[2];
-      size_t wsize[2];
-
-      gsize[0] = chunkSize*((image->columns+chunkSize-1)/chunkSize);
-      gsize[1] = image->rows;
-      wsize[0] = chunkSize;
-      wsize[1] = 1;
-
-      clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, blurRowKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-        goto cleanup;
-      }
-      RecordProfileData(device,blurRowKernel,event);
-    }
-
-
-    {
-      chunkSize = 256;
-      fGain = (float) gain;
-      fThreshold = (float) threshold;
-
-      i = 0;
-      clStatus=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&tempImageBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),&number_channels);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(ChannelType),&image->channel_mask);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&imageRows);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,(chunkSize+kernelWidth-1)*sizeof(cl_float4),NULL);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,kernelWidth*sizeof(float),NULL);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fGain);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(float),(void *)&fThreshold);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskBlurColumnKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
-
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-        goto cleanup;
-      }
-    }
-
-    /* launch the kernel */
-    {
-      size_t gsize[2];
-      size_t wsize[2];
-
-      gsize[0] = image->columns;
-      gsize[1] = chunkSize*((image->rows+chunkSize-1)/chunkSize);
-      wsize[0] = 1;
-      wsize[1] = chunkSize;
-
-      clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskBlurColumnKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-        goto cleanup;
-      }
-      RecordProfileData(device,unsharpMaskBlurColumnKernel,event);
-    }
-
-  }
-
-  /* get result */
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
-  {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
-
-cleanup:
-
-  image_view=DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view=DestroyCacheView(filteredImage_view);
-
-  if (imageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (filteredImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (tempImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(tempImageBuffer);
-  if (imageKernelBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageKernelBuffer);
-  if (blurRowKernel!=NULL)
-    RelinquishOpenCLKernel(blurRowKernel);
-  if (unsharpMaskBlurColumnKernel!=NULL)
-    RelinquishOpenCLKernel(unsharpMaskBlurColumnKernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (outputReady == MagickFalse)
-  {
-    if (filteredImage != NULL)
-    {
-      DestroyImage(filteredImage);
-      filteredImage = NULL;
-    }
-  }
+  if (tempImageBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(tempImageBuffer);
+  if (imageKernelBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(imageKernelBuffer);
+  if (blurRowKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(blurRowKernel);
+  if (unsharpMaskBlurColumnKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(unsharpMaskBlurColumnKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
+    filteredImage=DestroyImage(filteredImage);
 
   return(filteredImage);
 }
@@ -5720,29 +5131,18 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image,
   MagickCLEnv clEnv,const double radius,const double sigma,const double gain,
   const double threshold,ExceptionInfo *exception)
 {
-  CacheView
-    *filteredImage_view,
-    *image_view;
-
-  cl_command_queue
-    queue;
-
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
     unsharpMaskKernel;
 
-  cl_event
-    event;
-
   cl_mem
     filteredImageBuffer,
     imageBuffer,
     imageKernelBuffer;
 
   cl_uint
-    i,
     imageColumns,
     imageRows,
     kernelWidth,
@@ -5761,129 +5161,80 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image,
   MagickCLDevice
     device;
 
-  void
-    *filteredPixels;
-
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  filteredPixels = NULL;
-  imageBuffer = NULL;
-  filteredImageBuffer = NULL;
-  imageKernelBuffer = NULL;
-  unsharpMaskKernel = NULL;
-  queue = NULL;
-  outputReady = MagickFalse;
+  size_t
+    gsize[2],
+    i,
+    lsize[2];
 
-  device = RequestOpenCLDevice(clEnv);
-  queue = AcquireOpenCLCommandQueue(device);
+  filteredImage=NULL;
+  imageKernelBuffer=NULL;
+  unsharpMaskKernel=NULL;
+  outputReady=MagickFalse;
 
-  image_view=AcquireAuthenticCacheView(image,exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
-  filteredImage = CloneImage(image,image->columns,image->rows,MagickTrue,exception);
+  filteredImage=CloneImage(image,0,0,MagickTrue,exception);
   if (filteredImage == (Image *) NULL)
     goto cleanup;
-
-  filteredImage_view=AcquireAuthenticCacheView(filteredImage,exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
-  if (filteredImageBuffer == (void *) NULL)
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
+  if (filteredImageBuffer == (cl_mem) NULL)
     goto cleanup;
 
-  imageKernelBuffer=createKernelInfo(clEnv,device,queue,radius,sigma,
-    &kernelWidth,exception);
+  imageKernelBuffer=createKernelInfo(device,radius,sigma,&kernelWidth,
+    exception);
 
+  unsharpMaskKernel=AcquireOpenCLKernel(device,"UnsharpMask");
+  if (unsharpMaskKernel == NULL)
   {
-    /* get the opencl kernel */
-    {
-      unsharpMaskKernel = AcquireOpenCLKernel(device, "UnsharpMask");
-      if (unsharpMaskKernel == NULL)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-        goto cleanup;
-      };
-    }
-
-    {
-      imageColumns = (cl_uint) image->columns;
-      imageRows = (cl_uint) image->rows;
-      number_channels = (cl_uint) image->number_channels;
-      fGain = (float) gain;
-      fThreshold = (float) threshold;
-
-      /* set the kernel arguments */
-      i = 0;
-      clStatus=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&number_channels);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageRows);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernelWidth)),(void *) NULL);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold);
-      clStatus|=clEnv->library->clSetKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-        goto cleanup;
-      }
-    }
-
-    /* launch the kernel */
-    {
-      size_t gsize[2];
-      size_t wsize[2];
-
-      gsize[0] = ((image->columns + 7) / 8) * 8;
-      gsize[1] = ((image->rows + 31) / 32) * 32;
-      wsize[0] = 8;
-      wsize[1] = 32;
-
-      clStatus = clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, unsharpMaskKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-        goto cleanup;
-      }
-      RecordProfileData(device,unsharpMaskKernel,event);
-    }
-  }
-
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
+
+  imageColumns=(cl_uint) image->columns;
+  imageRows=(cl_uint) image->rows;
+  number_channels=(cl_uint) image->number_channels;
+  fGain=(float) gain;
+  fThreshold=(float) threshold;
+
+  i=0;
+  status =SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&number_channels);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(ChannelType),(void *)&image->channel_mask);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&imageKernelBuffer);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&kernelWidth);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageColumns);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_uint),(void *)&imageRows);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_float4)*(8 * (32 + kernelWidth)),(void *) NULL);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fGain);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(float),(void *)&fThreshold);
+  status|=SetOpenCLKernelArg(unsharpMaskKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+  if (status != CL_SUCCESS)
+  {
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
     goto cleanup;
+  }
 
-  outputReady=SyncCacheViewAuthenticPixels(filteredImage_view,exception);
+  gsize[0]=((image->columns + 7) / 8)*8;
+  gsize[1]=((image->rows + 31) / 32)*32;
+  lsize[0]=8;
+  lsize[1]=32;
+  outputReady=EnqueueOpenCLKernel(unsharpMaskKernel,2,(const size_t *) NULL,
+    gsize,lsize,image,filteredImage,exception);
 
 cleanup:
 
-  image_view=DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view=DestroyCacheView(filteredImage_view);
-
-  if (imageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (filteredImageBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (imageKernelBuffer!=NULL)
-    clEnv->library->clReleaseMemObject(imageKernelBuffer);
-  if (unsharpMaskKernel!=NULL)
-    RelinquishOpenCLKernel(unsharpMaskKernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (outputReady == MagickFalse)
-  {
-    if (filteredImage != NULL)
-    {
-      DestroyImage(filteredImage);
-      filteredImage = NULL;
-    }
-  }
+  if (imageKernelBuffer != (cl_mem) NULL)
+    ReleaseOpenCLMemObject(imageKernelBuffer);
+  if (unsharpMaskKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(unsharpMaskKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
+    filteredImage=DestroyImage(filteredImage);
 
   return(filteredImage);
 }
@@ -5920,26 +5271,33 @@ MagickPrivate Image *AccelerateUnsharpMaskImage(const Image *image,
 static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv,
   const double threshold,ExceptionInfo *exception)
 {
-  CacheView
-    *filteredImage_view,
-    *image_view;
+  const cl_int
+    PASSES=5;
 
-  cl_command_queue
-    queue;
+  const int
+    TILESIZE=64,
+    PAD=1<<(PASSES-1),
+    SIZE=TILESIZE-2*PAD;
+
+  cl_float
+    thresh;
 
   cl_int
-    clStatus;
+    status;
 
   cl_kernel
     denoiseKernel;
 
-  cl_event
-    event;
-
   cl_mem
     filteredImageBuffer,
     imageBuffer;
 
+  cl_uint
+    number_channels,
+    width,
+    height,
+    max_channels;
+
   Image
     *filteredImage;
 
@@ -5949,133 +5307,73 @@ static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv,
   MagickCLDevice
     device;
 
-  void
-    *filteredPixels;
-
-  unsigned int
-    i;
+  size_t
+    gsize[2],
+    i,
+    lsize[2];
 
-  filteredImage = NULL;
-  filteredImage_view = NULL;
-  filteredImageBuffer = NULL;
-  filteredPixels = NULL;
-  denoiseKernel = NULL;
-  outputReady = MagickFalse;
+  filteredImage=NULL;
+  denoiseKernel=NULL;
+  outputReady=MagickFalse;
 
-  device = RequestOpenCLDevice(clEnv);
-  queue = AcquireOpenCLCommandQueue(device);
-
-  /* Create and initialize OpenCL buffers. */
-  image_view = AcquireAuthenticCacheView(image, exception);
-  imageBuffer=createReadBuffer(image,image_view,clEnv,device,exception);
+  device=RequestOpenCLDevice(clEnv);
+  imageBuffer=GetAuthenticOpenCLBuffer(image,device,exception);
   if (imageBuffer == (cl_mem) NULL)
     goto cleanup;
-
-  /* create output */
   filteredImage=CloneImage(image,0,0,MagickTrue,exception);
   if (filteredImage == (Image *) NULL)
     goto cleanup;
-  if (SetImageStorageClass(filteredImage, DirectClass, exception) != MagickTrue)
-  {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "CloneImage failed.", "'%s'", ".");
-    goto cleanup;
-  }
-  filteredImage_view = AcquireAuthenticCacheView(filteredImage, exception);
-  filteredImageBuffer=createWriteBuffer(filteredImage,filteredImage_view,clEnv,
-    device,filteredPixels,exception);
+  filteredImageBuffer=GetAuthenticOpenCLBuffer(filteredImage,device,exception);
   if (filteredImageBuffer == (cl_mem) NULL)
     goto cleanup;
 
-  /* get the opencl kernel */
-  denoiseKernel = AcquireOpenCLKernel(device,"WaveletDenoise");
-  if (denoiseKernel == NULL)
-  {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "AcquireOpenCLKernel failed.", "'%s'", ".");
-    goto cleanup;
-  }
-
-  // Process image
+  denoiseKernel=AcquireOpenCLKernel(device,"WaveletDenoise");
+  if (denoiseKernel == (cl_kernel) NULL)
   {
-    const int PASSES = 5;
-    cl_uint number_channels = (cl_uint)image->number_channels;
-    cl_uint width = (cl_uint)image->columns;
-    cl_uint height = (cl_uint)image->rows;
-    cl_uint max_channels = number_channels;
-    if ((max_channels == 4) || (max_channels == 2))
-      max_channels=max_channels-1;
-    cl_float thresh = threshold;
-
-    /* set the kernel arguments */
-    i = 0;
-    clStatus = clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_mem), (void *)&imageBuffer);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_mem), (void *)&filteredImageBuffer);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&number_channels);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&max_channels);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_float), (void *)&thresh);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_int), (void *)&PASSES);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&width);
-    clStatus |= clEnv->library->clSetKernelArg(denoiseKernel, i++, sizeof(cl_uint), (void *)&height);
-    if (clStatus != CL_SUCCESS)
-    {
-      (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clSetKernelArg failed.", "'%s'", ".");
-      goto cleanup;
-    }
-
-    {
-      const int TILESIZE = 64;
-      const int PAD = 1 << (PASSES - 1);
-      const int SIZE = TILESIZE - 2 * PAD;
-
-      size_t gsize[2];
-      size_t wsize[2];
-
-      gsize[0] = ((width + (SIZE - 1)) / SIZE) * TILESIZE;
-      gsize[1] = ((height + (SIZE - 1)) / SIZE) * 4;
-      wsize[0] = TILESIZE;
-      wsize[1] = 4;
-
-      clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, denoiseKernel, 2, NULL, gsize, wsize, 0, NULL, &event);
-      if (clStatus != CL_SUCCESS)
-      {
-        (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
-        goto cleanup;
-      }
-    }
-    RecordProfileData(device,denoiseKernel,event);
-  }
-
-  if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"AcquireOpenCLKernel failed.",".");
+    goto cleanup;
+  }
+
+  number_channels=(cl_uint)image->number_channels;
+  width=(cl_uint)image->columns;
+  height=(cl_uint)image->rows;
+  max_channels=number_channels;
+  if ((max_channels == 4) || (max_channels == 2))
+    max_channels=max_channels-1;
+  thresh=threshold;
+
+  i=0;
+  status =SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_mem),(void *)&imageBuffer);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_mem),(void *)&filteredImageBuffer);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_uint),(void *)&number_channels);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_uint),(void *)&max_channels);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_float),(void *)&thresh);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_int),(void *)&PASSES);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_uint),(void *)&width);
+  status|=SetOpenCLKernelArg(denoiseKernel,i++,sizeof(cl_uint),(void *)&height);
+  if (status != CL_SUCCESS)
   {
-    (void)OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "Reading output image from CL buffer failed.", "'%s'", ".");
+    (void) OpenCLThrowMagickException(device,exception,GetMagickModule(),
+      ResourceLimitWarning,"SetOpenCLKernelArg failed.",".");
     goto cleanup;
   }
 
-  outputReady = SyncCacheViewAuthenticPixels(filteredImage_view, exception);
+  gsize[0]=((width+(SIZE-1))/SIZE)*TILESIZE;
+  gsize[1]=((height+(SIZE-1))/SIZE)*4;
+  lsize[0]=TILESIZE;
+  lsize[1]=4;
+  outputReady=EnqueueOpenCLKernel(denoiseKernel,2,(const size_t *) NULL,gsize,
+    lsize,image,filteredImage,exception);
 
 cleanup:
 
-  image_view = DestroyCacheView(image_view);
-  if (filteredImage_view != NULL)
-    filteredImage_view = DestroyCacheView(filteredImage_view);
-
-  if (imageBuffer != NULL)
-    clEnv->library->clReleaseMemObject(imageBuffer);
-  if (filteredImageBuffer != NULL)
-    clEnv->library->clReleaseMemObject(filteredImageBuffer);
-  if (denoiseKernel != NULL)
-    RelinquishOpenCLKernel(denoiseKernel);
-  if (queue != NULL)
-    RelinquishOpenCLCommandQueue(device,queue);
-  if (device != NULL)
-    ReleaseOpenCLDevice(clEnv,device);
-  if (outputReady == MagickFalse)
-  {
-    if (filteredImage != NULL)
-    {
-      DestroyImage(filteredImage);
-      filteredImage = NULL;
-    }
-  }
+  if (denoiseKernel != (cl_kernel) NULL)
+    ReleaseOpenCLKernel(denoiseKernel);
+  if (device != (MagickCLDevice) NULL)
+    ReleaseOpenCLDevice(device);
+  if ((outputReady == MagickFalse) && (filteredImage != (Image *) NULL))
+    filteredImage=DestroyImage(filteredImage);
 
   return(filteredImage);
 }
@@ -6103,230 +5401,4 @@ MagickPrivate Image *AccelerateWaveletDenoiseImage(const Image *image,
 
   return(filteredImage);
 }
-
-#else  /* MAGICKCORE_OPENCL_SUPPORT  */
-
-MagickPrivate Image *AccelerateAddNoiseImage(const Image *magick_unused(image),
-  const NoiseType magick_unused(noise_type),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(noise_type);
-  magick_unreferenced(exception);
-  return((Image *) NULL);
-}
-
-MagickPrivate Image *AccelerateBlurImage(const Image *magick_unused(image),
-  const double magick_unused(radius),const double magick_unused(sigma),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(radius);
-  magick_unreferenced(sigma);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate MagickBooleanType AccelerateCompositeImage(
-  Image *magick_unused(image),const CompositeOperator magick_unused(compose),
-  const Image *magick_unused(composite),
-  const float magick_unused(destination_dissolve),
-  const float magick_unused(source_dissolve),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(compose);
-  magick_unreferenced(composite);
-  magick_unreferenced(destination_dissolve);
-  magick_unreferenced(source_dissolve);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate MagickBooleanType AccelerateContrastImage(
-  Image* magick_unused(image),const MagickBooleanType magick_unused(sharpen),
-  ExceptionInfo* magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(sharpen);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate MagickBooleanType AccelerateContrastStretchImage(
-  Image *magick_unused(image),const double magick_unused(black_point),
-  const double magick_unused(white_point),
-  ExceptionInfo* magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(black_point);
-  magick_unreferenced(white_point);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate Image *AccelerateConvolveImage(const Image *magick_unused(image),
-  const KernelInfo *magick_unused(kernel),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(kernel);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate MagickBooleanType AccelerateEqualizeImage(
-  Image* magick_unused(image),ExceptionInfo* magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate Image *AccelerateDespeckleImage(const Image* magick_unused(image),
-  ExceptionInfo* magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate MagickBooleanType AccelerateFunctionImage(
-  Image *magick_unused(image),
-  const MagickFunction magick_unused(function),
-  const size_t magick_unused(number_parameters),
-  const double *magick_unused(parameters),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(function);
-  magick_unreferenced(number_parameters);
-  magick_unreferenced(parameters);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate MagickBooleanType AccelerateGrayscaleImage(
-  Image *magick_unused(image),const PixelIntensityMethod magick_unused(method),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(method);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate Image *AccelerateLocalContrastImage(
-  const Image *magick_unused(image),const double magick_unused(radius),
-  const double magick_unused(strength),ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(radius);
-  magick_unreferenced(strength);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate MagickBooleanType AccelerateModulateImage(
-  Image *magick_unused(image),const double magick_unused(percent_brightness),
-  const double magick_unused(percent_hue),
-  const double magick_unused(percent_saturation),
-  ColorspaceType magick_unused(colorspace),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(percent_brightness);
-  magick_unreferenced(percent_hue);
-  magick_unreferenced(percent_saturation);
-  magick_unreferenced(colorspace);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate Image *AccelerateMotionBlurImage(
-  const Image *magick_unused(image),const double *magick_unused(kernel),
-  const size_t magick_unused(width),const OffsetInfo *magick_unused(offset),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(kernel);
-  magick_unreferenced(width);
-  magick_unreferenced(offset);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate MagickBooleanType AccelerateRandomImage(
-  Image *magick_unused(image),ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(exception);
-
-  return(MagickFalse);
-}
-
-MagickPrivate Image *AccelerateResizeImage(const Image *magick_unused(image),
-  const size_t magick_unused(resizedColumns),
-  const size_t magick_unused(resizedRows),
-  const ResizeFilter *magick_unused(resizeFilter),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(resizedColumns);
-  magick_unreferenced(resizedRows);
-  magick_unreferenced(resizeFilter);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate Image *AccelerateRotationalBlurImage(
-  const Image *magick_unused(image),const double magick_unused(angle),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(angle);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate Image *AccelerateUnsharpMaskImage(
-  const Image *magick_unused(image),const double magick_unused(radius),
-  const double magick_unused(sigma),const double magick_unused(gain),
-  const double magick_unused(threshold),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(radius);
-  magick_unreferenced(sigma);
-  magick_unreferenced(gain);
-  magick_unreferenced(threshold);
-  magick_unreferenced(exception);
-
-  return((Image *) NULL);
-}
-
-MagickPrivate Image *AccelerateWaveletDenoiseImage(
-  const Image *magick_unused(image),const double magick_unused(threshold),
-  ExceptionInfo *magick_unused(exception))
-{
-  magick_unreferenced(image);
-  magick_unreferenced(threshold);
-  magick_unreferenced(exception);
-
-  return((Image *)NULL);
-}
 #endif /* MAGICKCORE_OPENCL_SUPPORT */
index 87bd9ac54c9227e24b959652127081aaf502cc15..9cf9c33c4a144abefa590124e2c4b53410c66cc5 100644 (file)
@@ -20,6 +20,7 @@
 
 #include "MagickCore/cache.h"
 #include "MagickCore/distribute-cache.h"
+#include "MagickCore/opencl-private.h"
 #include "MagickCore/pixel.h"
 #include "MagickCore/random_.h"
 #include "MagickCore/thread-private.h"
@@ -219,6 +220,9 @@ typedef struct _CacheInfo
 
   size_t
     signature;
+
+  MagickCLCacheInfo
+    opencl;
 } CacheInfo;
 
 extern MagickPrivate Cache
@@ -279,6 +283,14 @@ extern MagickPrivate void
   ResetPixelCacheEpoch(void),
   SetPixelCacheMethods(Cache,CacheMethods *);
 
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+extern MagickPrivate cl_mem
+  GetAuthenticOpenCLBuffer(const Image *,MagickCLDevice,ExceptionInfo *);
+
+extern MagickPrivate void
+  SyncAuthenticOpenCLBuffer(const Image *);
+#endif
+
 #if defined(__cplusplus) || defined(c_plusplus)
 }
 #endif
index ba4afd7e6be81637d418692e3addc7ba78ddf91c..c9eacbe64da07ee5dba79a49949cd783f561da11 100644 (file)
@@ -152,11 +152,14 @@ MagickExport CacheView *AcquireVirtualCacheView(const Image *image,
   CacheView
     *magick_restrict cache_view;
 
+  magick_unreferenced(exception);
   assert(image != (Image *) NULL);
   assert(image->signature == MagickCoreSignature);
   if (image->debug != MagickFalse)
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
-  (void) exception;
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+  SyncAuthenticOpenCLBuffer(image);
+#endif
   cache_view=(CacheView *) MagickAssumeAligned(AcquireAlignedMemory(1,
     sizeof(*cache_view)));
   if (cache_view == (CacheView *) NULL)
index e535b5a0de1277aeadd7e711adab602a651d5a3c..111927952305b5720655ef76ba18b465d285b6b2 100644 (file)
@@ -138,6 +138,11 @@ static Quantum
   *SetPixelCacheNexusPixels(const CacheInfo *,const MapMode,
     const RectangleInfo *,NexusInfo *,ExceptionInfo *) magick_hot_spot;
 
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+static void
+  CopyOpenCLBuffer(CacheInfo *magick_restrict);
+#endif
+
 #if defined(__cplusplus) || defined(c_plusplus)
 }
 #endif
@@ -866,14 +871,19 @@ static inline void RelinquishPixelCachePixels(CacheInfo *cache_info)
   {
     case MemoryCache:
     {
-      if (cache_info->mapped == MagickFalse)
-        cache_info->pixels=(Quantum *) RelinquishAlignedMemory(
-          cache_info->pixels);
-      else
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+      if (cache_info->opencl != (MagickCLCacheInfo) NULL)
         {
-          (void) UnmapBlob(cache_info->pixels,(size_t) cache_info->length);
+          cache_info->opencl=RelinquishMagickCLCacheInfo(cache_info->opencl,
+            MagickTrue);
           cache_info->pixels=(Quantum *) NULL;
+          break;
         }
+#endif
+      if (cache_info->mapped == MagickFalse)
+        cache_info->pixels=RelinquishAlignedMemory(cache_info->pixels);
+      else
+        (void) UnmapBlob(cache_info->pixels,(size_t) cache_info->length);
       RelinquishMagickResource(MemoryResource,cache_info->length);
       break;
     }
@@ -1106,6 +1116,64 @@ static void *GetAuthenticMetacontentFromCache(const Image *image)
   assert(id < (int) cache_info->number_threads);
   return(cache_info->nexus_info[id]->metacontent);
 }
+
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
++   G e t A u t h e n t i c O p e n C L B u f f e r                           %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  GetAuthenticOpenCLBuffer() returns an OpenCL buffer used to execute OpenCL
+%  operations.
+%
+%  The format of the GetAuthenticOpenCLBuffer() method is:
+%
+%      cl_mem GetAuthenticOpenCLBuffer(const Image *image,
+%        MagickCLDevice device,ExceptionInfo *exception)
+%
+%  A description of each parameter follows:
+%
+%    o image: the image.
+%
+%    o device: the device to use.
+%
+%    o exception: return any errors or warnings in this structure.
+%
+*/
+MagickPrivate cl_mem GetAuthenticOpenCLBuffer(const Image *image,
+  MagickCLDevice device,ExceptionInfo *exception)
+{
+  CacheInfo
+    *magick_restrict cache_info;
+
+  cl_int
+    status;
+
+  assert(image != (const Image *) NULL);
+  assert(device != (const MagickCLDevice) NULL);
+  cache_info=(CacheInfo *) image->cache;
+  if (cache_info->type == UndefinedCache)
+    SyncImagePixelCache((Image *) image,exception);
+  if (cache_info->type != MemoryCache || cache_info->mapped != MagickFalse)
+    return((cl_mem) NULL);
+  if ((cache_info->opencl != (MagickCLCacheInfo) NULL) &&
+      (cache_info->opencl->device->context != device->context))
+    cache_info->opencl=CopyMagickCLCacheInfo(cache_info->opencl);
+  if (cache_info->opencl == (MagickCLCacheInfo) NULL)
+    {
+      assert(cache_info->pixels != NULL);
+      cache_info->opencl=AcquireMagickCLCacheInfo(device,cache_info->pixels,
+        cache_info->length);
+    }
+  return(cache_info->opencl->buffer);
+}
+#endif
 \f
 /*
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
@@ -1265,7 +1333,8 @@ MagickExport Quantum *GetAuthenticPixelQueue(const Image *image)
 %   G e t A u t h e n t i c P i x e l s                                       %
 %                                                                             %
 %                                                                             %
-%                                                                             % %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 %
 %  GetAuthenticPixels() obtains a pixel region for read/write access. If the
 %  region is successfully accessed, a pointer to a Quantum array
@@ -1531,6 +1600,9 @@ static Cache GetImagePixelCache(Image *image,const MagickBooleanType clone,
   LockSemaphoreInfo(image->semaphore);
   assert(image->cache != (Cache) NULL);
   cache_info=(CacheInfo *) image->cache;
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+  CopyOpenCLBuffer(cache_info);
+#endif
   destroy=MagickFalse;
   if ((cache_info->reference_count > 1) || (cache_info->mode == ReadMode))
     {
@@ -3681,6 +3753,9 @@ MagickExport MagickBooleanType PersistPixelCache(Image *image,
   page_size=GetMagickPageSize();
   cache_info=(CacheInfo *) image->cache;
   assert(cache_info->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+  CopyOpenCLBuffer(cache_info);
+#endif
   if (attach != MagickFalse)
     {
       /*
@@ -4824,6 +4899,56 @@ MagickPrivate VirtualPixelMethod SetPixelCacheVirtualMethod(Image *image,
     }
   return(method);
 }
+
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
++   S y n c A u t h e n t i c O p e n C L B u f f e r                         %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  SyncAuthenticOpenCLBuffer() makes sure that all the OpenCL operations have
+%  been completed and updates the host memory.
+%
+%  The format of the SyncAuthenticOpenCLBuffer() method is:
+%
+%      void SyncAuthenticOpenCLBuffer(const Image *image)
+%
+%  A description of each parameter follows:
+%
+%    o image: the image.
+%
+*/
+static void CopyOpenCLBuffer(CacheInfo *magick_restrict cache_info)
+{
+  assert(cache_info != (CacheInfo *) NULL);
+  assert(cache_info->signature == MagickCoreSignature);
+  if ((cache_info->type != MemoryCache) ||
+      (cache_info->opencl == (MagickCLCacheInfo) NULL))
+    return;
+  /*
+  Ensure single threaded access to OpenCL environment.
+  */
+  LockSemaphoreInfo(cache_info->semaphore);
+  cache_info->opencl=CopyMagickCLCacheInfo(cache_info->opencl);
+  UnlockSemaphoreInfo(cache_info->semaphore);
+}
+
+MagickPrivate void SyncAuthenticOpenCLBuffer(const Image *image)
+{
+  CacheInfo
+    *magick_restrict cache_info;
+
+  assert(image != (const Image *) NULL);
+  cache_info=(CacheInfo *) image->cache;
+  CopyOpenCLBuffer(cache_info);
+}
+#endif
 \f
 /*
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
index 2062cfd18bfdf2d36178296c1b8de284ea66a4d3..e18879ba5ebcb6abb0d2df53910dc3b796a14f08 100644 (file)
@@ -785,9 +785,11 @@ MagickExport Image *BlurImage(const Image *image,const double radius,
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   blur_image=AccelerateBlurImage(image,radius,sigma,exception);
   if (blur_image != (Image *) NULL)
     return(blur_image);
+#endif
   (void) FormatLocaleString(geometry,MagickPathExtent,
     "blur:%.20gx%.20g;blur:%.20gx%.20g+90",radius,sigma,radius,sigma);
   kernel_info=AcquireKernelInfo(geometry,exception);
@@ -831,9 +833,11 @@ MagickExport Image *ConvolveImage(const Image *image,
   Image
     *convolve_image;
 
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   convolve_image=AccelerateConvolveImage(image,kernel_info,exception);
   if (convolve_image != (Image *) NULL)
     return(convolve_image);
+#endif
 
   convolve_image=MorphologyImage(image,ConvolveMorphology,1,kernel_info,
     exception);
@@ -1007,9 +1011,11 @@ MagickExport Image *DespeckleImage(const Image *image,ExceptionInfo *exception)
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   despeckle_image=AccelerateDespeckleImage(image,exception);
   if (despeckle_image != (Image *) NULL)
     return(despeckle_image);
+#endif
   despeckle_image=CloneImage(image,0,0,MagickTrue,exception);
   if (despeckle_image == (Image *) NULL)
     return((Image *) NULL);
@@ -1706,9 +1712,11 @@ MagickExport Image *LocalContrastImage(const Image *image,const double radius,
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   contrast_image=AccelerateLocalContrastImage(image,radius,strength,exception);
   if (contrast_image != (Image *) NULL)
     return(contrast_image);
+#endif
   contrast_image=CloneImage(image,0,0,MagickTrue,exception);
   if (contrast_image == (Image *) NULL)
     return((Image *) NULL);
@@ -2808,9 +2816,11 @@ MagickExport Image *RotationalBlurImage(const Image *image,const double angle,
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   blur_image=AccelerateRotationalBlurImage(image,angle,exception);
   if (blur_image != (Image *) NULL)
     return(blur_image);
+#endif
   blur_image=CloneImage(image,image->columns,image->rows,MagickTrue,exception);
   if (blur_image == (Image *) NULL)
     return((Image *) NULL);
@@ -3903,10 +3913,12 @@ MagickExport Image *UnsharpMaskImage(const Image *image,const double radius,
   if (image->debug != MagickFalse)
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   unsharp_image=AccelerateUnsharpMaskImage(image,radius,sigma,gain,threshold,
     exception);
   if (unsharp_image != (Image *) NULL)
     return(unsharp_image);
+#endif
   unsharp_image=BlurImage(image,radius,sigma,exception);
   if (unsharp_image == (Image *) NULL)
     return((Image *) NULL);
index f055df1c59564641b022040a48eff9c7c4232c26..5d49ffc30f502c5faf5a1caafc79ae5f84e66959 100644 (file)
@@ -887,8 +887,10 @@ MagickExport MagickBooleanType ContrastImage(Image *image,
 
   assert(image != (Image *) NULL);
   assert(image->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   if (AccelerateContrastImage(image,sharpen,exception) != MagickFalse)
     return(MagickTrue);
+#endif
   if (image->debug != MagickFalse)
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   sign=sharpen != MagickFalse ? 1 : -1;
@@ -1525,8 +1527,10 @@ MagickExport MagickBooleanType EqualizeImage(Image *image,
   */
   assert(image != (Image *) NULL);
   assert(image->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   if (AccelerateEqualizeImage(image,exception) != MagickFalse)
     return(MagickTrue);
+#endif
   if (image->debug != MagickFalse)
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   equalize_map=(double *) AcquireQuantumMemory(MaxMap+1UL,
@@ -1978,12 +1982,14 @@ MagickExport MagickBooleanType GrayscaleImage(Image *image,
       if (SetImageStorageClass(image,DirectClass,exception) == MagickFalse)
         return(MagickFalse);
     }
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   if (AccelerateGrayscaleImage(image,method,exception) != MagickFalse)
     {
       image->intensity=method;
       image->type=GrayscaleType;
       return(SetImageColorspace(image,GRAYColorspace,exception));
     }
+#endif
   /*
     Grayscale image.
   */
@@ -3274,9 +3280,11 @@ MagickExport MagickBooleanType ModulateImage(Image *image,const char *modulate,
   /*
     Modulate image.
   */
-  if(AccelerateModulateImage(image,percent_brightness,percent_hue,
-     percent_saturation,colorspace,exception) != MagickFalse)
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
+  if (AccelerateModulateImage(image,percent_brightness,percent_hue,
+        percent_saturation,colorspace,exception) != MagickFalse)
     return(MagickTrue);
+#endif
   status=MagickTrue;
   progress=0;
   image_view=AcquireAuthenticCacheView(image,exception);
index 70b63d095c3c0171bb3da36f9020b250189ab783..09d7a2a7a07ee650cc0366820c7f79e2299dc822 100644 (file)
@@ -305,9 +305,11 @@ MagickExport Image *AddNoiseImage(const Image *image,const NoiseType noise_type,
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   noise_image=AccelerateAddNoiseImage(image,noise_type,exception);
   if (noise_image != (Image *) NULL)
     return(noise_image);
+#endif
   noise_image=CloneImage(image,image->columns,image->rows,MagickTrue,exception);
   if (noise_image == (Image *) NULL)
     return((Image *) NULL);
@@ -5847,10 +5849,11 @@ MagickExport Image *WaveletDenoiseImage(const Image *image,
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
-  noise_image=(Image *) NULL;
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   noise_image=AccelerateWaveletDenoiseImage(image,threshold,exception);
   if (noise_image != (Image *) NULL)
     return(noise_image);
+#endif
   noise_image=CloneImage(image,0,0,MagickTrue,exception);
   if (noise_image == (Image *) NULL)
     return((Image *) NULL);
index c7e4b7269512287f7a9fa1e8f3f2967fd6ade639..49190b2a1321d4fb1afeb1fd3e204d2eb85efcf0 100644 (file)
@@ -30,15 +30,28 @@ extern "C" {
 #endif
 
 #if !defined(MAGICKCORE_OPENCL_SUPPORT)
-  typedef void* cl_context;
-  typedef void* cl_command_queue;
-  typedef void* cl_device_id;
-  typedef void* cl_event;
-  typedef void* cl_kernel;
-  typedef void* cl_mem;
-  typedef void* cl_platform_id;
-  typedef void* cl_device_type;
+typedef void* MagickCLCacheInfo;
 #else
+typedef struct _MagickCLCacheInfo
+{
+  cl_event
+    *events;
+
+  cl_mem
+    buffer;
+
+  cl_uint
+    event_count;
+
+  MagickCLDevice
+    device;
+
+  MagickSizeType
+    length;
+
+  Quantum
+    *pixels;
+}* MagickCLCacheInfo;
 
 /*
   Define declarations.
@@ -184,12 +197,7 @@ typedef CL_API_ENTRY cl_int
     CL_API_SUFFIX__VERSION_1_0;
 
 
-/* Profiling APIs */
-typedef CL_API_ENTRY cl_int
-  (CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)(cl_event event,
-    cl_profiling_info param_name,size_t param_value_size,void *param_value,
-    size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
-
+/* Events APIs */
 typedef CL_API_ENTRY cl_int
   (CL_API_CALL *MAGICKpfn_clWaitForEvents)(cl_uint num_events,
     const cl_event *event_list) CL_API_SUFFIX__VERSION_1_0;
@@ -198,8 +206,24 @@ typedef CL_API_ENTRY cl_int
   (CL_API_CALL *MAGICKpfn_clReleaseEvent)(cl_event event)
     CL_API_SUFFIX__VERSION_1_0;
 
+typedef CL_API_ENTRY cl_int
+  (CL_API_CALL *MAGICKpfn_clRetainEvent)(cl_event event)
+    CL_API_SUFFIX__VERSION_1_0;
 
-/* Finish APIs, only here for GetAndLockRandSeedBuffer */
+typedef CL_API_ENTRY cl_int
+  (CL_API_CALL *MAGICKpfn_clSetEventCallback)(cl_event event,
+    cl_int command_exec_callback_type,void (CL_CALLBACK *MAGICKpfn_notify)(
+      cl_event,cl_int,void *),void *user_data) CL_API_SUFFIX__VERSION_1_1;
+
+
+/* Profiling APIs */
+typedef CL_API_ENTRY cl_int
+  (CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)(cl_event event,
+    cl_profiling_info param_name,size_t param_value_size,void *param_value,
+    size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+
+
+/* Finish APIs */
 typedef CL_API_ENTRY cl_int
   (CL_API_CALL *MAGICKpfn_clFinish)(cl_command_queue command_queue)
     CL_API_SUFFIX__VERSION_1_0;
@@ -240,9 +264,12 @@ struct MagickLibraryRec
   MAGICKpfn_clEnqueueUnmapMemObject   clEnqueueUnmapMemObject;
   MAGICKpfn_clEnqueueNDRangeKernel    clEnqueueNDRangeKernel;
 
-  MAGICKpfn_clGetEventProfilingInfo   clGetEventProfilingInfo;
   MAGICKpfn_clWaitForEvents           clWaitForEvents;
   MAGICKpfn_clReleaseEvent            clReleaseEvent;
+  MAGICKpfn_clRetainEvent             clRetainEvent;
+  MAGICKpfn_clSetEventCallback        clSetEventCallback;
+
+  MAGICKpfn_clGetEventProfilingInfo   clGetEventProfilingInfo;
 
   MAGICKpfn_clFinish                  clFinish;
 };
@@ -296,7 +323,7 @@ struct _MagickCLDevice
     command_queues_index;
 };
 
-struct _MagickCLEnv
+typedef struct _MagickCLEnv
 {
   cl_context
     *contexts;
@@ -323,9 +350,7 @@ struct _MagickCLEnv
   size_t
     number_contexts,
     number_devices;
-};
-
-#endif
+} *MagickCLEnv;
 
 #if defined(MAGICKCORE_HDRI_SUPPORT)
 #define CLOptions "-cl-single-precision-constant -cl-mad-enable -DMAGICKCORE_HDRI_SUPPORT=1 "\
@@ -372,15 +397,28 @@ struct _MagickCLEnv
 extern MagickPrivate cl_command_queue
   AcquireOpenCLCommandQueue(MagickCLDevice);
 
+extern MagickPrivate cl_int
+  SetOpenCLKernelArg(cl_kernel,cl_uint,size_t,const void *);
+
 extern MagickPrivate cl_kernel
   AcquireOpenCLKernel(MagickCLDevice,const char *);
 
+extern MagickPrivate cl_mem
+  CreateOpenCLBuffer(MagickCLDevice,cl_mem_flags,size_t,void *);
+
 extern MagickPrivate MagickBooleanType
+  EnqueueOpenCLKernel(cl_kernel,cl_uint,const size_t *,const size_t *,
+    const size_t *,const Image *,const Image *,ExceptionInfo *),
   InitializeOpenCL(MagickCLEnv,ExceptionInfo *),
   OpenCLThrowMagickException(MagickCLDevice,ExceptionInfo *,
     const char *,const char *,const size_t,const ExceptionType,const char *,
     const char *,...);
 
+extern MagickPrivate MagickCLCacheInfo
+  AcquireMagickCLCacheInfo(MagickCLDevice,Quantum *,const MagickSizeType),
+  CopyMagickCLCacheInfo(MagickCLCacheInfo),
+  RelinquishMagickCLCacheInfo(MagickCLCacheInfo,const MagickBooleanType);
+
 extern MagickPrivate MagickCLDevice
   RequestOpenCLDevice(MagickCLEnv);
 
@@ -394,9 +432,13 @@ extern MagickPrivate void
   DumpOpenCLProfileData(),
   OpenCLTerminus(),
   RecordProfileData(MagickCLDevice,cl_kernel,cl_event),
-  ReleaseOpenCLDevice(MagickCLEnv,MagickCLDevice),
+  ReleaseOpenCLDevice(MagickCLDevice),
+  ReleaseOpenCLKernel(cl_kernel),
+  ReleaseOpenCLMemObject(cl_mem),
   RelinquishOpenCLCommandQueue(MagickCLDevice,cl_command_queue),
-  RelinquishOpenCLKernel(cl_kernel);
+  RetainOpenCLEvent(cl_event);
+
+#endif
 
 #if defined(__cplusplus) || defined(c_plusplus)
 }
index 24ede2c426732b14817054b6e308e29c6d705a02..703757fe0ee14c3c185dc1a99191ac3bec4e0eed 100644 (file)
@@ -43,6 +43,7 @@
 #include "MagickCore/studio.h"
 #include "MagickCore/artifact.h"
 #include "MagickCore/cache.h"
+#include "MagickCore/cache-private.h"
 #include "MagickCore/color.h"
 #include "MagickCore/compare.h"
 #include "MagickCore/constitute.h"
@@ -145,38 +146,6 @@ static void
 extern const char
   *accelerateKernels, *accelerateKernels2;
 
-/*
-  static declarations.
-*/
-static const char *kernelNames[] =
-{
-  "AddNoise",
-  "BlurColumn",
-  "BlurRow",
-  "Composite",
-  "Contrast",
-  "ContrastStretch",
-  "Convolve",
-  "ConvolveOptimized",
-  "ComputeFunction",
-  "Equalize",
-  "GrayScale",
-  "Histogram",
-  "HullPass1",
-  "HullPass2",
-  "LocalContrastBlurApplyColumn",
-  "LocalContrastBlurRow",
-  "Modulate",
-  "MotionBlur",
-  "ResizeHorizontal",
-  "ResizeVertical",
-  "RotationalBlur",
-  "UnsharpMask",
-  "UnsharpMaskBlurColumn",
-  "WaveletDenoise",
-  "NONE"
-};
-
 /* OpenCL library */
 MagickLibrary
   *openCL_library;
@@ -187,7 +156,7 @@ MagickCLEnv
 MagickThreadType
   test_thread_id=0;
 SemaphoreInfo
-  *default_CLEnv_Lock;
+  *openCL_lock;
 
 /* Cached location of the OpenCL cache files */
 char
@@ -464,6 +433,99 @@ static size_t StringSignature(const char* string)
   return(signature);
 }
 
+static MagickCLCacheInfo DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
+{
+  ssize_t
+    i;
+
+  for (i=0; i < (ssize_t) info->event_count; i++)
+    openCL_library->clReleaseEvent(info->events[i]);
+  info->events=RelinquishMagickMemory(info->events);
+  if (info->buffer != (cl_mem) NULL)
+    {
+      openCL_library->clReleaseMemObject(info->buffer);
+      info->buffer=(cl_mem) NULL;
+    }
+  ReleaseOpenCLDevice(info->device);
+  return((MagickCLCacheInfo) RelinquishMagickMemory(info));
+}
+
+/*
+  Provide call to OpenCL library methods
+*/
+
+MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
+  cl_mem_flags flags, size_t size, void *host_ptr)
+{
+  return(openCL_library->clCreateBuffer(device->context, flags, size, host_ptr,
+    (cl_int *) NULL));
+}
+
+MagickPrivate void ReleaseOpenCLKernel(cl_kernel kernel)
+{
+  (void) openCL_library->clReleaseKernel(kernel);
+}
+
+MagickPrivate void ReleaseOpenCLMemObject(cl_mem memobj)
+{
+  (void) openCL_library->clReleaseMemObject(memobj);
+}
+
+MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,cl_uint arg_index,
+  size_t arg_size,const void *arg_value)
+{
+  return(openCL_library->clSetKernelArg(kernel,arg_index,arg_size,arg_value));
+}
+
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
++   A c q u i r e M a g i c k C L C a c h e I n f o                           %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  AcquireMagickCLCacheInfo() acquires an OpenCL cache info structure.
+%
+%  The format of the AcquireMagickCLCacheInfo method is:
+%
+%      MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
+%        Quantum *pixels,const MagickSizeType length)
+%
+%  A description of each parameter follows:
+%
+%    o device: the OpenCL device.
+%
+%    o pixels: the pixel buffer of the image.
+%
+%    o length: the length of the pixel buffer.
+%
+*/
+
+MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
+  Quantum *pixels,const MagickSizeType length)
+{
+  MagickCLCacheInfo
+    info;
+
+  info=(MagickCLCacheInfo) AcquireMagickMemory(sizeof(*info));
+  if (info == (MagickCLCacheInfo) NULL)
+    ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
+  (void) ResetMagickMemory(info,0,sizeof(*info));
+  LockSemaphoreInfo(openCL_lock);
+  device->requested++;
+  UnlockSemaphoreInfo(openCL_lock);
+  info->device=device;
+  info->length=length;
+  info->pixels=pixels;
+  info->buffer=openCL_library->clCreateBuffer(device->context,
+    CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,NULL);
+  return(info);
+}
+
 /*
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 %                                                                             %
@@ -542,7 +604,7 @@ static MagickCLEnv AcquireMagickCLEnv(void)
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   A c q u i r e O p e n C L C o m m a n d Q u e u e                         %
++   A c q u i r e O p e n C L C o m m a n d Q u e u e                         %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -593,7 +655,7 @@ MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   A c q u i r e O p e n C L K e r n e l                                     %
++   A c q u i r e O p e n C L K e r n e l                                     %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -953,7 +1015,7 @@ static void AutoSelectOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
 %    o exception: return any errors or warnings
 */
 
-static double RunOpenCLBenchmark()
+static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
 {
   AccelerateTimer
     timer;
@@ -994,6 +1056,21 @@ static double RunOpenCLBenchmark()
     resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
       exception);
 
+    /* 
+      We need this to get a proper performance benchmark, the operations
+      are executed asynchronous.
+    */
+    if (is_cpu == MagickFalse)
+      {
+        CacheInfo
+          *cache_info;
+
+        cache_info=(CacheInfo *) resizedImage->cache;
+        if (cache_info->opencl != (MagickCLCacheInfo) NULL)
+          openCL_library->clWaitForEvents(cache_info->opencl->event_count,
+            cache_info->opencl->events);
+      }
+
     if (i > 0)
       StopAccelerateTimer(&timer);
 
@@ -1013,7 +1090,7 @@ static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
 {
   testEnv->devices[0]=device;
   default_CLEnv=testEnv;
-  device->score=RunOpenCLBenchmark();
+  device->score=RunOpenCLBenchmark(MagickFalse);
   default_CLEnv=clEnv;
   testEnv->devices[0]=(MagickCLDevice) NULL;
 }
@@ -1117,7 +1194,7 @@ static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
 
   testEnv->enabled=MagickFalse;
   default_CLEnv=testEnv;
-  clEnv->cpu_score=RunOpenCLBenchmark();
+  clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
   default_CLEnv=clEnv;
 
   testEnv=RelinquishMagickCLEnv(testEnv);
@@ -1307,7 +1384,48 @@ static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   D u m p O p e n C L P r o f i l e D a t a                                 %
++   C o p y M a g i c k C L C a c h e I n f o                                 %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  CopyMagickCLCacheInfo() copies the memory from the device into host memory.
+%
+%  The format of the CopyMagickCLCacheInfo method is:
+%
+%      void CopyMagickCLCacheInfo(MagickCLCacheInfo info)
+%
+%  A description of each parameter follows:
+%
+%    o info: the OpenCL cache info.
+%
+*/
+MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
+{
+  cl_command_queue
+    queue;
+
+  Quantum
+    *pixels;
+
+  if (info == (MagickCLCacheInfo) NULL || info->event_count == 0)
+    return((MagickCLCacheInfo) NULL);
+  queue=AcquireOpenCLCommandQueue(info->device);
+  pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
+    CL_MAP_READ | CL_MAP_WRITE,0,info->length,info->event_count,info->events,
+    NULL,NULL);
+  assert(pixels == info->pixels);
+  RelinquishOpenCLCommandQueue(info->device,queue);
+  return(RelinquishMagickCLCacheInfo(info,MagickFalse));
+}
+
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
++   D u m p O p e n C L P r o f i l e D a t a                                 %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -1379,8 +1497,8 @@ MagickPrivate void DumpOpenCLProfileData()
 
       profile=device->profile_records[j];
       strcpy(indent,"                    ");
-      strncpy(indent,kernelNames[j],MagickMin(strlen(kernelNames[j]),
-        strlen(indent)-1));
+      strncpy(indent,profile->kernel_name,MagickMin(strlen(
+        profile->kernel_name),strlen(indent)-1));
       sprintf(buf,"%s %7d %7d %7d %7d",indent,(int) (profile->total/
         profile->count),(int) profile->count,(int) profile->min,
         (int) profile->max);
@@ -1392,13 +1510,158 @@ MagickPrivate void DumpOpenCLProfileData()
   }
   fclose(log);
 }
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
++   E n q u e u e O p e n C L K e r n e l                                     %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
+%  events with the images.
+%
+%  The format of the EnqueueOpenCLKernel method is:
+%
+%      MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
+%        const size_t *global_work_offset,const size_t *global_work_size,
+%        const size_t *local_work_size,const Image *input_image,
+%        const Image *output_image,ExceptionInfo *exception)
+%
+%  A description of each parameter follows:
+%
+%    o kernel: the OpenCL kernel.
+%
+%    o work_dim: the number of dimensions used to specify the global work-items
+%                and work-items in the work-group.
+%
+%    o offset: can be used to specify an array of work_dim unsigned values
+%              that describe the offset used to calculate the global ID of a
+%              work-item.
+%
+%    o gsize: points to an array of work_dim unsigned values that describe the
+%             number of global work-items in work_dim dimensions that will
+%             execute the kernel function.
+%
+%    o lsize: points to an array of work_dim unsigned values that describe the
+%             number of work-items that make up a work-group that will execute
+%             the kernel specified by kernel.
+%
+%    o input_image: the input image of the operation.
+%
+%    o output_image: the output or secondairy image of the operation.
+%
+%    o exception: return any errors or warnings in this structure.
+%
+*/
+
+extern void RegisterCacheEvent(MagickCLCacheInfo info,cl_event event)
+{
+  assert(info != (MagickCLCacheInfo) NULL);
+  assert(event != (cl_event) NULL);
+  if (info->events == (cl_event *) NULL)
+    {
+      info->events=AcquireMagickMemory(sizeof(*info->events));
+      info->event_count=1;
+    }
+  else
+    info->events=ResizeQuantumMemory(info->events,++info->event_count,
+      sizeof(*info->events));
+  if (info->events == (cl_event *) NULL)
+    ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
+  info->events[info->event_count-1]=event;
+  openCL_library->clRetainEvent(event);
+}
+
+MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,
+  cl_uint work_dim,const size_t *offset,const size_t *gsize,
+  const size_t *lsize,const Image *input_image,const Image *output_image,
+  ExceptionInfo *exception)
+{
+  CacheInfo
+    *output_info,
+    *input_info;
+
+  cl_command_queue
+    queue;
+
+  cl_event
+    event,
+    *events;
+
+  cl_int
+    status;
+
+  cl_uint
+    event_count;
+
+  assert(input_image != (const Image *) NULL);
+  input_info=(CacheInfo *) input_image->cache;
+  assert(input_info != (CacheInfo *) NULL);
+  assert(input_info->opencl != (MagickCLCacheInfo) NULL);
+  queue=AcquireOpenCLCommandQueue(input_info->opencl->device);
+  if (queue == (cl_command_queue) NULL)
+    return(MagickFalse);
+  event_count=input_info->opencl->event_count;
+  events=input_info->opencl->events;
+  output_info=(CacheInfo *) NULL;
+  if (output_image != (const Image *) NULL)
+    {
+      output_info=(CacheInfo *) output_image->cache;
+      assert(output_info != (CacheInfo *) NULL);
+      assert(output_info->opencl != (MagickCLCacheInfo) NULL);
+      if (output_info->opencl->event_count > 0)
+        {
+          ssize_t
+            i;
+
+          event_count+=output_info->opencl->event_count;
+          events=AcquireQuantumMemory(event_count,sizeof(*events));
+          if (events == (cl_event *) NULL)
+            {
+              RelinquishOpenCLCommandQueue(input_info->opencl->device,queue);
+              return(MagickFalse);
+            }
+          for (i=0; i < (ssize_t) event_count; i++)
+          {
+            if (i < input_info->opencl->event_count)
+              events[i]=input_info->opencl->events[i];
+            else
+              events[i]=output_info->opencl->events[i-
+                input_info->opencl->event_count];
+          }
+        }
+    }
+  status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
+    gsize,lsize,event_count,events,&event);
+  RelinquishOpenCLCommandQueue(input_info->opencl->device,queue);
+  if ((output_info != (CacheInfo *) NULL) &&
+      (output_info->opencl->event_count > 0))
+    events=(cl_event *) RelinquishMagickMemory(events);
+  if (status != CL_SUCCESS)
+    {
+      (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
+      GetMagickModule(),ResourceLimitWarning,"clEnqueueNDRangeKernel failed.",
+      "'%s'",".");
+      return(MagickFalse);
+    }
+  RegisterCacheEvent(input_info->opencl,event);
+  if (output_info != (CacheInfo *) NULL)
+    RegisterCacheEvent(output_info->opencl,event);
+  RecordProfileData(input_info->opencl->device,kernel,event);
+  openCL_library->clReleaseEvent(event);
+  return(MagickTrue);
+}
 
 /*
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   G e t C u r r u n t O p e n C L E n v                                     %
++   G e t C u r r u n t O p e n C L E n v                                     %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -1423,13 +1686,13 @@ MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
       return(default_CLEnv);
   }
 
-  if (default_CLEnv_Lock == (SemaphoreInfo *) NULL)
-    ActivateSemaphoreInfo(&default_CLEnv_Lock);
+  if (openCL_lock == (SemaphoreInfo *) NULL)
+    ActivateSemaphoreInfo(&openCL_lock);
 
-  LockSemaphoreInfo(default_CLEnv_Lock);
+  LockSemaphoreInfo(openCL_lock);
   if (default_CLEnv == (MagickCLEnv) NULL)
     default_CLEnv=AcquireMagickCLEnv();
-  UnlockSemaphoreInfo(default_CLEnv_Lock);
+  UnlockSemaphoreInfo(openCL_lock);
 
   return(default_CLEnv);
 }
@@ -1826,7 +2089,7 @@ static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   I n i t i a l i z e O p e n C L                                           %
++   I n i t i a l i z e O p e n C L                                           %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2117,9 +2380,12 @@ static MagickBooleanType BindOpenCLFunctions()
   BIND(clEnqueueUnmapMemObject);
   BIND(clEnqueueNDRangeKernel);
 
-  BIND(clGetEventProfilingInfo);
   BIND(clWaitForEvents);
   BIND(clReleaseEvent);
+  BIND(clRetainEvent);
+  BIND(clSetEventCallback);
+
+  BIND(clGetEventProfilingInfo);
 
   BIND(clFinish);
 
@@ -2146,7 +2412,7 @@ static MagickBooleanType LoadOpenCLLibrary(void)
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   O p e n C L T e r m i n u s                                               %
++   O p e n C L T e r m i n u s                                               %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2169,8 +2435,8 @@ MagickPrivate void OpenCLTerminus()
     RelinquishSemaphoreInfo(&cache_directory_lock);
   if (default_CLEnv != (MagickCLEnv) NULL)
     default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
-  if (default_CLEnv_Lock != (SemaphoreInfo *) NULL)
-    RelinquishSemaphoreInfo(&default_CLEnv_Lock);
+  if (openCL_lock != (SemaphoreInfo *) NULL)
+    RelinquishSemaphoreInfo(&openCL_lock);
   if (openCL_library != (MagickLibrary *) NULL)
     openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
 }
@@ -2180,7 +2446,7 @@ MagickPrivate void OpenCLTerminus()
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   O p e n C L T h r o w M a g i c k E x c e p t i o n                       %
++   O p e n C L T h r o w M a g i c k E x c e p t i o n                       %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2263,7 +2529,7 @@ MagickPrivate MagickBooleanType OpenCLThrowMagickException(
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   R e c o r d P r o f i l e D a t a                                         %
++   R e c o r d P r o f i l e D a t a                                         %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2306,17 +2572,11 @@ MagickPrivate void RecordProfileData(MagickCLDevice device,
     length;
 
   if (device->profile_kernels == MagickFalse)
-    {
-      openCL_library->clReleaseEvent(event);
-      return;
-    }
+    return;
   status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
     &length);
   if (status != CL_SUCCESS)
-    {
-      openCL_library->clReleaseEvent(event);
-      return;
-    }
+    return;
   name=AcquireQuantumMemory(length,sizeof(*name));
   (void) openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
     name,NULL);
@@ -2326,7 +2586,6 @@ MagickPrivate void RecordProfileData(MagickCLDevice device,
     CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
   status&=openCL_library->clGetEventProfilingInfo(event,
     CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
-  openCL_library->clReleaseEvent(event);
   if (status != CL_SUCCESS)
     {
       name=DestroyString(name);
@@ -2375,7 +2634,7 @@ MagickPrivate void RecordProfileData(MagickCLDevice device,
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   R e l e a s e  M a g i c k C L D e v i c e                                %
++   R e l e a s e  M a g i c k C L D e v i c e                                %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2385,21 +2644,75 @@ MagickPrivate void RecordProfileData(MagickCLDevice device,
 %
 %  The format of the ReleaseOpenCLDevice method is:
 %
-%      void ReleaseOpenCLDevice(MagickCLEnv clEnv,MagickCLDevice device)
+%      void ReleaseOpenCLDevice(MagickCLDevice device)
 %
 %  A description of each parameter follows:
 %
-%    o clEnv: the OpenCL environment.
-%
 %    o device: the OpenCL device to be released.
 %
 */
 
-MagickPrivate void ReleaseOpenCLDevice(MagickCLEnv clEnv,MagickCLDevice device)
+MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
 {
-  LockSemaphoreInfo(clEnv->lock);
+  assert(device != (MagickCLDevice) NULL);
+  LockSemaphoreInfo(openCL_lock);
   device->requested--;
-  UnlockSemaphoreInfo(clEnv->lock);
+  UnlockSemaphoreInfo(openCL_lock);
+}
+
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
++   R e l i n q u i s h M a g i c k C L C a c h e I n f o                     %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  RelinquishMagickCLCacheInfo() frees memory acquired with
+%  AcquireMagickCLCacheInfo()
+%
+%  The format of the RelinquishMagickCLCacheInfo method is:
+%
+%      MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
+%        const MagickBooleanType relinquish_pixels)
+%
+%  A description of each parameter follows:
+%
+%    o info: the OpenCL cache info.
+%
+%    o relinquish_pixels: the pixels will be relinquish when set to true.
+%
+*/
+
+static void CL_API_CALL DestroyMagickCLCacheInfoDelayed(
+  cl_event magick_unused(event),cl_int magick_unused(event_command_exec_status),
+  void *user_data)
+{
+  MagickCLCacheInfo
+    info;
+
+  magick_unreferenced(event);
+  magick_unreferenced(event_command_exec_status);
+  info=(MagickCLCacheInfo) user_data;
+  (void) RelinquishAlignedMemory(info->pixels);
+  RelinquishMagickResource(MemoryResource,info->length);
+  DestroyMagickCLCacheInfo(info);
+}
+
+MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
+  MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
+{
+  if (info == (MagickCLCacheInfo) NULL)
+    return((MagickCLCacheInfo) NULL);
+  if (relinquish_pixels)
+    openCL_library->clSetEventCallback(info->events[info->event_count-1],
+      CL_COMPLETE,&DestroyMagickCLCacheInfoDelayed,info);
+  else
+    DestroyMagickCLCacheInfo(info);
+  return((MagickCLCacheInfo) NULL);
 }
 
 /*
@@ -2489,7 +2802,7 @@ static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   R e l i n q u i s h O p e n C L C o m m a n d Q u e u e                   %
++   R e l i n q u i s h O p e n C L C o m m a n d Q u e u e                   %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2521,45 +2834,16 @@ MagickPrivate void RelinquishOpenCLCommandQueue(MagickCLDevice device,
   LockSemaphoreInfo(device->lock);
   if ((device->profile_kernels != MagickFalse) ||
       (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES - 1))
-  {
-    UnlockSemaphoreInfo(device->lock);
-    (void)openCL_library->clReleaseCommandQueue(queue);
-  }
+    {
+      UnlockSemaphoreInfo(device->lock);
+      openCL_library->clFinish(queue);
+      (void) openCL_library->clReleaseCommandQueue(queue);
+    }
   else
-  {
-    device->command_queues[++device->command_queues_index] = queue;
-    UnlockSemaphoreInfo(device->lock);
-  }
-}
-
-/*
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-%                                                                             %
-%                                                                             %
-%                                                                             %
-%   R e l i n q u i s h O p e n C L K e r n e l                               %
-%                                                                             %
-%                                                                             %
-%                                                                             %
-%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
-%
-%  RelinquishOpenCLKernel() releases an OpenCL kernel
-%
-%  The format of the RelinquishOpenCLKernel method is:
-%
-%    void RelinquishOpenCLKernel(cl_kernel kernel)
-%
-%  A description of each parameter follows:
-%
-%    o kernel: the OpenCL kernel object to be released.
-%
-%
-*/
-
-MagickPrivate void RelinquishOpenCLKernel(cl_kernel kernel)
-{
-  if (kernel != (cl_kernel) NULL)
-    (void) openCL_library->clReleaseKernel(kernel);
+    {
+      device->command_queues[++device->command_queues_index] = queue;
+      UnlockSemaphoreInfo(device->lock);
+    }
 }
 
 /*
@@ -2567,7 +2851,7 @@ MagickPrivate void RelinquishOpenCLKernel(cl_kernel kernel)
 %                                                                             %
 %                                                                             %
 %                                                                             %
-%   R e q u e s t O p e n C L D e v i c e                                     %
++   R e q u e s t O p e n C L D e v i c e                                     %
 %                                                                             %
 %                                                                             %
 %                                                                             %
@@ -2609,7 +2893,7 @@ MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
 
   device=(MagickCLDevice) NULL;
   best_score=0.0;
-  LockSemaphoreInfo(clEnv->lock);
+  LockSemaphoreInfo(openCL_lock);
   for (i = 0; i < clEnv->number_devices; i++)
   {
     if (clEnv->devices[i]->enabled == MagickFalse)
@@ -2625,7 +2909,7 @@ MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
   }
   if (device != (MagickCLDevice)NULL)
     device->requested++;
-  UnlockSemaphoreInfo(clEnv->lock);
+  UnlockSemaphoreInfo(openCL_lock);
 
   return(device);
 }
index 005a399090aaafcf85d53d6ddcf243b19f557b16..2342ec840013d226c51c2df70a2e8f525faf7aef 100644 (file)
@@ -42,7 +42,6 @@ typedef struct _KernelProfileRecord
 }* KernelProfileRecord;
 
 typedef struct _MagickCLDevice* MagickCLDevice;
-typedef struct _MagickCLEnv* MagickCLEnv;
 
 extern MagickExport const char
   *GetOpenCLDeviceName(const MagickCLDevice),
index bf5c92043e3f9e2ca1b9fc17a15ec474fa50a289..53e7e4620e8f390ec48afb81694ed16d695fb57a 100644 (file)
@@ -2873,6 +2873,7 @@ MagickExport Image *ResizeImage(const Image *image,const size_t columns,
           ((x_factor*y_factor) > 1.0))
         filter_type=MitchellFilter;
   resize_filter=AcquireResizeFilter(image,filter_type,MagickFalse,exception);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   resize_image=AccelerateResizeImage(image,columns,rows,resize_filter,
     exception);
   if (resize_image != (Image *) NULL)
@@ -2880,6 +2881,7 @@ MagickExport Image *ResizeImage(const Image *image,const size_t columns,
       resize_filter=DestroyResizeFilter(resize_filter);
       return(resize_image);
     }
+#endif
   resize_image=CloneImage(image,columns,rows,MagickTrue,exception);
   if (resize_image == (Image *) NULL)
     {
index 901989a54902e2767c6f091d5462a00b17ff6a7a..4fb916aad19cfdf117ad4944daa72c9ecd4118fc 100644 (file)
@@ -1011,9 +1011,11 @@ MagickExport MagickBooleanType FunctionImage(Image *image,
     (void) LogMagickEvent(TraceEvent,GetMagickModule(),"%s",image->filename);
   assert(exception != (ExceptionInfo *) NULL);
   assert(exception->signature == MagickCoreSignature);
+#if defined(MAGICKCORE_OPENCL_SUPPORT)
   if (AccelerateFunctionImage(image,function,number_parameters,parameters,
         exception) != MagickFalse)
     return(MagickTrue);
+#endif
   if (SetImageStorageClass(image,DirectClass,exception) == MagickFalse)
     return(MagickFalse);
   status=MagickTrue;