MagickCore/opencl.c

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%                   OOO   PPPP   EEEEE  N   N   CCCC  L                       %
%                  O   O  P   P  E      NN  N  C      L                       %
%                  O   O  PPPP   EEE    N N N  C      L                       %
%                  O   O  P      E      N  NN  C      L                       %
%                   OOO   P      EEEEE  N   N   CCCC  LLLLL                   %
%                                                                             %
%                                                                             %
%                         MagickCore OpenCL Methods                           %
%                                                                             %
%                              Software Design                                %
%                                   Cristy                                    %
%                                 March 2000                                  %
%                                                                             %
%                                                                             %
%  Copyright 1999-2021 ImageMagick Studio LLC, a non-profit organization      %
%  dedicated to making software imaging solutions freely available.           %
%                                                                             %
%  You may not use this file except in compliance with the License.  You may  %
%  obtain a copy of the License at                                            %
%                                                                             %
%    https://imagemagick.org/script/license.php                               %
%                                                                             %
%  Unless required by applicable law or agreed to in writing, software        %
%  distributed under the License is distributed on an "AS IS" BASIS,          %
%  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.   %
%  See the License for the specific language governing permissions and        %
%  limitations under the License.                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%
%
*/

/*
  Include declarations.
*/
#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"
#include "MagickCore/configure.h"
#include "MagickCore/distort.h"
#include "MagickCore/draw.h"
#include "MagickCore/effect.h"
#include "MagickCore/exception.h"
#include "MagickCore/exception-private.h"
#include "MagickCore/fx.h"
#include "MagickCore/gem.h"
#include "MagickCore/geometry.h"
#include "MagickCore/image.h"
#include "MagickCore/image-private.h"
#include "MagickCore/layer.h"
#include "MagickCore/mime-private.h"
#include "MagickCore/memory_.h"
#include "MagickCore/memory-private.h"
#include "MagickCore/monitor.h"
#include "MagickCore/montage.h"
#include "MagickCore/morphology.h"
#include "MagickCore/nt-base.h"
#include "MagickCore/nt-base-private.h"
#include "MagickCore/opencl.h"
#include "MagickCore/opencl-private.h"
#include "MagickCore/option.h"
#include "MagickCore/policy.h"
#include "MagickCore/property.h"
#include "MagickCore/quantize.h"
#include "MagickCore/quantum.h"
#include "MagickCore/random_.h"
#include "MagickCore/random-private.h"
#include "MagickCore/resample.h"
#include "MagickCore/resource_.h"
#include "MagickCore/splay-tree.h"
#include "MagickCore/semaphore.h"
#include "MagickCore/statistic.h"
#include "MagickCore/string_.h"
#include "MagickCore/string-private.h"
#include "MagickCore/token.h"
#include "MagickCore/utility.h"
#include "MagickCore/utility-private.h"

#if defined(MAGICKCORE_OPENCL_SUPPORT)
#if defined(MAGICKCORE_LTDL_DELEGATE)
#include "ltdl.h"
#endif

#ifndef MAGICKCORE_WINDOWS_SUPPORT
#include <dlfcn.h>
#endif

#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
#define MAGICKCORE_OPENCL_MACOSX  1
#endif

/*
  Define declarations.
*/
#define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"

/*
  Typedef declarations.
*/
typedef struct
{
  long long freq;
  long long clocks;
  long long start;
} AccelerateTimer;

typedef struct
{
  char
    *name,
    *platform_name,
    *vendor_name,
    *version;

  cl_uint
    max_clock_frequency,
    max_compute_units;

  double
    score;
} MagickCLDeviceBenchmark;

/*
  Forward declarations.
*/

static MagickBooleanType
  HasOpenCLDevices(MagickCLEnv,ExceptionInfo *),
  LoadOpenCLLibrary(void);

static MagickCLDevice
  RelinquishMagickCLDevice(MagickCLDevice);

static MagickCLEnv
  RelinquishMagickCLEnv(MagickCLEnv);

static void
  BenchmarkOpenCLDevices(MagickCLEnv);

extern const char
  *accelerateKernels, *accelerateKernels2;

/* OpenCL library */
MagickLibrary
  *openCL_library;

/* Default OpenCL environment */
MagickCLEnv
  default_CLEnv;
MagickThreadType
  test_thread_id=0;
SemaphoreInfo
  *openCL_lock;

/* Cached location of the OpenCL cache files */
char
  *cache_directory;
SemaphoreInfo
  *cache_directory_lock;

static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
  MagickCLDevice b)
{
  if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
      (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
      (LocaleCompare(a->name,b->name) == 0) &&
      (LocaleCompare(a->version,b->version) == 0) &&
      (a->max_clock_frequency == b->max_clock_frequency) &&
      (a->max_compute_units == b->max_compute_units))
    return(MagickTrue);

  return(MagickFalse);
}

static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
  MagickCLDeviceBenchmark *b)
{
  if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
      (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
      (LocaleCompare(a->name,b->name) == 0) &&
      (LocaleCompare(a->version,b->version) == 0) &&
      (a->max_clock_frequency == b->max_clock_frequency) &&
      (a->max_compute_units == b->max_compute_units))
    return(MagickTrue);

  return(MagickFalse);
}

static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
{
  size_t
    i;

  if (clEnv->devices != (MagickCLDevice *) NULL)
    {
      for (i = 0; i < clEnv->number_devices; i++)
        clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
      clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
    }
  clEnv->number_devices=0;
}

static inline MagickBooleanType MagickCreateDirectory(const char *path)
{
  int
    status;

#ifdef MAGICKCORE_WINDOWS_SUPPORT
  status=mkdir(path);
#else
  status=mkdir(path,0777);
#endif
  return(status == 0 ? MagickTrue : MagickFalse);
}

static inline void InitAccelerateTimer(AccelerateTimer *timer)
{
#ifdef _WIN32
  QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
#else
  timer->freq=(long long)1.0E3;
#endif
  timer->clocks=0;
  timer->start=0;
}

static inline double ReadAccelerateTimer(AccelerateTimer *timer)
{
  return (double)timer->clocks/(double)timer->freq;
}

static inline void StartAccelerateTimer(AccelerateTimer* timer)
{
#ifdef _WIN32
  QueryPerformanceCounter((LARGE_INTEGER*)&timer->start);
#else
  struct timeval
    s;
  gettimeofday(&s,0);
  timer->start=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
    (long long)1.0E3;
#endif
}

static inline void StopAccelerateTimer(AccelerateTimer *timer)
{
  long long
    n;

  n=0;
#ifdef _WIN32
  QueryPerformanceCounter((LARGE_INTEGER*)&(n));
#else
  struct timeval
    s;
  gettimeofday(&s,0);
  n=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
    (long long)1.0E3;
#endif
  n-=timer->start;
  timer->start=0;
  timer->clocks+=n;
}

static const char *GetOpenCLCacheDirectory()
{
  if (cache_directory == (char *) NULL)
    {
      if (cache_directory_lock == (SemaphoreInfo *) NULL)
        ActivateSemaphoreInfo(&cache_directory_lock);
      LockSemaphoreInfo(cache_directory_lock);
      if (cache_directory == (char *) NULL)
        {
          char
            *home,
            path[MagickPathExtent],
            *temp;

          MagickBooleanType
            status;

          struct stat
            attributes;

          temp=(char *) NULL;
          home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
          if (home == (char *) NULL)
            {
              home=GetEnvironmentValue("XDG_CACHE_HOME");
#if defined(MAGICKCORE_WINDOWS_SUPPORT) || defined(__MINGW32__)
              if (home == (char *) NULL)
                home=GetEnvironmentValue("LOCALAPPDATA");
              if (home == (char *) NULL)
                home=GetEnvironmentValue("APPDATA");
              if (home == (char *) NULL)
                home=GetEnvironmentValue("USERPROFILE");
#endif
            }

          if (home != (char *) NULL)
            {
              /* first check if $HOME exists */
              (void) FormatLocaleString(path,MagickPathExtent,"%s",home);
              status=GetPathAttributes(path,&attributes);
              if (status == MagickFalse)
                status=MagickCreateDirectory(path);

              /* first check if $HOME/ImageMagick exists */
              if (status != MagickFalse)
                {
                  (void) FormatLocaleString(path,MagickPathExtent,
                    "%s%sImageMagick",home,DirectorySeparator);

                  status=GetPathAttributes(path,&attributes);
                  if (status == MagickFalse)
                    status=MagickCreateDirectory(path);
                }

              if (status != MagickFalse)
                {
                  temp=(char*) AcquireCriticalMemory(strlen(path)+1);
                  CopyMagickString(temp,path,strlen(path)+1);
                }
              home=DestroyString(home);
            }
          else
            {
              home=GetEnvironmentValue("HOME");
              if (home != (char *) NULL)
                {
                  /* first check if $HOME/.cache exists */
                  (void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache",
                    home,DirectorySeparator);
                  status=GetPathAttributes(path,&attributes);
                  if (status == MagickFalse)
                    status=MagickCreateDirectory(path);

                  /* first check if $HOME/.cache/ImageMagick exists */
                  if (status != MagickFalse)
                    {
                      (void) FormatLocaleString(path,MagickPathExtent,
                        "%s%s.cache%sImageMagick",home,DirectorySeparator,
                        DirectorySeparator);
                      status=GetPathAttributes(path,&attributes);
                      if (status == MagickFalse)
                        status=MagickCreateDirectory(path);
                    }

                  if (status != MagickFalse)
                    {
                      temp=(char*) AcquireCriticalMemory(strlen(path)+1);
                      CopyMagickString(temp,path,strlen(path)+1);
                    }
                  home=DestroyString(home);
                }
            }
          if (temp == (char *) NULL)
            {
              temp=AcquireString("?");
              (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
                "Cannot use cache directory: \"%s\"",path);
            }
          else
            (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
              "Using cache directory: \"%s\"",temp);
          cache_directory=temp;
        }
      UnlockSemaphoreInfo(cache_directory_lock);
    }
  if (*cache_directory == '?')
    return((const char *) NULL);
  return(cache_directory);
}

static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
{
  MagickCLDevice
    device;

  size_t
    i,
    j;

  (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
    "Selecting device for type: %d",(int) type);
  for (i = 0; i < clEnv->number_devices; i++)
    clEnv->devices[i]->enabled=MagickFalse;

  for (i = 0; i < clEnv->number_devices; i++)
  {
    device=clEnv->devices[i];
    if (device->type != type)
      continue;

    device->enabled=MagickTrue;
    (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
      "Selected device: %s",device->name);
    for (j = i+1; j < clEnv->number_devices; j++)
    {
      MagickCLDevice
        other_device;

      other_device=clEnv->devices[j];
      if (IsSameOpenCLDevice(device,other_device))
        other_device->enabled=MagickTrue;
    }
  }
}

static size_t StringSignature(const char* string)
{
  size_t
    n,
    i,
    j,
    signature,
    stringLength;

  union
  {
    const char* s;
    const size_t* u;
  } p;

  stringLength=(size_t) strlen(string);
  signature=stringLength;
  n=stringLength/sizeof(size_t);
  p.s=string;
  for (i = 0; i < n; i++)
    signature^=p.u[i];
  if (n * sizeof(size_t) != stringLength)
    {
      char
        padded[4];

      j=n*sizeof(size_t);
      for (i = 0; i < 4; i++, j++)
      {
        if (j < stringLength)
          padded[i]=p.s[j];
        else
          padded[i]=0;
      }
      p.s=padded;
      signature^=p.u[0];
    }
  return(signature);
}

static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
{
  ssize_t
    i;

  for (i=0; i < (ssize_t) info->event_count; i++)
    openCL_library->clReleaseEvent(info->events[i]);
  info->events=(cl_event *) RelinquishMagickMemory(info->events);
  if (info->buffer != (cl_mem) NULL)
    openCL_library->clReleaseMemObject(info->buffer);
  RelinquishSemaphoreInfo(&info->events_semaphore);
  ReleaseOpenCLDevice(info->device);
  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 void RetainOpenCLMemObject(cl_mem memobj)
{
  (void) openCL_library->clRetainMemObject(memobj);
}

MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,size_t arg_index,
  size_t arg_size,const void *arg_value)
{
  return(openCL_library->clSetKernelArg(kernel,(cl_uint) 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)
{
  cl_int
    status;

  MagickCLCacheInfo
    info;

  info=(MagickCLCacheInfo) AcquireCriticalMemory(sizeof(*info));
  (void) memset(info,0,sizeof(*info));
  LockSemaphoreInfo(openCL_lock);
  device->requested++;
  UnlockSemaphoreInfo(openCL_lock);
  info->device=device;
  info->length=length;
  info->pixels=pixels;
  info->events_semaphore=AcquireSemaphoreInfo();
  info->buffer=openCL_library->clCreateBuffer(device->context,
    CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,
    &status);
  if (status == CL_SUCCESS)
    return(info);
  DestroyMagickCLCacheInfo(info);
  return((MagickCLCacheInfo) NULL);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   A c q u i r e M a g i c k C L D e v i c e                                 %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  AcquireMagickCLDevice() acquires an OpenCL device
%
%  The format of the AcquireMagickCLDevice method is:
%
%      MagickCLDevice AcquireMagickCLDevice()
%
*/

static MagickCLDevice AcquireMagickCLDevice()
{
  MagickCLDevice
    device;

  device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device));
  if (device != NULL)
  {
    (void) memset(device,0,sizeof(*device));
    ActivateSemaphoreInfo(&device->lock);
    device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
    device->command_queues_index=-1;
    device->enabled=MagickTrue;
  }
  return(device);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   A c q u i r e M a g i c k C L E n v                                       %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
% AcquireMagickCLEnv() allocates the MagickCLEnv structure
%
*/

static MagickCLEnv AcquireMagickCLEnv(void)
{
  const char
    *option;

  MagickCLEnv
    clEnv;

  clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv));
  if (clEnv != (MagickCLEnv) NULL)
  {
    (void) memset(clEnv,0,sizeof(*clEnv));
    ActivateSemaphoreInfo(&clEnv->lock);
    clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
    clEnv->enabled=MagickFalse;
    option=getenv("MAGICK_OCL_DEVICE");
    if (option != (const char *) NULL)
      {
        if ((IsStringTrue(option) != MagickFalse) ||
            (strcmp(option,"GPU") == 0) ||
            (strcmp(option,"CPU") == 0))
          clEnv->enabled=MagickTrue;
      }
  }
  return clEnv;
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   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                         %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  AcquireOpenCLCommandQueue() acquires an OpenCL command queue
%
%  The format of the AcquireOpenCLCommandQueue method is:
%
%      cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
%
*/

MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
{
  cl_command_queue
    queue;

  cl_command_queue_properties
    properties;

  assert(device != (MagickCLDevice) NULL);
  LockSemaphoreInfo(device->lock);
  if ((device->profile_kernels == MagickFalse) &&
      (device->command_queues_index >= 0))
  {
    queue=device->command_queues[device->command_queues_index--];
    UnlockSemaphoreInfo(device->lock);
  }
  else
  {
    UnlockSemaphoreInfo(device->lock);
    properties=0;
    if (device->profile_kernels != MagickFalse)
      properties=CL_QUEUE_PROFILING_ENABLE;
    queue=openCL_library->clCreateCommandQueue(device->context,
      device->deviceID,properties,(cl_int *) NULL);
  }
  return(queue);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   A c q u i r e O p e n C L K e r n e l                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  AcquireOpenCLKernel() acquires an OpenCL kernel
%
%  The format of the AcquireOpenCLKernel method is:
%
%      cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
%        MagickOpenCLProgram program, const char* kernelName)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o program: the OpenCL program module that the kernel belongs to.
%
%    o kernelName:  the name of the kernel
%
*/

MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
  const char *kernel_name)
{
  cl_kernel
    kernel;

  assert(device != (MagickCLDevice) NULL);
  (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),"Using kernel: %s",
    kernel_name);
  kernel=openCL_library->clCreateKernel(device->program,kernel_name,
    (cl_int *) NULL);
  return(kernel);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   A u t o S e l e c t O p e n C L D e v i c e s                             %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  AutoSelectOpenCLDevices() determines the best device based on the
%  information from the micro-benchmark.
%
%  The format of the AutoSelectOpenCLDevices method is:
%
%      void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o exception: return any errors or warnings in this structure.
%
*/

static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
{
  char
    keyword[MagickPathExtent],
    *token;

  const char
    *q;

  MagickCLDeviceBenchmark
    *device_benchmark;

  size_t
    i,
    extent;

  if (xml == (char *) NULL)
    return;
  device_benchmark=(MagickCLDeviceBenchmark *) NULL;
  token=AcquireString(xml);
  extent=strlen(token)+MagickPathExtent;
  for (q=(char *) xml; *q != '\0'; )
  {
    /*
      Interpret XML.
    */
    (void) GetNextToken(q,&q,extent,token);
    if (*token == '\0')
      break;
    (void) CopyMagickString(keyword,token,MagickPathExtent);
    if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
      {
        /*
          Doctype element.
        */
        while ((LocaleNCompare(q,"]>",2) != 0) && (*q != '\0'))
          (void) GetNextToken(q,&q,extent,token);
        continue;
      }
    if (LocaleNCompare(keyword,"<!--",4) == 0)
      {
        /*
          Comment element.
        */
        while ((LocaleNCompare(q,"->",2) != 0) && (*q != '\0'))
          (void) GetNextToken(q,&q,extent,token);
        continue;
      }
    if (LocaleCompare(keyword,"<device") == 0)
      {
        /*
          Device element.
        */
        device_benchmark=(MagickCLDeviceBenchmark *) AcquireQuantumMemory(1,
          sizeof(*device_benchmark));
        if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
          break;
        (void) memset(device_benchmark,0,sizeof(*device_benchmark));
        device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
        continue;
      }
    if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
      continue;
    if (LocaleCompare(keyword,"/>") == 0)
      {
        if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
          {
            if (LocaleCompare(device_benchmark->name,"CPU") == 0)
              clEnv->cpu_score=device_benchmark->score;
            else
              {
                MagickCLDevice
                  device;

                /*
                  Set the score for all devices that match this device.
                */
                for (i = 0; i < clEnv->number_devices; i++)
                {
                  device=clEnv->devices[i];
                  if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
                    device->score=device_benchmark->score;
                }
              }
          }

        device_benchmark->platform_name=RelinquishMagickMemory(
          device_benchmark->platform_name);
        device_benchmark->vendor_name=RelinquishMagickMemory(
          device_benchmark->vendor_name);
        device_benchmark->name=RelinquishMagickMemory(device_benchmark->name);
        device_benchmark->version=RelinquishMagickMemory(
          device_benchmark->version);
        device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
          device_benchmark);
        continue;
      }
    (void) GetNextToken(q,(const char **) NULL,extent,token);
    if (*token != '=')
      continue;
    (void) GetNextToken(q,&q,extent,token);
    (void) GetNextToken(q,&q,extent,token);
    switch (*keyword)
    {
      case 'M':
      case 'm':
      {
        if (LocaleCompare((char *) keyword,"maxClockFrequency") == 0)
          {
            device_benchmark->max_clock_frequency=StringToInteger(token);
            break;
          }
        if (LocaleCompare((char *) keyword,"maxComputeUnits") == 0)
          {
            device_benchmark->max_compute_units=StringToInteger(token);
            break;
          }
        break;
      }
      case 'N':
      case 'n':
      {
        if (LocaleCompare((char *) keyword,"name") == 0)
          device_benchmark->name=ConstantString(token);
        break;
      }
      case 'P':
      case 'p':
      {
        if (LocaleCompare((char *) keyword,"platform") == 0)
          device_benchmark->platform_name=ConstantString(token);
        break;
      }
      case 'S':
      case 's':
      {
        if (LocaleCompare((char *) keyword,"score") == 0)
          device_benchmark->score=StringToDouble(token,(char **) NULL);
        break;
      }
      case 'V':
      case 'v':
      {
        if (LocaleCompare((char *) keyword,"vendor") == 0)
          device_benchmark->vendor_name=ConstantString(token);
        if (LocaleCompare((char *) keyword,"version") == 0)
          device_benchmark->version=ConstantString(token);
        break;
      }
      default:
        break;
    }
  }
  token=(char *) RelinquishMagickMemory(token);
  device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
    device_benchmark);
}

static MagickBooleanType CanWriteProfileToFile(const char *filename)
{
  FILE
    *profileFile;

  profileFile=fopen(filename,"ab");

  if (profileFile == (FILE *) NULL)
    {
      (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
        "Unable to save profile to: \"%s\"",filename);
      return(MagickFalse);
    }

  fclose(profileFile);
  return(MagickTrue);
}

static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
{
  char
    filename[MagickPathExtent];

  StringInfo
    *option;

  size_t
    i;

  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
    GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);

  /*
    We don't run the benchmark when we can not write out a device profile. The
    first GPU device will be used.
  */
#if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
  if (CanWriteProfileToFile(filename) == MagickFalse)
#endif
    {
      for (i = 0; i < clEnv->number_devices; i++)
        clEnv->devices[i]->score=1.0;

      SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
      return(MagickFalse);
    }

  option=ConfigureFileToStringInfo(filename);
  LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(option));
  option=DestroyStringInfo(option);
  return(MagickTrue);
}

static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
{
  const char
    *option;

  double
    best_score;

  MagickBooleanType
    benchmark;

  size_t
    i;

  option=getenv("MAGICK_OCL_DEVICE");
  if (option != (const char *) NULL)
    {
      if (strcmp(option,"GPU") == 0)
        SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
      else if (strcmp(option,"CPU") == 0)
        SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
    }

  if (LoadOpenCLBenchmarks(clEnv) == MagickFalse)
    return;

  benchmark=MagickFalse;
  if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
    benchmark=MagickTrue;
  else
    {
      for (i = 0; i < clEnv->number_devices; i++)
      {
        if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
        {
          benchmark=MagickTrue;
          break;
        }
      }
    }

  if (benchmark != MagickFalse)
    BenchmarkOpenCLDevices(clEnv);

  best_score=clEnv->cpu_score;
  for (i = 0; i < clEnv->number_devices; i++)
    best_score=MagickMin(clEnv->devices[i]->score,best_score);

  for (i = 0; i < clEnv->number_devices; i++)
  {
    if (clEnv->devices[i]->score != best_score)
      clEnv->devices[i]->enabled=MagickFalse;
  }
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   B e n c h m a r k O p e n C L D e v i c e s                               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  BenchmarkOpenCLDevices() benchmarks the OpenCL devices and the CPU to help
%  the automatic selection of the best device.
%
%  The format of the BenchmarkOpenCLDevices method is:
%
%    void BenchmarkOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o exception: return any errors or warnings
*/

static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
{
  AccelerateTimer
    timer;

  ExceptionInfo
    *exception;

  Image
    *inputImage;

  ImageInfo
    *imageInfo;

  size_t
    i;

  exception=AcquireExceptionInfo();
  imageInfo=AcquireImageInfo();
  CloneString(&imageInfo->size,"2048x1536");
  CopyMagickString(imageInfo->filename,"xc:none",MagickPathExtent);
  inputImage=ReadImage(imageInfo,exception);

  InitAccelerateTimer(&timer);

  for (i=0; i<=2; i++)
  {
    Image
      *bluredImage,
      *resizedImage,
      *unsharpedImage;

    if (i > 0)
      StartAccelerateTimer(&timer);

    bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
    unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
      exception);
    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);

    if (bluredImage != (Image *) NULL)
      DestroyImage(bluredImage);
    if (unsharpedImage != (Image *) NULL)
      DestroyImage(unsharpedImage);
    if (resizedImage != (Image *) NULL)
      DestroyImage(resizedImage);
  }
  DestroyImage(inputImage);
  return(ReadAccelerateTimer(&timer));
}

static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
  MagickCLDevice device)
{
  testEnv->devices[0]=device;
  default_CLEnv=testEnv;
  device->score=RunOpenCLBenchmark(MagickFalse);
  default_CLEnv=clEnv;
  testEnv->devices[0]=(MagickCLDevice) NULL;
}

static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
{
  char
    filename[MagickPathExtent];

  FILE
    *cache_file;

  MagickCLDevice
    device;

  size_t
    i,
    j;

  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
    GetOpenCLCacheDirectory(),DirectorySeparator,
    IMAGEMAGICK_PROFILE_FILE);

  cache_file=fopen_utf8(filename,"wb");
  if (cache_file == (FILE *) NULL)
    return;
  fwrite("<devices>\n",sizeof(char),10,cache_file);
  fprintf(cache_file,"  <device name=\"CPU\" score=\"%.4g\"/>\n",
    clEnv->cpu_score);
  for (i = 0; i < clEnv->number_devices; i++)
  {
    MagickBooleanType
      duplicate;

    device=clEnv->devices[i];
    duplicate=MagickFalse;
    for (j = 0; j < i; j++)
    {
      if (IsSameOpenCLDevice(clEnv->devices[j],device))
      {
        duplicate=MagickTrue;
        break;
      }
    }

    if (duplicate)
      continue;

    if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
      fprintf(cache_file,"  <device platform=\"%s\" vendor=\"%s\" name=\"%s\"\
 version=\"%s\" maxClockFrequency=\"%d\" maxComputeUnits=\"%d\"\
 score=\"%.4g\"/>\n",
        device->platform_name,device->vendor_name,device->name,device->version,
        (int)device->max_clock_frequency,(int)device->max_compute_units,
        device->score);
  }
  fwrite("</devices>",sizeof(char),10,cache_file);

  fclose(cache_file);
}

static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
{
  MagickCLDevice
    device;

  MagickCLEnv
    testEnv;

  size_t
    i,
    j;

  (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
    "Starting benchmark");
  testEnv=AcquireMagickCLEnv();
  testEnv->library=openCL_library;
  testEnv->devices=(MagickCLDevice *) AcquireCriticalMemory(
    sizeof(MagickCLDevice));
  testEnv->number_devices=1;
  testEnv->benchmark_thread_id=GetMagickThreadId();
  testEnv->initialized=MagickTrue;

  for (i = 0; i < clEnv->number_devices; i++)
    clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;

  for (i = 0; i < clEnv->number_devices; i++)
  {
    device=clEnv->devices[i];
    if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
      RunDeviceBenckmark(clEnv,testEnv,device);

    /* Set the score on all the other devices that are the same */
    for (j = i+1; j < clEnv->number_devices; j++)
    {
      MagickCLDevice
        other_device;

      other_device=clEnv->devices[j];
      if (IsSameOpenCLDevice(device,other_device))
        other_device->score=device->score;
    }
  }

  testEnv->enabled=MagickFalse;
  default_CLEnv=testEnv;
  clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
  default_CLEnv=clEnv;

  testEnv=RelinquishMagickCLEnv(testEnv);
  CacheOpenCLBenchmarks(clEnv);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   C o m p i l e O p e n C L K e r n e l                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  CompileOpenCLKernel() compiles the kernel for the specified device. The
%  kernel will be cached on disk to reduce the compilation time.
%
%  The format of the CompileOpenCLKernel method is:
%
%      MagickBooleanType AcquireOpenCLKernel(MagickCLDevice clEnv,
%        unsigned int signature,const char *kernel,const char *options,
%        ExceptionInfo *exception)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
%
%    o kernel: the source code of the kernel.
%
%    o options: options for the compiler.
%
%    o signature: a number to uniquely identify the kernel
%
%    o exception: return any errors or warnings in this structure.
%
*/

static void CacheOpenCLKernel(MagickCLDevice device,char *filename,
  ExceptionInfo *exception)
{
  cl_uint
    status;

  size_t
    binaryProgramSize;

  unsigned char
    *binaryProgram;

  status=openCL_library->clGetProgramInfo(device->program,
    CL_PROGRAM_BINARY_SIZES,sizeof(size_t),&binaryProgramSize,NULL);
  if (status != CL_SUCCESS)
    return;
  binaryProgram=(unsigned char*) AcquireQuantumMemory(1,binaryProgramSize);
  if (binaryProgram == (unsigned char *) NULL)
    {
      (void) ThrowMagickException(exception,GetMagickModule(),
        ResourceLimitError,"MemoryAllocationFailed","`%s'",filename);
      return;
    }
  status=openCL_library->clGetProgramInfo(device->program,
    CL_PROGRAM_BINARIES,sizeof(unsigned char*),&binaryProgram,NULL);
  if (status == CL_SUCCESS)
    {
      (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
        "Creating cache file: \"%s\"",filename);
      (void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
    }
  binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
}

static MagickBooleanType LoadCachedOpenCLKernels(MagickCLDevice device,
  const char *filename)
{
  cl_int
    binaryStatus,
    status;

  ExceptionInfo
    *sans_exception;

  size_t
    length;

  unsigned char
    *binaryProgram;

  sans_exception=AcquireExceptionInfo();
  binaryProgram=(unsigned char *) FileToBlob(filename,~0UL,&length,
    sans_exception);
  sans_exception=DestroyExceptionInfo(sans_exception);
  if (binaryProgram == (unsigned char *) NULL)
    return(MagickFalse);
  (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
    "Loaded cached kernels: \"%s\"",filename);
  device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
    &device->deviceID,&length,(const unsigned char**)&binaryProgram,
    &binaryStatus,&status);
  binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
  return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
    MagickTrue);
}

static void LogOpenCLBuildFailure(MagickCLDevice device,const char *kernel,
  ExceptionInfo *exception)
{
  char
    filename[MagickPathExtent],
    *log;

  size_t
    log_size;

  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
    GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.cl");

  (void) remove_utf8(filename);
  (void) BlobToFile(filename,kernel,strlen(kernel),exception);

  openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
    CL_PROGRAM_BUILD_LOG,0,NULL,&log_size);
  log=(char*)AcquireCriticalMemory(log_size);
  openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
    CL_PROGRAM_BUILD_LOG,log_size,log,&log_size);

  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
    GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.log");

  (void) remove_utf8(filename);
  (void) BlobToFile(filename,log,log_size,exception);
  log=(char*)RelinquishMagickMemory(log);
}

static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
  const char *kernel,const char *options,size_t signature,
  ExceptionInfo *exception)
{
  char
    deviceName[MagickPathExtent],
    filename[MagickPathExtent],
    *ptr;

  cl_int
    status;

  MagickBooleanType
    loaded;

  size_t
    length;

  (void) CopyMagickString(deviceName,device->name,MagickPathExtent);
  ptr=deviceName;
  /* Strip out illegal characters for file names */
  while (*ptr != '\0')
  {
    if ((*ptr == ' ') || (*ptr == '\\') || (*ptr == '/') || (*ptr == ':') ||
        (*ptr == '*') || (*ptr == '?') || (*ptr == '"') || (*ptr == '<') ||
        (*ptr == '>' || *ptr == '|'))
      *ptr = '_';
    ptr++;
  }
  (void) FormatLocaleString(filename,MagickPathExtent,
    "%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
    DirectorySeparator,"magick_opencl",deviceName,(unsigned int) signature,
    (double) sizeof(char*)*8);
  loaded=LoadCachedOpenCLKernels(device,filename);
  if (loaded == MagickFalse)
    {
      /* Binary CL program unavailable, compile the program from source */
      length=strlen(kernel);
      device->program=openCL_library->clCreateProgramWithSource(
        device->context,1,&kernel,&length,&status);
      if (status != CL_SUCCESS)
        return(MagickFalse);
    }

  status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
    options,NULL,NULL);
  if (status != CL_SUCCESS)
  {
    (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
      "clBuildProgram failed.","(%d)",(int)status);
    LogOpenCLBuildFailure(device,kernel,exception);
    return(MagickFalse);
  }

  /* Save the binary to a file to avoid re-compilation of the kernels */
  if (loaded == MagickFalse)
    CacheOpenCLKernel(device,filename,exception);

  return(MagickTrue);
}

static cl_event* CopyOpenCLEvents(MagickCLCacheInfo first,
  MagickCLCacheInfo second,cl_uint *event_count)
{
  cl_event
    *events;

  size_t
    i;

  size_t
    j;

  assert(first != (MagickCLCacheInfo) NULL);
  assert(event_count != (cl_uint *) NULL);
  events=(cl_event *) NULL;
  LockSemaphoreInfo(first->events_semaphore);
  if (second != (MagickCLCacheInfo) NULL)
    LockSemaphoreInfo(second->events_semaphore);
  *event_count=first->event_count;
  if (second != (MagickCLCacheInfo) NULL)
    *event_count+=second->event_count;
  if (*event_count > 0)
    {
      events=AcquireQuantumMemory(*event_count,sizeof(*events));
      if (events == (cl_event *) NULL)
        *event_count=0;
      else
        {
          j=0;
          for (i=0; i < first->event_count; i++, j++)
            events[j]=first->events[i];
          if (second != (MagickCLCacheInfo) NULL)
            {
              for (i=0; i < second->event_count; i++, j++)
                events[j]=second->events[i];
            }
        }
    }
  UnlockSemaphoreInfo(first->events_semaphore);
  if (second != (MagickCLCacheInfo) NULL)
    UnlockSemaphoreInfo(second->events_semaphore);
  return(events);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   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;

  cl_event
    *events;

  cl_uint
    event_count;

  Quantum
    *pixels;

  if (info == (MagickCLCacheInfo) NULL)
    return((MagickCLCacheInfo) NULL);
  events=CopyOpenCLEvents(info,(MagickCLCacheInfo) NULL,&event_count);
  if (events != (cl_event *) NULL)
    {
      queue=AcquireOpenCLCommandQueue(info->device);
      pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
        CL_MAP_READ | CL_MAP_WRITE,0,info->length,event_count,events,
        (cl_event *) NULL,(cl_int *) NULL);
      assert(pixels == info->pixels);
      ReleaseOpenCLCommandQueue(info->device,queue);
      events=(cl_event *) RelinquishMagickMemory(events);
    }
  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                                 %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  DumpOpenCLProfileData() dumps the kernel profile data.
%
%  The format of the DumpProfileData method is:
%
%      void DumpProfileData()
%
*/

MagickPrivate void DumpOpenCLProfileData()
{
#define OpenCLLog(message) \
   fwrite(message,sizeof(char),strlen(message),log); \
   fwrite("\n",sizeof(char),1,log);

  char
    buf[4096],
    filename[MagickPathExtent],
    indent[160];

  FILE
    *log;

  size_t
    i,
    j;

  if (default_CLEnv == (MagickCLEnv) NULL)
    return;

  for (i = 0; i < default_CLEnv->number_devices; i++)
    if (default_CLEnv->devices[i]->profile_kernels != MagickFalse)
      break;
  if (i == default_CLEnv->number_devices)
    return;

  (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
    GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");

  log=fopen_utf8(filename,"wb");
  if (log == (FILE *) NULL)
    return;
  for (i = 0; i < default_CLEnv->number_devices; i++)
  {
    MagickCLDevice
      device;

    device=default_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("                     -------   -----     ---     ---");
    j=0;
    while (device->profile_records[j] != (KernelProfileRecord) NULL)
    {
      KernelProfileRecord
        profile;

      profile=device->profile_records[j];
      strcpy(indent,"                    ");
      CopyMagickString(indent,profile->kernel_name,MagickMin(strlen(
        profile->kernel_name),strlen(indent)));
      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);
}
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   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.
%
*/

static MagickBooleanType RegisterCacheEvent(MagickCLCacheInfo info,
  cl_event event)
{
  assert(info != (MagickCLCacheInfo) NULL);
  assert(event != (cl_event) NULL);
  if (openCL_library->clRetainEvent(event) != CL_SUCCESS)
    {
      openCL_library->clWaitForEvents(1,&event);
      return(MagickFalse);
    }
  LockSemaphoreInfo(info->events_semaphore);
  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;
  UnlockSemaphoreInfo(info->events_semaphore);
  return(MagickTrue);
}

MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
  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,
  MagickBooleanType flush,ExceptionInfo *exception)
{
  CacheInfo
    *output_info,
    *input_info;

  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);
  output_info=(CacheInfo *) NULL;
  if (output_image == (const Image *) NULL)
    events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
      &event_count);
  else
    {
      output_info=(CacheInfo *) output_image->cache;
      assert(output_info != (CacheInfo *) NULL);
      assert(output_info->opencl != (MagickCLCacheInfo) NULL);
      events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
        &event_count);
    }
  status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
    gsize,lsize,event_count,events,&event);
  /* This can fail due to memory issues and calling clFinish might help. */
  if ((status != CL_SUCCESS) && (event_count > 0))
    {
      openCL_library->clFinish(queue);
      status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
        offset,gsize,lsize,event_count,events,&event);
    }
  events=(cl_event *) RelinquishMagickMemory(events);
  if (status != CL_SUCCESS)
    {
      (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
        GetMagickModule(),ResourceLimitWarning,
        "clEnqueueNDRangeKernel failed.","'%s'",".");
      return(MagickFalse);
    }
  if (flush != MagickFalse)
    openCL_library->clFlush(queue);
  if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
    {
      if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
        {
          if (output_info != (CacheInfo *) NULL)
            (void) RegisterCacheEvent(output_info->opencl,event);
        }
    }
  openCL_library->clReleaseEvent(event);
  return(MagickTrue);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   G e t C u r r e n t O p e n C L E n v                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetCurrentOpenCLEnv() returns the current OpenCL env
%
%  The format of the GetCurrentOpenCLEnv method is:
%
%      MagickCLEnv GetCurrentOpenCLEnv()
%
*/

MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
{
  if (default_CLEnv != (MagickCLEnv) NULL)
  {
    if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
        (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
      return((MagickCLEnv) NULL);
    else
      return(default_CLEnv);
  }

  if (GetOpenCLCacheDirectory() == (char *) NULL)
    return((MagickCLEnv) NULL);

  if (openCL_lock == (SemaphoreInfo *) NULL)
    ActivateSemaphoreInfo(&openCL_lock);

  LockSemaphoreInfo(openCL_lock);
  if (default_CLEnv == (MagickCLEnv) NULL)
    default_CLEnv=AcquireMagickCLEnv();
  UnlockSemaphoreInfo(openCL_lock);

  return(default_CLEnv);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e B e n c h m a r k D u r a t i o n           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
%  device. The score is determined by the duration of the micro benchmark so
%  that means a lower score is better than a higher score.
%
%  The format of the GetOpenCLDeviceBenchmarkScore method is:
%
%      double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
*/

MagickExport double GetOpenCLDeviceBenchmarkScore(
  const MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
  return(device->score);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e E n a b l e d                               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceEnabled() returns true if the device is enabled.
%
%  The format of the GetOpenCLDeviceEnabled method is:
%
%      MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
*/

MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
  const MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return(MagickFalse);
  return(device->enabled);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e N a m e                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceName() returns the name of the device.
%
%  The format of the GetOpenCLDeviceName method is:
%
%      const char *GetOpenCLDeviceName(const MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
*/

MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return((const char *) NULL);
  return(device->name);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e V e n d o r N a m e                         %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceVendorName() returns the vendor name of the device.
%
%  The format of the GetOpenCLDeviceVendorName method is:
%
%      const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
*/

MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return((const char *) NULL);
  return(device->vendor_name);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e s                                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
%  value of length to the number of devices that are available.
%
%  The format of the GetOpenCLDevices method is:
%
%      const MagickCLDevice *GetOpenCLDevices(size_t *length,
%        ExceptionInfo *exception)
%
%  A description of each parameter follows:
%
%    o length: the number of device.
%
%    o exception: return any errors or warnings in this structure.
%
*/

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);
    }
  InitializeOpenCL(clEnv,exception);
  if (length != (size_t *) NULL)
    *length=clEnv->number_devices;
  return(clEnv->devices);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e T y p e                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceType() returns the type of the device.
%
%  The format of the GetOpenCLDeviceType method is:
%
%      MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
*/

MagickExport MagickCLDeviceType GetOpenCLDeviceType(
  const MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return(UndefinedCLDeviceType);
  if (device->type == CL_DEVICE_TYPE_GPU)
    return(GpuCLDeviceType);
  if (device->type == CL_DEVICE_TYPE_CPU)
    return(CpuCLDeviceType);
  return(UndefinedCLDeviceType);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L D e v i c e V e r s i o n                               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceVersion() returns the version of the device.
%
%  The format of the GetOpenCLDeviceName method is:
%
%      const char *GetOpenCLDeviceVersion(MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
*/

MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return((const char *) NULL);
  return(device->version);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   G e t O p e n C L E n a b l e d                                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
%
%  The format of the GetOpenCLEnabled method is:
%
%      MagickBooleanType GetOpenCLEnabled()
%
*/

MagickExport MagickBooleanType GetOpenCLEnabled(void)
{
  MagickCLEnv
    clEnv;

  clEnv=GetCurrentOpenCLEnv();
  if (clEnv == (MagickCLEnv) NULL)
    return(MagickFalse);
  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;
      LockSemaphoreInfo(device->lock);
      while (device->profile_records[*length] != (KernelProfileRecord) NULL)
        *length=*length+1;
      UnlockSemaphoreInfo(device->lock);
    }
  return(device->profile_records);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   H a s O p e n C L D e v i c e s                                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  HasOpenCLDevices() checks if the OpenCL environment has devices that are
%  enabled and compiles the kernel for the device when necessary. False will be
%  returned if no enabled devices could be found
%
%  The format of the HasOpenCLDevices method is:
%
%    MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
%      ExceptionInfo exception)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o exception: return any errors or warnings in this structure.
%
*/

static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
  ExceptionInfo *exception)
{
  char
    *accelerateKernelsBuffer,
    options[MagickPathExtent];

  MagickStatusType
    status;

  size_t
    i;

  size_t
    signature;

  /* Check if there are enabled devices */
  for (i = 0; i < clEnv->number_devices; i++)
  {
    if ((clEnv->devices[i]->enabled != MagickFalse))
      break;
  }
  if (i == clEnv->number_devices)
    return(MagickFalse);

  /* Check if we need to compile a kernel for one of the devices */
  status=MagickTrue;
  for (i = 0; i < clEnv->number_devices; i++)
  {
    if ((clEnv->devices[i]->enabled != MagickFalse) &&
        (clEnv->devices[i]->program == (cl_program) NULL))
    {
      status=MagickFalse;
      break;
    }
  }
  if (status != MagickFalse)
    return(MagickTrue);

  /* Get additional options */
  (void) FormatLocaleString(options,MagickPathExtent,CLOptions,
    (float)QuantumRange,(float)QuantumScale,(float)CLCharQuantumScale,
    (float)MagickEpsilon,(float)MagickPI,(unsigned int)MaxMap,
    (unsigned int)MAGICKCORE_QUANTUM_DEPTH);

  signature=StringSignature(options);
  accelerateKernelsBuffer=(char*) AcquireQuantumMemory(1,
    strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
  if (accelerateKernelsBuffer == (char*) NULL)
    return(MagickFalse);
  sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
  signature^=StringSignature(accelerateKernelsBuffer);

  status=MagickTrue;
  for (i = 0; i < clEnv->number_devices; i++)
  {
    MagickCLDevice
      device;

    size_t
      device_signature;

    device=clEnv->devices[i];
    if ((device->enabled == MagickFalse) ||
        (device->program != (cl_program) NULL))
      continue;

    LockSemaphoreInfo(device->lock);
    if (device->program != (cl_program) NULL)
    {
      UnlockSemaphoreInfo(device->lock);
      continue;
    }
    device_signature=signature;
    device_signature^=StringSignature(device->platform_name);
    status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
      device_signature,exception);
    UnlockSemaphoreInfo(device->lock);
    if (status == MagickFalse)
      break;
  }
  accelerateKernelsBuffer=RelinquishMagickMemory(accelerateKernelsBuffer);
  return(status);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   I n i t i a l i z e O p e n C L                                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  InitializeOpenCL() is used to initialize the OpenCL environment. This method
%  makes sure the devices are propertly initialized and benchmarked.
%
%  The format of the InitializeOpenCL method is:
%
%    MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
%
%  A description of each parameter follows:
%
%    o exception: return any errors or warnings in this structure.
%
*/

static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
{
  char
    version[MagickPathExtent];

  cl_uint
    num;

  if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
        MagickPathExtent,version,NULL) != CL_SUCCESS)
    return(0);
  if (strncmp(version,"OpenCL 1.0 ",11) == 0)
    return(0);
  if (clEnv->library->clGetDeviceIDs(platform,
        CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
    return(0);
  return(num);
}

static inline char *GetOpenCLPlatformString(cl_platform_id platform,
  cl_platform_info param_name)
{
  char
    *value;

  size_t
    length;

  openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
  value=AcquireCriticalMemory(length*sizeof(*value));
  openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
  return(value);
}

static inline char *GetOpenCLDeviceString(cl_device_id device,
  cl_device_info param_name)
{
  char
    *value;

  size_t
    length;

  openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
  value=AcquireCriticalMemory(length*sizeof(*value));
  openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
  return(value);
}

static void LoadOpenCLDevices(MagickCLEnv clEnv)
{
  cl_context_properties
    properties[3];

  cl_device_id
    *devices;

  cl_int
    status;

  cl_platform_id
    *platforms;

  cl_uint
    i,
    j,
    next,
    number_devices,
    number_platforms;

  size_t
    length;

  number_platforms=0;
  if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
    return;
  if (number_platforms == 0)
    return;
  platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
    sizeof(cl_platform_id));
  if (platforms == (cl_platform_id *) NULL)
    return;
  if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
    {
       platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
       return;
    }
  for (i = 0; i < number_platforms; i++)
  {
    char
      *platform_name;

    number_devices=0;
    platform_name=GetOpenCLPlatformString(platforms[i],CL_PLATFORM_NAME);
    /* NVIDIA is disabled by default due to reported access violation */
    if (strncmp(platform_name,"NVIDIA",6) != 0)
      {
        number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
        clEnv->number_devices+=number_devices;
      }
    platform_name=(char *) RelinquishMagickMemory(platform_name);
    if (number_devices == 0)
      platforms[i]=(cl_platform_id) NULL;
  }
  if (clEnv->number_devices == 0)
    {
      platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
      return;
    }
  clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
    sizeof(MagickCLDevice));
  if (clEnv->devices == (MagickCLDevice *) NULL)
    {
      RelinquishMagickCLDevices(clEnv);
      platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
      return;
    }
  (void) memset(clEnv->devices,0,clEnv->number_devices*sizeof(MagickCLDevice));
  devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
    sizeof(cl_device_id));
  if (devices == (cl_device_id *) NULL)
    {
      platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
      RelinquishMagickCLDevices(clEnv);
      return;
    }
  (void) memset(devices,0,clEnv->number_devices*sizeof(cl_device_id));
  clEnv->number_contexts=(size_t) number_platforms;
  clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
    sizeof(cl_context));
  if (clEnv->contexts == (cl_context *) NULL)
    {
      devices=(cl_device_id *) RelinquishMagickMemory(devices);
      platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
      RelinquishMagickCLDevices(clEnv);
      return;
    }
  (void) memset(clEnv->contexts,0,clEnv->number_contexts*sizeof(cl_context));
  next=0;
  for (i = 0; i < number_platforms; i++)
  {
    if (platforms[i] == (cl_platform_id) NULL)
      continue;

    status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
      CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
    if (status != CL_SUCCESS)
      continue;

    properties[0]=CL_CONTEXT_PLATFORM;
    properties[1]=(cl_context_properties) platforms[i];
    properties[2]=0;
    clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
      devices,NULL,NULL,&status);
    if (status != CL_SUCCESS)
      continue;

    for (j = 0; j < number_devices; j++,next++)
    {
      MagickCLDevice
        device;

      device=AcquireMagickCLDevice();
      if (device == (MagickCLDevice) NULL)
        break;

      device->context=clEnv->contexts[i];
      device->deviceID=devices[j];

      device->platform_name=GetOpenCLPlatformString(platforms[i],
        CL_PLATFORM_NAME);

      device->vendor_name=GetOpenCLPlatformString(platforms[i],
        CL_PLATFORM_VENDOR);

      device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);

      device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);

      openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
        sizeof(cl_uint),&device->max_clock_frequency,NULL);

      openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
        sizeof(cl_uint),&device->max_compute_units,NULL);

      openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
        sizeof(cl_device_type),&device->type,NULL);

      openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
        sizeof(cl_ulong),&device->local_memory_size,NULL);

      clEnv->devices[next]=device;
      (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
        "Found device: %s (%s)",device->name,device->platform_name);
    }
  }
  if (next != clEnv->number_devices)
    RelinquishMagickCLDevices(clEnv);
  platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
  devices=(cl_device_id *) RelinquishMagickMemory(devices);
}

MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
  ExceptionInfo *exception)
{
  register
    size_t i;

  LockSemaphoreInfo(clEnv->lock);
  if (clEnv->initialized != MagickFalse)
    {
      UnlockSemaphoreInfo(clEnv->lock);
      return(HasOpenCLDevices(clEnv,exception));
    }
  if (LoadOpenCLLibrary() != MagickFalse)
    {
      clEnv->library=openCL_library;
      LoadOpenCLDevices(clEnv);
      if (clEnv->number_devices > 0)
        AutoSelectOpenCLDevices(clEnv);
    }
  clEnv->initialized=MagickTrue;
  /* NVIDIA is disabled by default due to reported access violation */
  for (i=0; i < (ssize_t) clEnv->number_devices; i++)
  {
    if (strncmp(clEnv->devices[i]->platform_name,"NVIDIA",6) == 0)
      clEnv->devices[i]->enabled=MagickFalse;
  }
  UnlockSemaphoreInfo(clEnv->lock);
  return(HasOpenCLDevices(clEnv,exception));
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   L o a d O p e n C L L i b r a r y                                         %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  LoadOpenCLLibrary() load and binds the OpenCL library.
%
%  The format of the LoadOpenCLLibrary method is:
%
%    MagickBooleanType LoadOpenCLLibrary(void)
%
*/

void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
{
  if ((library == (void *) NULL) || (functionName == (const char *) NULL))
    return (void *) NULL;
#ifdef MAGICKCORE_WINDOWS_SUPPORT
    return (void *) GetProcAddress((HMODULE)library,functionName);
#else
    return (void *) dlsym(library,functionName);
#endif
}

static MagickBooleanType BindOpenCLFunctions()
{
#ifdef MAGICKCORE_OPENCL_MACOSX
#define BIND(X) openCL_library->X= &X;
#else
  (void) memset(openCL_library,0,sizeof(MagickLibrary));
#ifdef MAGICKCORE_WINDOWS_SUPPORT
  openCL_library->library=(void *)LoadLibraryA("OpenCL.dll");
#else
  openCL_library->library=(void *)dlopen("libOpenCL.so",RTLD_NOW);
#endif
#define BIND(X) \
  if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
    return(MagickFalse);
#endif

  if (openCL_library->library == (void*) NULL)
    return(MagickFalse);

  BIND(clGetPlatformIDs);
  BIND(clGetPlatformInfo);

  BIND(clGetDeviceIDs);
  BIND(clGetDeviceInfo);

  BIND(clCreateBuffer);
  BIND(clReleaseMemObject);
  BIND(clRetainMemObject);

  BIND(clCreateContext);
  BIND(clReleaseContext);

  BIND(clCreateCommandQueue);
  BIND(clReleaseCommandQueue);
  BIND(clFlush);
  BIND(clFinish);

  BIND(clCreateProgramWithSource);
  BIND(clCreateProgramWithBinary);
  BIND(clReleaseProgram);
  BIND(clBuildProgram);
  BIND(clGetProgramBuildInfo);
  BIND(clGetProgramInfo);

  BIND(clCreateKernel);
  BIND(clReleaseKernel);
  BIND(clSetKernelArg);
  BIND(clGetKernelInfo);

  BIND(clEnqueueReadBuffer);
  BIND(clEnqueueMapBuffer);
  BIND(clEnqueueUnmapMemObject);
  BIND(clEnqueueNDRangeKernel);

  BIND(clGetEventInfo);
  BIND(clWaitForEvents);
  BIND(clReleaseEvent);
  BIND(clRetainEvent);
  BIND(clSetEventCallback);

  BIND(clGetEventProfilingInfo);

  return(MagickTrue);
}

static MagickBooleanType LoadOpenCLLibrary(void)
{
  openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
  if (openCL_library == (MagickLibrary *) NULL)
    return(MagickFalse);

  if (BindOpenCLFunctions() == MagickFalse)
    {
      openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
      return(MagickFalse);
    }

  return(MagickTrue);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   O p e n C L T e r m i n u s                                               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  OpenCLTerminus() destroys the OpenCL component.
%
%  The format of the OpenCLTerminus method is:
%
%      OpenCLTerminus(void)
%
*/

MagickPrivate void OpenCLTerminus()
{
  DumpOpenCLProfileData();
  if (cache_directory != (char *) NULL)
    cache_directory=DestroyString(cache_directory);
  if (cache_directory_lock != (SemaphoreInfo *) NULL)
    RelinquishSemaphoreInfo(&cache_directory_lock);
  if (default_CLEnv != (MagickCLEnv) NULL)
    default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
  if (openCL_lock != (SemaphoreInfo *) NULL)
    RelinquishSemaphoreInfo(&openCL_lock);
  if (openCL_library != (MagickLibrary *) NULL)
    {
      if (openCL_library->library != (void *) NULL)
        (void) lt_dlclose(openCL_library->library);
      openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
    }
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   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                       %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  OpenCLThrowMagickException logs an OpenCL exception as determined by the log
%  configuration file.  If an error occurs, MagickFalse is returned
%  otherwise MagickTrue.
%
%  The format of the OpenCLThrowMagickException method is:
%
%      MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
%        const char *module,const char *function,const size_t line,
%        const ExceptionType severity,const char *tag,const char *format,...)
%
%  A description of each parameter follows:
%
%    o exception: the exception info.
%
%    o filename: the source module filename.
%
%    o function: the function name.
%
%    o line: the line number of the source module.
%
%    o severity: Specifies the numeric error category.
%
%    o tag: the locale tag.
%
%    o format: the output format.
%
*/

MagickPrivate MagickBooleanType OpenCLThrowMagickException(
  MagickCLDevice device,ExceptionInfo *exception,const char *module,
  const char *function,const size_t line,const ExceptionType severity,
  const char *tag,const char *format,...)
{
  MagickBooleanType
    status;

  assert(device != (MagickCLDevice) NULL);
  assert(exception != (ExceptionInfo *) NULL);
  assert(exception->signature == MagickCoreSignature);
  (void) exception;
  status=MagickTrue;
  if (severity != 0)
  {
    if (device->type == CL_DEVICE_TYPE_CPU)
    {
      /* Workaround for Intel OpenCL CPU runtime bug */
      /* Turn off OpenCL when a problem is detected! */
      if (strncmp(device->platform_name,"Intel",5) == 0)
        default_CLEnv->enabled=MagickFalse;
    }
  }

#ifdef OPENCLLOG_ENABLED
  {
    va_list
      operands;
    va_start(operands,format);
    status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
      format,operands);
    va_end(operands);
  }
#else
  magick_unreferenced(module);
  magick_unreferenced(function);
  magick_unreferenced(line);
  magick_unreferenced(tag);
  magick_unreferenced(format);
#endif

  return(status);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   R e c o r d P r o f i l e D a t a                                         %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RecordProfileData() records profile data.
%
%  The format of the RecordProfileData method is:
%
%      void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
%        cl_event event)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device that did the operation.
%
%    o event: the event that contains the profiling data.
%
*/

MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
  cl_kernel kernel,cl_event event)
{
  char
    *name;

  cl_int
    status;

  cl_ulong
    elapsed,
    end,
    start;

  KernelProfileRecord
    profile_record;

  size_t
    i,
    length;

  if (device->profile_kernels == MagickFalse)
    return(MagickFalse);
  status=openCL_library->clWaitForEvents(1,&event);
  if (status != CL_SUCCESS)
    return(MagickFalse);
  status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
    &length);
  if (status != CL_SUCCESS)
    return(MagickTrue);
  name=AcquireQuantumMemory(length,sizeof(*name));
  if (name == (char *) NULL)
    return(MagickTrue);
  start=end=elapsed=0;
  status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
    name,(size_t *) NULL);
  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);
  if (status != CL_SUCCESS)
    {
      name=DestroyString(name);
      return(MagickTrue);
    }
  start/=1000; /* usecs */
  end/=1000;   
  elapsed=end-start;
  LockSemaphoreInfo(device->lock);
  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) == 0)
          {
            profile_record=device->profile_records[i];
            break;
          }
        i++;
      }
    }
  if (profile_record != (KernelProfileRecord) NULL)
    name=DestroyString(name);
  else
    {
      profile_record=AcquireCriticalMemory(sizeof(*profile_record));
      (void) memset(profile_record,0,sizeof(*profile_record));
      profile_record->kernel_name=name;
      device->profile_records=ResizeQuantumMemory(device->profile_records,(i+2),
        sizeof(*device->profile_records));
      if (device->profile_records == (KernelProfileRecord *) NULL)
        ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
      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);
  return(MagickTrue);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+  R e l e a s e O p e n C L C o m m a n d Q u e u e                          %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  ReleaseOpenCLCommandQueue() releases the OpenCL command queue
%
%  The format of the ReleaseOpenCLCommandQueue method is:
%
%      void ReleaseOpenCLCommandQueue(MagickCLDevice device,
%        cl_command_queue queue)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
%
%    o queue: the OpenCL queue to be released.
*/

MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
  cl_command_queue queue)
{
  if (queue == (cl_command_queue) NULL)
    return;

  assert(device != (MagickCLDevice) NULL);
  LockSemaphoreInfo(device->lock);
  if ((device->profile_kernels != MagickFalse) ||
      (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
    {
      UnlockSemaphoreInfo(device->lock);
      openCL_library->clFinish(queue);
      (void) openCL_library->clReleaseCommandQueue(queue);
    }
  else
    {
      openCL_library->clFlush(queue);
      device->command_queues[++device->command_queues_index]=queue;
      UnlockSemaphoreInfo(device->lock);
    }
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   R e l e a s e  M a g i c k C L D e v i c e                                %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  ReleaseOpenCLDevice() returns the OpenCL device to the environment
%
%  The format of the ReleaseOpenCLDevice method is:
%
%      void ReleaseOpenCLDevice(MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device to be released.
%
*/

MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
{
  assert(device != (MagickCLDevice) NULL);
  LockSemaphoreInfo(openCL_lock);
  device->requested--;
  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 DestroyMagickCLCacheInfoAndPixels(
  cl_event magick_unused(event),
  cl_int magick_unused(event_command_exec_status),void *user_data)
{
  MagickCLCacheInfo
    info;

  Quantum
    *pixels;

  ssize_t
    i;

  magick_unreferenced(event);
  magick_unreferenced(event_command_exec_status);
  info=(MagickCLCacheInfo) user_data;
  for (i=(ssize_t)info->event_count-1; i >= 0; i--)
  {
    cl_int
      event_status;

    cl_uint
      status;

    status=openCL_library->clGetEventInfo(info->events[i],
      CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(event_status),&event_status,
      NULL);
    if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
      {
        openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
          &DestroyMagickCLCacheInfoAndPixels,info);
        return;
      }
  }
  pixels=info->pixels;
  RelinquishMagickResource(MemoryResource,info->length);
  DestroyMagickCLCacheInfo(info);
  (void) RelinquishAlignedMemory(pixels);
}

MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
  MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
{
  if (info == (MagickCLCacheInfo) NULL)
    return((MagickCLCacheInfo) NULL);
  if (relinquish_pixels != MagickFalse)
    DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
  else
    DestroyMagickCLCacheInfo(info);
  return((MagickCLCacheInfo) NULL);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   R e l i n q u i s h M a g i c k C L D e v i c e                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RelinquishMagickCLDevice() releases the OpenCL device
%
%  The format of the RelinquishMagickCLDevice method is:
%
%      MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device to be released.
%
*/

static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
{
  if (device == (MagickCLDevice) NULL)
    return((MagickCLDevice) NULL);

  device->platform_name=RelinquishMagickMemory(device->platform_name);
  device->vendor_name=RelinquishMagickMemory(device->vendor_name);
  device->name=RelinquishMagickMemory(device->name);
  device->version=RelinquishMagickMemory(device->version);
  if (device->program != (cl_program) NULL)
    (void) openCL_library->clReleaseProgram(device->program);
  while (device->command_queues_index >= 0)
    (void) openCL_library->clReleaseCommandQueue(
      device->command_queues[device->command_queues_index--]);
  RelinquishSemaphoreInfo(&device->lock);
  return((MagickCLDevice) RelinquishMagickMemory(device));
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   R e l i n q u i s h M a g i c k C L E n v                                 %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RelinquishMagickCLEnv() releases the OpenCL environment
%
%  The format of the RelinquishMagickCLEnv method is:
%
%      MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment to be released.
%
*/

static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
{
  if (clEnv == (MagickCLEnv) NULL)
    return((MagickCLEnv) NULL);

  RelinquishSemaphoreInfo(&clEnv->lock);
  RelinquishMagickCLDevices(clEnv);
  if (clEnv->contexts != (cl_context *) NULL)
    {
      ssize_t
        i;

      for (i=0; i < clEnv->number_contexts; i++)
        if (clEnv->contexts[i] != (cl_context) NULL)
          (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
      clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
    }
  return((MagickCLEnv) RelinquishMagickMemory(clEnv));
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   R e q u e s t O p e n C L D e v i c e                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RequestOpenCLDevice() returns one of the enabled OpenCL devices.
%
%  The format of the RequestOpenCLDevice method is:
%
%      MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
*/

MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
{
  MagickCLDevice
    device;

  double
    score,
    best_score;

  size_t
    i;

  if (clEnv == (MagickCLEnv) NULL)
    return((MagickCLDevice) NULL);

  if (clEnv->number_devices == 1)
  {
    if (clEnv->devices[0]->enabled)
      return(clEnv->devices[0]);
    else
      return((MagickCLDevice) NULL);
  }

  device=(MagickCLDevice) NULL;
  best_score=0.0;
  LockSemaphoreInfo(openCL_lock);
  for (i = 0; i < clEnv->number_devices; i++)
  {
    if (clEnv->devices[i]->enabled == MagickFalse)
      continue;

    score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
      clEnv->devices[i]->requested);
    if ((device == (MagickCLDevice) NULL) || (score < best_score))
    {
      device=clEnv->devices[i];
      best_score=score;
    }
  }
  if (device != (MagickCLDevice)NULL)
    device->requested++;
  UnlockSemaphoreInfo(openCL_lock);

  return(device);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   S e t O p e n C L D e v i c e E n a b l e d                               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
%
%  The format of the SetOpenCLDeviceEnabled method is:
%
%      void SetOpenCLDeviceEnabled(MagickCLDevice device,
%        MagickBooleanType value)
%
%  A description of each parameter follows:
%
%    o device: the OpenCL device.
%
%    o value: determines if the device should be enabled or disabled.
*/

MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
  const MagickBooleanType value)
{
  if (device == (MagickCLDevice) NULL)
    return;
  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;
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%   S e t O p e n C L E n a b l e d                                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
%
%  The format of the SetOpenCLEnabled method is:
%
%      void SetOpenCLEnabled(MagickBooleanType)
%
%  A description of each parameter follows:
%
%    o value: specify true to enable OpenCL acceleration
*/

MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
{
  MagickCLEnv
    clEnv;

  clEnv=GetCurrentOpenCLEnv();
  if (clEnv == (MagickCLEnv) NULL)
    return(MagickFalse);
  clEnv->enabled=value;
  return(clEnv->enabled);
}

#else

MagickExport double GetOpenCLDeviceBenchmarkScore(
  const MagickCLDevice magick_unused(device))
{
  magick_unreferenced(device);
  return(0.0);
}

MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
  const MagickCLDevice magick_unused(device))
{
  magick_unreferenced(device);
  return(MagickFalse);
}

MagickExport const char *GetOpenCLDeviceName(
  const MagickCLDevice magick_unused(device))
{
  magick_unreferenced(device);
  return((const char *) NULL);
}

MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
  ExceptionInfo *magick_unused(exception))
{
  magick_unreferenced(exception);
  if (length != (size_t *) NULL)
    *length=0;
  return((MagickCLDevice *) NULL);
}

MagickExport MagickCLDeviceType GetOpenCLDeviceType(
  const MagickCLDevice magick_unused(device))
{
  magick_unreferenced(device);
  return(UndefinedCLDeviceType);
}

MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
  const MagickCLDevice magick_unused(device),size_t *length)
{
  magick_unreferenced(device);
  if (length != (size_t *) NULL)
    *length=0;
  return((const KernelProfileRecord *) NULL);
}

MagickExport const char *GetOpenCLDeviceVersion(
  const MagickCLDevice magick_unused(device))
{
  magick_unreferenced(device);
  return((const char *) NULL);
}

MagickExport MagickBooleanType GetOpenCLEnabled(void)
{
  return(MagickFalse);
}

MagickExport void SetOpenCLDeviceEnabled(
  MagickCLDevice magick_unused(device),
  const MagickBooleanType magick_unused(value))
{
  magick_unreferenced(device);
  magick_unreferenced(value);
}

MagickExport MagickBooleanType SetOpenCLEnabled(
  const MagickBooleanType magick_unused(value))
{
  magick_unreferenced(value);
  return(MagickFalse);
}

MagickExport void SetOpenCLKernelProfileEnabled(
  MagickCLDevice magick_unused(device),
  const MagickBooleanType magick_unused(value))
{
  magick_unreferenced(device);
  magick_unreferenced(value);
}
#endif