diff --git a/filters/convolve.c b/filters/convolve.c
index 8005c74..415180f 100644
--- a/filters/convolve.c
+++ b/filters/convolve.c
@@ -81,45 +81,27 @@
 
 #if defined(MAGICKCORE_OPENCL_SUPPORT)
 
+#if defined(MAGICKCORE_HDRI_SUPPORT)
+#define CLOptions "-DMAGICKCORE_HDRI_SUPPORT=1 -DCLQuantumType=float " \
+  "-DCLPixelType=float4 -DQuantumRange=%g"
+#define CLPixelPacket  cl_ufloat4
+#else
 #if (MAGICKCORE_QUANTUM_DEPTH == 8)
-#if !defined(MAGICKCORE_HDRI_SUPPORT)
-#define CLQuantumString  "uchar"
-#define CLPixelPacketString  "uchar4"
+#define CLOptions \
+  "-DCLQuantumType=uchar -DCLPixelType=uchar4 -DQuantumRange=%g"
 #define CLPixelPacket  cl_uchar4
-#else
-#define CLQuantumString  "float"
-#define CLPixelPacketString  "float4"
-#define CLPixelPacket  cl_float4
-#endif
 #elif (MAGICKCORE_QUANTUM_DEPTH == 16)
-#if !defined(MAGICKCORE_HDRI_SUPPORT)
-#define CLQuantumString  "ushort"
-#define CLPixelPacketString  "ushort4"
+#define CLOptions \
+  "-DCLQuantumType=ushort -DCLPixelType=ushort4 -DQuantumRange=%g"
 #define CLPixelPacket  cl_ushort4
-#else
-#define CLQuantumString  "float"
-#define CLPixelPacketString  "float4"
-#define CLPixelPacket  cl_float4
-#endif
 #elif (MAGICKCORE_QUANTUM_DEPTH == 32)
-#if !defined(MAGICKCORE_HDRI_SUPPORT)
-#define CLQuantumString  "uint"
-#define CLPixelPacketString  "uint4"
+#define CLOptions \
+  "-DCLQuantumType=uint -DCLPixelType=uint4 -DQuantumRange=%g"
 #define CLPixelPacket  cl_uint4
-#else
-#define CLQuantumString  "float"
-#define CLPixelPacketString  "float4"
-#define CLPixelPacket  cl_float4
-#endif
-#else
-#if !defined(MAGICKCORE_HDRI_SUPPORT)
-#define CLQuantumString  "ulong"
-#define CLPixelPacketString  "ulong4"
+#elif (MAGICKCORE_QUANTUM_DEPTH == 32)
+#define CLOptions \
+  "-DCLQuantumType=ulong -DCLPixelType=ulong4 -DQuantumRange=%g"
 #define CLPixelPacket  cl_ulong4
-#else
-#define CLQuantumString  "float"
-#define CLPixelPacketString  "float4"
-#define CLPixelPacket  cl_float4
 #endif
 #endif
 
@@ -146,7 +128,9 @@
 
   cl_uint
     width,
-    height,
+    height;
+
+  cl_bool
     matte;
 
   cl_mem
@@ -155,65 +139,139 @@
 
 static char
   *convolve_program =
-    "#define QuantumScale  (1.0/QuantumRange)\n"
-    "\n"
-    "static uint AuthenticPixel(const int value,const uint range)\n"
+    "static long AuthenticPixel(const long offset,const ulong range)\n"
     "{\n"
-    "  if (value < 0)\n"
+    "  if (offset < 0)\n"
     "    return(0);\n"
-    "  if (value >= range)\n"
+    "  if (offset >= range)\n"
     "    return(range-1);\n"
-    "  return(value);\n"
+    "  return(offset);\n"
     "}\n"
     "\n"
-    "static " CLQuantumString " AuthenticQuantum(const float value)\n"
+    "static CLQuantumType AuthenticQuantum(const double value)\n"
     "{\n"
+    "#if !defined(MAGICKCORE_HDRI_SUPPORT)\n"
     "  if (value < 0)\n"
     "    return(0);\n"
     "  if (value >= QuantumRange)\n"
     "    return(QuantumRange);\n"
-    "  return(value+0.5);\n"
+    "  return((CLQuantumType) (value+0.5));\n"
+    "#else\n"
+    "  return((CLQuantumType) value)\n"
+    "#endif\n"
     "}\n"
     "\n"
-    "__kernel void Convolve(const __global " CLPixelPacketString " *input,\n"
-    "  __constant float *mask,const uint width,const uint height,\n"
-    "  const uint matte,__global " CLPixelPacketString " *output)\n"
+    "__kernel void Convolve(const __global CLPixelType *input,\n"
+    "  __constant double *mask,const uint width,const uint height,\n"
+    "  const bool matte,__global CLPixelType *output)\n"
     "{\n"
-    "  const uint columns = get_global_size(0);\n"
-    "  const uint rows = get_global_size(1);\n"
+    "  const ulong columns = get_global_size(0);\n"
+    "  const ulong rows = get_global_size(1);\n"
     "\n"
-    "  const int x = get_global_id(0);\n"
-    "  const int y = get_global_id(1);\n"
+    "  const long x = get_global_id(0);\n"
+    "  const long y = get_global_id(1);\n"
     "\n"
-    "  float4 sum = { 0.0, 0.0, 0.0, 0.0 };\n"
-    "  float gamma = 0.0;\n"
-    "  const int hstep = (width-1)/2;\n"
-    "  const int vstep = (height-1)/2;\n"
-    "  uint i = 0;\n"
+    "  double4 sum = { 0.0, 0.0, 0.0, 0.0 };\n"
+    "  double alpha = 0.0;\n"
+    "  double gamma = 0.0;\n"
+    "  double scale = (1.0/QuantumRange);\n"
+    "  const long mid_width = (width-1)/2;\n"
+    "  const long mid_height = (height-1)/2;\n"
+    "  ulong i = 0;\n"
+    "  ulong index = 0;\n"
     "\n"
-    "  for (int v=(-vstep); v <= vstep; v++)\n"
-    "  {\n"
-    "    for (int u=(-hstep); u <= vstep; u++)\n"
+    "  int state = 0;\n"
+    "  if (matte != false)\n"
+    "    state=1;\n"
+    "  if ((x >= width) && (x < (columns-width-1)) &&\n"
+    "      (y >= height) && (y < (rows-height-1)))\n"
     "    {\n"
-    "      const uint index = AuthenticPixel(y+v,rows)*columns+\n"
-    "        AuthenticPixel(x+u,columns);\n"
-    "      float alpha = 1.0;\n"
-    "      if (matte != 0)\n"
-    "        alpha = (float) (QuantumScale*(QuantumRange-input[index].w));\n"
-    "      sum.x+=alpha*mask[i]*input[index].x;\n"
-    "      sum.y+=alpha*mask[i]*input[index].y;\n"
-    "      sum.z+=alpha*mask[i]*input[index].z;\n"
-    "      sum.w+=mask[i]*input[index].w;\n"
-    "      gamma+=alpha*mask[i];\n"
-    "      i++;\n"
+    "      state=2;\n"
+    "      if (matte != false)\n"
+    "        state=3;\n"
     "    }\n"
-    "  }\n"
+    "  switch (state)\n"
+    "  {"
+    "    case 0:"
+    "    {"
+    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
+    "      {\n"
+    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
+    "        {\n"
+    "          index=AuthenticPixel(y+v,rows)*columns+\n"
+    "            AuthenticPixel(x+u,columns);\n"
+    "          sum.x+=mask[i]*input[index].x;\n"
+    "          sum.y+=mask[i]*input[index].y;\n"
+    "          sum.z+=mask[i]*input[index].z;\n"
+    "          gamma+=mask[i];\n"
+    "          i++;\n"
+    "        }\n"
+    "      }\n"
+    "      break;"
+    "    }"
+    "    case 1:"
+    "    {"
+    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
+    "      {\n"
+    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
+    "        {\n"
+    "          index=AuthenticPixel(y+v,rows)*columns+\n"
+    "            AuthenticPixel(x+u,columns);\n"
+    "          alpha=scale*(QuantumRange-input[index].w);\n"
+    "          sum.x+=alpha*mask[i]*input[index].x;\n"
+    "          sum.y+=alpha*mask[i]*input[index].y;\n"
+    "          sum.z+=alpha*mask[i]*input[index].z;\n"
+    "          sum.w+=mask[i]*input[index].w;\n"
+    "          gamma+=alpha*mask[i];\n"
+    "          i++;\n"
+    "        }\n"
+    "      }\n"
+    "      break;"
+    "    }"
+    "    case 2:"
+    "    {"
+    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
+    "      {\n"
+    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
+    "        {\n"
+    "          index=(y+v)*columns+(x+u);\n"
+    "          sum.x+=mask[i]*input[index].x;\n"
+    "          sum.y+=mask[i]*input[index].y;\n"
+    "          sum.z+=mask[i]*input[index].z;\n"
+    "          gamma+=mask[i];\n"
+    "          i++;\n"
+    "        }\n"
+    "      }\n"
+    "      break;"
+    "    }"
+    "    case 3:"
+    "    {"
+    "      for (long v=(-mid_height); v <= mid_height; v++)\n"
+    "      {\n"
+    "        for (long u=(-mid_width); u <= mid_width; u++)\n"
+    "        {\n"
+    "          index=(y+v)*columns+(x+u);\n"
+    "          alpha=scale*(QuantumRange-input[index].w);\n"
+    "          sum.x+=alpha*mask[i]*input[index].x;\n"
+    "          sum.y+=alpha*mask[i]*input[index].y;\n"
+    "          sum.z+=alpha*mask[i]*input[index].z;\n"
+    "          sum.w+=mask[i]*input[index].w;\n"
+    "          gamma+=alpha*mask[i];\n"
+    "          i++;\n"
+    "        }\n"
+    "      }\n"
+    "      break;"
+    "    }"
+    "  }"
     "  gamma=1.0/((gamma <= 0.000001) && (gamma >= -0.000001) ? 1.0 : gamma);\n"
-    "  const uint index = y*columns+x;\n"
+    "  index=y*columns+x;\n"
     "  output[index].x=AuthenticQuantum(gamma*sum.x);\n"
     "  output[index].y=AuthenticQuantum(gamma*sum.y);\n"
     "  output[index].z=AuthenticQuantum(gamma*sum.z);\n"
-    "  output[index].w=AuthenticQuantum(sum.w);\n"
+    "  if (matte == false)\n"
+    "    output[index].w=input[index].w;\n"
+    "  else\n"
+    "    output[index].w=AuthenticQuantum(sum.w);\n"
     "}\n";
 
 static void OpenCLNotify(const char *message,const void *data,size_t length,
@@ -230,8 +288,8 @@
 }
 
 static MagickBooleanType BindCLParameters(CLInfo *cl_info,Image *image,
-  void *pixels,float *mask,const unsigned long width,const unsigned long height,
-  void *convolve_pixels)
+  void *pixels,double *mask,const unsigned long width,
+  const unsigned long height,void *convolve_pixels)
 {
   cl_int
     status;
@@ -249,7 +307,7 @@
   if (status != CL_SUCCESS)
     return(MagickFalse);
   cl_info->mask=clCreateBuffer(cl_info->context,CL_MEM_READ_ONLY |
-    CL_MEM_USE_HOST_PTR,width*height*sizeof(cl_float),mask,&status);
+    CL_MEM_USE_HOST_PTR,width*height*sizeof(cl_double),mask,&status);
   if ((cl_info->mask == (cl_mem) NULL) || (status != CL_SUCCESS))
     return(MagickFalse);
   status=clSetKernelArg(cl_info->kernel,1,sizeof(cl_mem),(void *)
@@ -266,8 +324,8 @@
     &cl_info->height);
   if (status != CL_SUCCESS)
     return(MagickFalse);
-  cl_info->matte=(cl_uint) image->matte;
-  status=clSetKernelArg(cl_info->kernel,4,sizeof(cl_uint),(void *)
+  cl_info->matte=(cl_bool) image->matte;
+  status=clSetKernelArg(cl_info->kernel,4,sizeof(cl_bool),(void *)
     &cl_info->matte);
   if (status != CL_SUCCESS)
     return(MagickFalse);
@@ -420,8 +478,8 @@
       DestroyCLInfo(cl_info);
       return((CLInfo *) NULL);
     }
-  (void) FormatMagickString(options,MaxTextExtent,"-DQuantumRange=%g",
-    (double) QuantumRange);
+  (void) FormatMagickString(options,MaxTextExtent,CLOptions,(double)
+    QuantumRange);
   status=clBuildProgram(cl_info->program,1,cl_info->devices,options,NULL,NULL);
   if ((cl_info->program == (cl_program) NULL) || (status != CL_SUCCESS))
     {
@@ -456,7 +514,7 @@
   return(cl_info);
 }
 
-static float *ParseMask(const char *value,unsigned long *order)
+static double *ParseMask(const char *value,unsigned long *order)
 {
   char
     token[MaxTextExtent];
@@ -464,7 +522,7 @@
   const char
     *p;
 
-  float
+  double
     *mask,
     normalize;
 
@@ -484,8 +542,8 @@
       GetMagickToken(p,&p,token);
   }
   *order=(unsigned long) sqrt((double) i+1.0);
-  mask=(float *) AcquireQuantumMemory(*order,*order*sizeof(*mask));
-  if (mask == (float *) NULL)
+  mask=(double *) AcquireQuantumMemory(*order,*order*sizeof(*mask));
+  if (mask == (double *) NULL)
     return(mask);
   p=(const char *) value;
   if (*p == '\'')
@@ -523,7 +581,7 @@
     "DelegateLibrarySupportNotBuiltIn","`%s' (OpenCL)",(*images)->filename);
 #else
   {
-    float
+    double
       *mask;
 
     Image
@@ -544,13 +602,13 @@
       Convolve image.
     */
     mask=ParseMask(argv[0],&order);
-    if (mask == (float *) NULL)
+    if (mask == (double *) NULL)
       (void) ThrowMagickException(exception,GetMagickModule(),
         ResourceLimitError,"MemoryAllocationFailed","`%s'",(*images)->filename);
     cl_info=GetCLInfo(*images,"Convolve",convolve_program,exception);
     if (cl_info == (CLInfo *) NULL)
       {
-        mask=(float *) RelinquishMagickMemory(mask);
+        mask=(double *) RelinquishMagickMemory(mask);
         return(MagickImageFilterSignature);
       }
     image=(*images);
@@ -596,7 +654,7 @@
       DestroyCLBuffers(cl_info);
       convolve_image=DestroyImage(convolve_image);
     }
-    mask=(float *) RelinquishMagickMemory(mask);
+    mask=(double *) RelinquishMagickMemory(mask);
     cl_info=DestroyCLInfo(cl_info);
   }
 #endif
diff --git a/magick/image-private.h b/magick/image-private.h
index f79c415..92a6b6a 100644
--- a/magick/image-private.h
+++ b/magick/image-private.h
@@ -26,6 +26,7 @@
 #define Magick2PI  6.28318530717958647692528676655900576839433879875020
 #define MagickPI2  1.57079632679489661923132169163975144209858469968755
 #define MagickSQ1_2  0.7071067811865475244008443621048490
+#define MagickSQ2  1.414213562373095
 #define MagickSQ2PI  2.50662827463100024161235523934010416269302368164062
 #define QuantumScale  ((double) 1.0/(double) QuantumRange)
 #define UndefinedTicksPerSecond  100L
diff --git a/magick/morphology.c b/magick/morphology.c
index 54a4e68..634c274 100644
--- a/magick/morphology.c
+++ b/magick/morphology.c
@@ -59,6 +59,7 @@
 #include "magick/gem.h"
 #include "magick/hashmap.h"
 #include "magick/image.h"
+#include "magick/image-private.h"
 #include "magick/list.h"
 #include "magick/memory_.h"
 #include "magick/monitor-private.h"
@@ -1126,6 +1127,7 @@
       if (image->colorspace == CMYKColorspace)
         q_indexes[x] = p_indexes[r];
 
+      result.index=0;
       switch (method) {
         case ConvolveMorphology:
           result=bias;
@@ -1513,6 +1515,7 @@
     limit = image->columns > image->rows ? image->columns : image->rows;
 
   /* Special morphology cases */
+  changed=MagickFalse;
   switch( method ) {
     case CloseMorphology:
       new_image = MorphologyImage(image, DialateMorphology, iterations, channel,
diff --git a/magick/option.h b/magick/option.h
index 72e11c9..0a42ffa 100644
--- a/magick/option.h
+++ b/magick/option.h
@@ -88,7 +88,9 @@
   MagickThresholdOptions,
   MagickTypeOptions,
   MagickValidateOptions,
-  MagickVirtualPixelOptions
+  MagickVirtualPixelOptions,
+  MagickKernelOptions,
+  MagickMorphologyOptions
 } MagickOption;
 
 typedef enum