text: /*
  %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  %                                                                             %
  %                                                                             %
  %                                                                             %
  %                   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
  