From 7d42b3c11794313e25ede648463812545161de2d Mon Sep 17 00:00:00 2001
From: dirk <dirk@git.imagemagick.org>
Date: Mon, 18 Apr 2016 23:12:54 +0200
Subject: [PATCH] Changed recording of kernel profiles. Some const fixes.

---
 MagickCore/accelerate.c     |  62 +++++-----
 MagickCore/opencl-private.h |  59 ++-------
 MagickCore/opencl.c         | 236 +++++++++++++++++++++++++++++-------
 MagickCore/opencl.h         |  25 +++-
 4 files changed, 258 insertions(+), 124 deletions(-)

diff --git a/MagickCore/accelerate.c b/MagickCore/accelerate.c
index f9c9a50f5..de9be6c2a 100644
--- a/MagickCore/accelerate.c
+++ b/MagickCore/accelerate.c
@@ -472,7 +472,7 @@ static MagickBooleanType LaunchHistogramKernel(MagickCLEnv clEnv,
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,HistogramKernel,event);
+  RecordProfileData(device,histogramKernel,event);
 
   outputReady = MagickTrue;
 
@@ -671,7 +671,7 @@ static Image *ComputeAddNoiseImage(const Image *image,MagickCLEnv clEnv,
     goto cleanup;
   }
 
-  RecordProfileData(device,AddNoiseKernel,event);
+  RecordProfileData(device,addNoiseKernel,event);
   if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
     goto cleanup;
 
@@ -889,7 +889,7 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv,
           (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
           goto cleanup;
         }
-        RecordProfileData(device,BlurRowKernel,event);
+        RecordProfileData(device,blurRowKernel,event);
       }
     }
 
@@ -932,7 +932,7 @@ static Image *ComputeBlurImage(const Image* image,MagickCLEnv clEnv,
           (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
           goto cleanup;
         }
-        RecordProfileData(device,BlurColumnKernel,event);
+        RecordProfileData(device,blurColumnKernel,event);
       }
     }
 
@@ -1052,7 +1052,7 @@ static MagickBooleanType LaunchCompositeKernel(MagickCLEnv clEnv,
   clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, compositeKernel, 2, NULL, 
     global_work_size, local_work_size, 0, NULL, &event);
 
-  RecordProfileData(device,CompositeKernel,event);
+  RecordProfileData(device,compositeKernel,event);
 
   RelinquishOpenCLKernel(compositeKernel);
 
@@ -1366,7 +1366,7 @@ static MagickBooleanType ComputeContrastImage(Image *image,MagickCLEnv clEnv,
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,ContrastKernel,event);
+  RecordProfileData(device,filterKernel,event);
 
   if (ALIGNED(inputPixels,CLPixelPacket)) 
   {
@@ -1947,7 +1947,7 @@ static MagickBooleanType ComputeContrastStretchImage(Image *image,
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,ContrastStretchKernel,event);
+  RecordProfileData(device,stretchKernel,event);
 
   /* read the data back */
   if (ALIGNED(inputPixels,CLPixelPacket)) 
@@ -2265,7 +2265,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,ConvolveOptimizedKernel,event);
+    RecordProfileData(device,clkernel,event);
   }
   else
   {
@@ -2311,7 +2311,7 @@ static Image *ComputeConvolveImage(const Image* image,MagickCLEnv clEnv,
       goto cleanup;
     }
   }
-  RecordProfileData(device,ConvolveKernel,event);
+  RecordProfileData(device,clkernel,event);
 
   if (ALIGNED(filteredPixels,CLPixelPacket)) 
   {
@@ -2609,7 +2609,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass1Kernel,event);
+    RecordProfileData(device,hullPass1,event);
 
     /* launch the kernel */
 	clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
@@ -2618,7 +2618,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass2Kernel,event);
+    RecordProfileData(device,hullPass2,event);
 
     if (k == 0)
       clStatus =clEnv->library->clSetKernelArg(hullPass1,0,sizeof(cl_mem),(void *)(tempImageBuffer));
@@ -2641,7 +2641,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass1Kernel,event);
+    RecordProfileData(device,hullPass1,event);
 
     /* launch the kernel */
 	clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
@@ -2650,7 +2650,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass2Kernel,event);
+    RecordProfileData(device,hullPass2,event);
 
     offset.s[0] = -X[k];
     offset.s[1] = -Y[k];
@@ -2671,7 +2671,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass1Kernel,event);
+    RecordProfileData(device,hullPass1,event);
 
     /* launch the kernel */
 	clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
@@ -2680,7 +2680,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass2Kernel,event);
+    RecordProfileData(device,hullPass2,event);
 
     offset.s[0] = X[k];
     offset.s[1] = Y[k];
@@ -2705,7 +2705,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass1Kernel,event);
+    RecordProfileData(device,hullPass1,event);
 
     /* launch the kernel */
 	clStatus = clEnv->library->clEnqueueNDRangeKernel(queue, hullPass2, 2, NULL, global_work_size, NULL, 0, NULL, &event);
@@ -2714,7 +2714,7 @@ static Image *ComputeDespeckleImage(const Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,HullPass2Kernel,event);
+    RecordProfileData(device,hullPass2,event);
   }
 
   if (ALIGNED(filteredPixels,CLPixelPacket)) 
@@ -3167,7 +3167,7 @@ static MagickBooleanType ComputeEqualizeImage(Image *image,MagickCLEnv clEnv,
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,EqualizeKernel,event);
+  RecordProfileData(device,equalizeKernel,event);
 
   /* read the data back */
   if (ALIGNED(inputPixels,CLPixelPacket)) 
@@ -3374,7 +3374,7 @@ static MagickBooleanType ComputeFunctionImage(Image *image,MagickCLEnv clEnv,
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,ComputeFunctionKernel,event);
+  RecordProfileData(device,clkernel,event);
 
   if (copyWriteBuffer(image,clEnv,device,queue,imageBuffer,pixels,exception) == MagickFalse)
       goto cleanup;
@@ -3526,7 +3526,7 @@ static MagickBooleanType ComputeGrayscaleImage(Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,GrayScaleKernel,event);
+    RecordProfileData(device,grayscaleKernel,event);
   }
 
   if (copyWriteBuffer(image,clEnv,device,queue,imageBuffer,inputPixels,exception) == MagickFalse)
@@ -3804,7 +3804,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
           (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
           goto cleanup;
         }
-        RecordProfileData(device,LocalContrastBlurRowKernel,event);
+        RecordProfileData(device,blurRowKernel,event);
       }
     }
 
@@ -3847,7 +3847,7 @@ static Image *ComputeLocalContrastImage(const Image *image,MagickCLEnv clEnv,
           (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
           goto cleanup;
         }
-        RecordProfileData(device,LocalContrastBlurApplyColumnKernel,event);
+        RecordProfileData(device,blurColumnKernel,event);
       }
     }
   }
@@ -4063,7 +4063,7 @@ static MagickBooleanType ComputeModulateImage(Image *image,MagickCLEnv clEnv,
       (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
       goto cleanup;
     }
-    RecordProfileData(device,ModulateKernel,event);
+    RecordProfileData(device,modulateKernel,event);
   }
 
   if (ALIGNED(inputPixels,CLPixelPacket)) 
@@ -4421,7 +4421,7 @@ static Image* ComputeMotionBlurImage(const Image *image,MagickCLEnv clEnv,
       "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,MotionBlurKernel,event);
+  RecordProfileData(device,motionBlurKernel,event);
 
   if (ALIGNED(filteredPixels,CLPixelPacket)) 
   {
@@ -4686,7 +4686,7 @@ RestoreMSCWarning
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,ResizeHorizontalKernel,event);
+  RecordProfileData(device,horizontalKernel,event);
   status = MagickTrue;
 
 
@@ -4882,7 +4882,7 @@ RestoreMSCWarning
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,ResizeVerticalKernel,event);
+  RecordProfileData(device,verticalKernel,event);
   status = MagickTrue;
 
 
@@ -5320,7 +5320,7 @@ static Image* ComputeRotationalBlurImage(const Image *image,MagickCLEnv clEnv,
     (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
     goto cleanup;
   }
-  RecordProfileData(device,RotationalBlurKernel,event);
+  RecordProfileData(device,rotationalBlurKernel,event);
 
   if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
   {
@@ -5553,7 +5553,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv,
         (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
         goto cleanup;
       }
-      RecordProfileData(device,BlurRowKernel,event);
+      RecordProfileData(device,blurRowKernel,event);
     }
 
 
@@ -5600,7 +5600,7 @@ static Image *ComputeUnsharpMaskImage(const Image *image,MagickCLEnv clEnv,
         (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
         goto cleanup;
       }
-      RecordProfileData(device,UnsharpMaskBlurColumnKernel,event);
+      RecordProfileData(device,unsharpMaskBlurColumnKernel,event);
     }
 
   }
@@ -5771,7 +5771,7 @@ static Image *ComputeUnsharpMaskImageSingle(const Image *image,
         (void) OpenCLThrowMagickException(device,exception, GetMagickModule(), ResourceLimitWarning, "clEnv->library->clEnqueueNDRangeKernel failed.", "'%s'", ".");
         goto cleanup;
       }
-      RecordProfileData(device,UnsharpMaskKernel,event);
+      RecordProfileData(device,unsharpMaskKernel,event);
     }
   }
 
@@ -5955,7 +5955,7 @@ static Image *ComputeWaveletDenoiseImage(const Image *image,MagickCLEnv clEnv,
         goto cleanup;
       }
     }
-    RecordProfileData(device,WaveletDenoiseKernel,event);
+    RecordProfileData(device,denoiseKernel,event);
   }
 
   if (copyWriteBuffer(filteredImage,clEnv,device,queue,filteredImageBuffer,filteredPixels,exception) == MagickFalse)
diff --git a/MagickCore/opencl-private.h b/MagickCore/opencl-private.h
index 58ec7c767..06b153d3b 100644
--- a/MagickCore/opencl-private.h
+++ b/MagickCore/opencl-private.h
@@ -29,35 +29,6 @@ Include declarations.
 extern "C" {
 #endif
 
-typedef enum
-{
-  AddNoiseKernel,
-  BlurColumnKernel,
-  BlurRowKernel,
-  CompositeKernel,
-  ContrastKernel,
-  ContrastStretchKernel,
-  ConvolveKernel,
-  ConvolveOptimizedKernel,
-  ComputeFunctionKernel,
-  EqualizeKernel,
-  GrayScaleKernel,
-  HistogramKernel,
-  HullPass1Kernel,
-  HullPass2Kernel,
-  LocalContrastBlurApplyColumnKernel,
-  LocalContrastBlurRowKernel,
-  ModulateKernel,
-  MotionBlurKernel,
-  ResizeHorizontalKernel,
-  ResizeVerticalKernel,
-  RotationalBlurKernel,
-  UnsharpMaskKernel,
-  UnsharpMaskBlurColumnKernel,
-  WaveletDenoiseKernel,
-  KERNEL_COUNT
-} ProfiledKernels;
-
 #if !defined(MAGICKCORE_OPENCL_SUPPORT)
   typedef void* cl_context;
   typedef void* cl_command_queue;
@@ -73,19 +44,8 @@ typedef enum
   Define declarations.
 */
 #define MAGICKCORE_OPENCL_UNDEFINED_SCORE -1.0
-#define MAGICKCORE_OPENCL_PROFILE_KERNELS 0
 #define MAGICKCORE_OPENCL_COMMAND_QUEUES 16
 
-#if MAGICKCORE_OPENCL_PROFILE_KERNELS
-typedef struct
-{
-  cl_ulong min;
-  cl_ulong max;
-  cl_ulong total;
-  cl_ulong count;
-} KernelProfileRecord;
-#endif
-
 /* Platform APIs */
 typedef CL_API_ENTRY cl_int
   (CL_API_CALL *MAGICKpfn_clGetPlatformIDs)(cl_uint num_entries,
@@ -190,6 +150,11 @@ typedef CL_API_ENTRY cl_int
   (CL_API_CALL *MAGICKpfn_clSetKernelArg)(cl_kernel kernel,cl_uint arg_index,
   size_t arg_size,const void * arg_value) CL_API_SUFFIX__VERSION_1_0;
 
+typedef CL_API_ENTRY cl_int
+  (CL_API_CALL *MAGICKpfn_clGetKernelInfo)(cl_kernel kernel,
+    cl_kernel_info param_name,size_t param_value_size,void *param_value,
+    size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+
 
 /* Enqueued Commands APIs */
 typedef CL_API_ENTRY cl_int
@@ -268,6 +233,7 @@ struct MagickLibraryRec
   MAGICKpfn_clCreateKernel            clCreateKernel;
   MAGICKpfn_clReleaseKernel           clReleaseKernel;
   MAGICKpfn_clSetKernelArg            clSetKernelArg;
+  MAGICKpfn_clGetKernelInfo           clGetKernelInfo;
 
   MAGICKpfn_clEnqueueReadBuffer       clEnqueueReadBuffer;
   MAGICKpfn_clEnqueueMapBuffer        clEnqueueMapBuffer;
@@ -315,8 +281,12 @@ struct _MagickCLDevice
   double
     score;
 
+  KernelProfileRecord
+    *profile_records;
+
   MagickBooleanType
-    enabled;
+    enabled,
+    profile_kernels;
 
   SemaphoreInfo
     *lock;
@@ -324,11 +294,6 @@ struct _MagickCLDevice
   ssize_t
     command_queues_index,
     created_queues;
-
-#if MAGICKCORE_OPENCL_PROFILE_KERNELS
-  KernelProfileRecord
-    profileRecords[KERNEL_COUNT];
-#endif
 };
 
 struct _MagickCLEnv
@@ -424,7 +389,7 @@ extern MagickPrivate unsigned long
 extern MagickPrivate void
   DumpOpenCLProfileData(),
   OpenCLTerminus(),
-  RecordProfileData(MagickCLDevice,ProfiledKernels,cl_event),
+  RecordProfileData(MagickCLDevice,cl_kernel,cl_event),
   RelinquishOpenCLCommandQueue(MagickCLDevice,cl_command_queue),
   RelinquishOpenCLKernel(cl_kernel);
 
diff --git a/MagickCore/opencl.c b/MagickCore/opencl.c
index 0152259cc..5d1417f39 100644
--- a/MagickCore/opencl.c
+++ b/MagickCore/opencl.c
@@ -577,9 +577,8 @@ MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
   {
     UnlockSemaphoreInfo(device->lock);
     properties=(cl_command_queue_properties) NULL;
-#if MAGICKCORE_OPENCL_PROFILE_KERNELS
-    properties=CL_QUEUE_PROFILING_ENABLE;
-#endif
+    if (device->profile_kernels != MagickFalse)
+      properties=CL_QUEUE_PROFILING_ENABLE;
     queue=openCL_library->clCreateCommandQueue(device->context,
       device->deviceID,properties,NULL);
   }
@@ -1324,7 +1323,6 @@ static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
 
 MagickPrivate void DumpOpenCLProfileData()
 {
-#if MAGICKCORE_OPENCL_PROFILE_KERNELS
 #define OpenCLLog(message) \
    fwrite(message,sizeof(char),strlen(message),log); \
    fwrite("\n",sizeof(char),1,log);
@@ -1346,6 +1344,12 @@ MagickPrivate void DumpOpenCLProfileData()
 
   clEnv=GetCurrentOpenCLEnv();
 
+  for (i = 0; i < clEnv->number_devices; i++)
+    if (clEnv->devices[i]->profile_kernels != MagickFalse)
+      break;
+  if (i == clEnv->number_devices)
+    return;
+
   (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
     GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
 
@@ -1357,33 +1361,36 @@ MagickPrivate void DumpOpenCLProfileData()
       device;
 
     device=clEnv->devices[i];
+    if ((device->profile_kernels == MagickFalse) ||
+        (device->profile_records == (KernelProfileRecord *) NULL))
+      continue;
+
     OpenCLLog("====================================================");
     fprintf(log,"Device:  %s\n",device->name);
     fprintf(log,"Version: %s\n",device->version);
     OpenCLLog("====================================================");
     OpenCLLog("                     average   calls     min     max");
     OpenCLLog("                     -------   -----     ---     ---");
-    for (j = 0; j < KERNEL_COUNT; j++)
+    j=0;
+    while (device->profile_records[j] != (KernelProfileRecord) NULL)
     {
       KernelProfileRecord
         profile;
 
-      profile=device->profileRecords[j];
-      if (profile.count == 0)
-        continue;
+      profile=device->profile_records[j];
       strcpy(indent,"                    ");
       strncpy(indent,kernelNames[j],min(strlen(kernelNames[j]),
         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);
+      sprintf(buf,"%s %7d %7d %7d %7d",indent,(int) (profile->total/
+        profile->count),(int) profile->count,(int) profile->min,
+        (int) profile->max);
       OpenCLLog(buf);
+      j++;
     }
     OpenCLLog("====================================================");
     fwrite("\n\n",sizeof(char),2,log);
   }
   fclose(log);
-#endif
 }
 
 /*
@@ -1599,25 +1606,31 @@ MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
 %
 %  The format of the GetOpenCLDevices method is:
 %
-%      MagickBooleanType GetOpenCLDevices()
+%      const MagickCLDevice *GetOpenCLDevices(size_t *length,
+%        ExceptionInfo *exception)
 %
 %  A description of each parameter follows:
 %
-%    o device: the OpenCL device.
+%    o length: the number of device.
+%
+%    o exception: return any errors or warnings in this structure.
+%
 */
 
-MagickExport const MagickCLDevice *GetOpenCLDevices(size_t *length)
+MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
+  ExceptionInfo *exception)
 {
   MagickCLEnv
     clEnv;
 
   clEnv=GetCurrentOpenCLEnv();
   if (clEnv == (MagickCLEnv) NULL)
-  {
-    if (length != (size_t *) NULL)
-      *length=0;
-    return((MagickCLDevice *) NULL);
-  }
+    {
+      if (length != (size_t *) NULL)
+        *length=0;
+      return((MagickCLDevice *) NULL);
+    }
+  InitializeOpenCL(clEnv,exception);
   if (length != (size_t *) NULL)
     *length=clEnv->number_devices;
   return(clEnv->devices);
@@ -1672,7 +1685,7 @@ MagickExport MagickCLDeviceType GetOpenCLDeviceType(
 %
 %  The format of the GetOpenCLDeviceName method is:
 %
-%      MagickBooleanType GetOpenCLDeviceVersion(MagickCLDevice device)
+%      const char *GetOpenCLDeviceVersion(MagickCLDevice device)
 %
 %  A description of each parameter follows:
 %
@@ -1716,6 +1729,48 @@ MagickExport MagickBooleanType GetOpenCLEnabled(void)
   return(clEnv->enabled);
 }
 
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%   G e t O p e n C L K e r n e l P r o f i l e R e c o r d s                 %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  GetOpenCLKernelProfileRecords() returns the profile records for the
+%  specified device and sets length to the number of profile records.
+%
+%  The format of the GetOpenCLKernelProfileRecords method is:
+%
+%      const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
+%
+%  A description of each parameter follows:
+%
+%    o length: the number of profiles records.
+*/
+
+MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
+  const MagickCLDevice device,size_t *length)
+{
+  if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
+      (KernelProfileRecord *) NULL))
+  {
+    if (length != (size_t *) NULL)
+      *length=0;
+    return((const KernelProfileRecord *) NULL);
+  }
+  if (length != (size_t *) NULL)
+    {
+      length=0;
+      while (device->profile_records[*length] != (KernelProfileRecord) NULL)
+        *length=*length+1;
+    }
+  return(device->profile_records);
+}
+
 /*
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 %                                                                             %
@@ -2077,6 +2132,7 @@ static MagickBooleanType BindOpenCLFunctions(MagickLibrary *openCL_library)
   BIND(clCreateKernel);
   BIND(clReleaseKernel);
   BIND(clSetKernelArg);
+  BIND(clGetKernelInfo);
 
   BIND(clEnqueueReadBuffer);
   BIND(clEnqueueMapBuffer);
@@ -2252,16 +2308,16 @@ MagickPrivate MagickBooleanType OpenCLThrowMagickException(
 %
 %    o device: the OpenCL device that did the operation.
 %
-%    o kernel: the kernel that was executed.
-%
 %    o event: the event that contains the profiling data.
 %
 */
 
 MagickPrivate void RecordProfileData(MagickCLDevice device,
-  ProfiledKernels kernel,cl_event event)
+  cl_kernel kernel,cl_event event)
 {
-#if MAGICKCORE_OPENCL_PROFILE_KERNELS
+  char
+    *name;
+
   cl_int
     status;
 
@@ -2270,31 +2326,75 @@ MagickPrivate void RecordProfileData(MagickCLDevice device,
     end,
     start;
 
+  KernelProfileRecord
+    profile_record;
+
+  size_t
+    i,
+    length;
+
+  if (device->profile_kernels == MagickFalse)
+    {
+      openCL_library->clReleaseEvent(event);
+      return;
+    }
+  status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
+    &length);
+  if (status != CL_SUCCESS)
+    {
+      openCL_library->clReleaseEvent(event);
+      return;
+    }
+  name=AcquireQuantumMemory(length,sizeof(*name));
+  (void) openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
+    name,NULL);
   start=end=elapsed=0;
   openCL_library->clWaitForEvents(1,&event);
   status=openCL_library->clGetEventProfilingInfo(event,
     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)
-    return;
+    {
+      name=DestroyString(name);
+      return;
+    }
   start/=1000; // usecs
   end/=1000;   // usecs
   elapsed=end-start;
   LockSemaphoreInfo(device->lock);
-  if ((elapsed < device->profileRecords[kernel].min) ||
-      (device->profileRecords[kernel].count == 0))
-    device->profileRecords[kernel].min=elapsed;
-  if (elapsed > device->profileRecords[kernel].max)
-    device->profileRecords[kernel].max = elapsed;
-  device->profileRecords[kernel].total += elapsed;
-  device->profileRecords[kernel].count += 1;
+  i=0;
+  profile_record=(KernelProfileRecord) NULL;
+  if (device->profile_records != (KernelProfileRecord *) NULL)
+    {
+      while (device->profile_records[i] != ((KernelProfileRecord) NULL))
+      {
+        if (LocaleCompare(device->profile_records[i]->kernel_name,name))
+          {
+            profile_record=device->profile_records[i];
+            break;
+          }
+        i++;
+      }
+    }
+  if (profile_record == ((KernelProfileRecord) NULL))
+    {
+      profile_record=AcquireMagickMemory(sizeof(*profile_record));
+      (void) ResetMagickMemory(profile_record,0,sizeof(*profile_record));
+      profile_record->kernel_name=AcquireString(name);
+      device->profile_records=ResizeMagickMemory(device->profile_records,i+2);
+      device->profile_records[i]=profile_record;
+      device->profile_records[i+1]=(KernelProfileRecord) NULL;
+    }
+  if ((elapsed < profile_record->min) || (profile_record->count == 0))
+    profile_record->min=elapsed;
+  if (elapsed > profile_record->max)
+    profile_record->max=elapsed;
+  profile_record->total+=elapsed;
+  profile_record->count+=1;
   UnlockSemaphoreInfo(device->lock);
-#else
-  magick_unreferenced(device);
-  magick_unreferenced(kernel);
-#endif
-  openCL_library->clReleaseEvent(event);
+  name=DestroyString(name);
 }
 
 /*
@@ -2459,7 +2559,7 @@ MagickPrivate void RelinquishOpenCLKernel(cl_kernel kernel)
 %
 %  The format of the SetOpenCLDeviceEnabled method is:
 %
-%      void SetOpenCLDeviceEnabled(const MagickCLDevice device,
+%      void SetOpenCLDeviceEnabled(MagickCLDevice device,
 %        MagickBooleanType value)
 %
 %  A description of each parameter follows:
@@ -2469,7 +2569,7 @@ MagickPrivate void RelinquishOpenCLKernel(cl_kernel kernel)
 %    o value: determines if the device should be enabled or disabled.
 */
 
-MagickExport void SetOpenCLDeviceEnabled(const MagickCLDevice device,
+MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
   const MagickBooleanType value)
 {
   if (device == (MagickCLDevice) NULL)
@@ -2477,6 +2577,41 @@ MagickExport void SetOpenCLDeviceEnabled(const MagickCLDevice device,
   device->enabled=value;
 }
 
+/*
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%   S e t O p e n C L K e r n e l P r o f i l e E n a b l e d                 %
+%                                                                             %
+%                                                                             %
+%                                                                             %
+%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
+%
+%  SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
+%  kernel profiling of a device.
+%
+%  The format of the SetOpenCLKernelProfileEnabled method is:
+%
+%      void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
+%        MagickBooleanType value)
+%
+%  A description of each parameter follows:
+%
+%    o device: the OpenCL device.
+%
+%    o value: determines if kernel profiling for the device should be enabled
+%             or disabled.
+*/
+
+MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
+  const MagickBooleanType value)
+{
+  if (device == (MagickCLDevice) NULL)
+    return;
+  device->profile_kernels=value;
+}
+
 /*
 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 %                                                                             %
@@ -2534,8 +2669,10 @@ MagickExport const char *GetOpenCLDeviceName(
   return((const char *) NULL);
 }
 
-MagickExport const MagickCLDevice *GetOpenCLDevices(size_t *length)
+MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
+  ExceptionInfo *magick_unused(exception))
 {
+  magick_unreferenced(exception);
   if (length != (size_t *) NULL)
     *length=0;
   return((MagickCLDevice *) NULL);
@@ -2548,6 +2685,14 @@ MagickExport MagickCLDeviceType GetOpenCLDeviceType(
   return(UndefinedCLDeviceType);
 }
 
+MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
+  size_t *length)
+{
+  if (length != (size_t *) NULL)
+    *length=0;
+  return((MagickCLDevice *) NULL);
+}
+
 MagickExport const char *GetOpenCLDeviceVersion(
   const MagickCLDevice magick_unused(device))
 {
@@ -2561,7 +2706,7 @@ MagickExport MagickBooleanType GetOpenCLEnabled(void)
 }
 
 MagickExport void SetOpenCLDeviceEnabled(
-  const MagickCLDevice magick_unused(device),
+  MagickCLDevice magick_unused(device),
   const MagickBooleanType magick_unused(value))
 {
   magick_unreferenced(device);
@@ -2575,4 +2720,11 @@ MagickExport MagickBooleanType SetOpenCLEnabled(
   return(MagickFalse);
 }
 
+MagickExport SetOpenCLKernelProfileEnabled(
+  MagickCLDevice magick_unused(device),
+  const MagickBooleanType magick_unused(value))
+{
+  magick_unreferenced(device);
+  magick_unreferenced(value);
+}
 #endif
\ No newline at end of file
diff --git a/MagickCore/opencl.h b/MagickCore/opencl.h
index 56d75e2b9..40a70ad17 100644
--- a/MagickCore/opencl.h
+++ b/MagickCore/opencl.h
@@ -29,6 +29,19 @@ typedef enum
   GpuCLDeviceType
 } MagickCLDeviceType;
 
+struct _KernelProfileRecord
+{
+  char
+    *kernel_name;
+
+  unsigned long
+    count,
+    max,
+    min,
+    total;
+};
+
+typedef struct _KernelProfileRecord* KernelProfileRecord;
 typedef struct _MagickCLDevice* MagickCLDevice;
 typedef struct _MagickCLEnv* MagickCLEnv;
 
@@ -36,12 +49,15 @@ extern MagickExport const char
   *GetOpenCLDeviceName(const MagickCLDevice),
   *GetOpenCLDeviceVersion(const MagickCLDevice);
 
-extern MagickExport const MagickCLDevice
-  *GetOpenCLDevices(size_t *);
+extern MagickExport const KernelProfileRecord
+  *GetOpenCLKernelProfilesRecords(const MagickCLDevice,size_t *);
 
 extern MagickExport double
   GetOpenCLDeviceBenchmarkScore(const MagickCLDevice);
 
+extern MagickExport MagickCLDevice
+  *GetOpenCLDevices(size_t *,ExceptionInfo *);
+
 extern MagickExport MagickCLDeviceType
   GetOpenCLDeviceType(const MagickCLDevice);
 
@@ -51,10 +67,11 @@ extern MagickExport MagickBooleanType
   SetOpenCLEnabled(const MagickBooleanType);
 
 extern MagickExport void
-  SetOpenCLDeviceEnabled(const MagickCLDevice,
+  SetOpenCLDeviceEnabled(MagickCLDevice,
+    const MagickBooleanType),
+  SetOpenCLKernelProfileEnabled(MagickCLDevice,
     const MagickBooleanType);
 
-
 #if defined(__cplusplus) || defined(c_plusplus)
 }
 #endif
-- 
2.40.0