root/magick/opencl.c

DEFINITIONS

1. startAccelerateTimer
2. stopAccelerateTimer
3. resetAccelerateTimer
4. initAccelerateTimer
5. readAccelerateTimer
6. RecordProfileData
7. DumpProfileData
8. OsLibraryLoad
9. OsLibraryGetFunctionAddress
10. OsLibraryUnload
11. AcquireMagickOpenCLEnv
12. RelinquishMagickOpenCLEnv
13. bindOpenCLFunctions
14. GetOpenCLLib
15. GetDefaultOpenCLEnv
16. LockDefaultOpenCLEnv
17. UnlockDefaultOpenCLEnv
18. SetDefaultOpenCLEnv
19. SetMagickOpenCLEnvParamInternal
20. SetMagickOpenCLEnvParam
21. GetMagickOpenCLEnvParam
22. GetOpenCLContext
23. getBinaryCLProgramName
24. saveBinaryCLProgram
25. loadBinaryCLProgram
26. stringSignature
27. CompileOpenCLKernels
28. InitOpenCLPlatformDevice
29. EnableOpenCLInternal
30. InitOpenCLEnvInternal
31. InitOpenCLEnv
32. AcquireOpenCLCommandQueue
33. RelinquishOpenCLCommandQueue
34. AcquireOpenCLKernel
35. RelinquishOpenCLKernel
36. GetOpenCLDeviceLocalMemorySize
37. GetOpenCLDeviceMaxMemAllocSize
38. releaseDeviceResource
39. releaseDSProfile
40. initDSProfile
41. profileDevices
42. writeProfileToFile
43. readProFile
44. findString
45. readProfileFromFile
46. getNumDeviceWithEmptyScore
47. AcceleratePerfEvaluator
48. AccelerateScoreSerializer
49. AccelerateScoreDeserializer
50. AccelerateScoreRelease
51. canWriteProfileToFile
52. autoSelectDevice
53. InitImageMagickOpenCL
54. OpenCLThrowMagickException
55. GetAndLockRandSeedBuffer
56. UnlockRandSeedBuffer
57. GetNumRandGenerators
58. GetRandNormalize
59. AcquireMagickOpenCLEnv
60. RelinquishMagickOpenCLEnv
61. GetDefaultOpenCLEnv
62. SetDefaultOpenCLEnv
63. SetMagickOpenCLEnvParam
64. GetMagickOpenCLEnvParam
65. InitOpenCLEnv
66. AcquireOpenCLCommandQueue
67. RelinquishCommandQueue
68. AcquireOpenCLKernel
69. RelinquishOpenCLKernel
70. GetOpenCLDeviceLocalMemorySize
71. InitImageMagickOpenCL
72. OpenCLThrowMagickException
73. GetAndLockRandSeedBuffer
74. UnlockRandSeedBuffer
75. GetNumRandGenerators
76. GetRandNormalize
77. GetOpenCLCachedFilesDirectory
78. OpenCLLog
79. OpenCLTerminus

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
%                   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-2016 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                                            %
%                                                                             %
%    http://www.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 "magick/studio.h"
#include "magick/artifact.h"
#include "magick/cache.h"
#include "magick/color.h"
#include "magick/compare.h"
#include "magick/constitute.h"
#include "magick/distort.h"
#include "magick/draw.h"
#include "magick/effect.h"
#include "magick/exception.h"
#include "magick/exception-private.h"
#include "magick/fx.h"
#include "magick/gem.h"
#include "magick/geometry.h"
#include "magick/image.h"
#include "magick/image-private.h"
#include "magick/layer.h"
#include "magick/mime-private.h"
#include "magick/memory_.h"
#include "magick/monitor.h"
#include "magick/montage.h"
#include "magick/morphology.h"
#include "magick/nt-base.h"
#include "magick/nt-base-private.h"
#include "magick/opencl.h"
#include "magick/opencl-private.h"
#include "magick/option.h"
#include "magick/policy.h"
#include "magick/property.h"
#include "magick/quantize.h"
#include "magick/quantum.h"
#include "magick/random_.h"
#include "magick/random-private.h"
#include "magick/resample.h"
#include "magick/resource_.h"
#include "magick/splay-tree.h"
#include "magick/semaphore.h"
#include "magick/statistic.h"
#include "magick/string_.h"
#include "magick/token.h"
#include "magick/utility.h"

#ifdef MAGICKCORE_CLPERFMARKER
#include "CLPerfMarker.h"
#endif


#if defined(MAGICKCORE_OPENCL_SUPPORT)

#ifdef MAGICKCORE_HAVE_OPENCL_CL_H
#define MAGICKCORE_OPENCL_MACOSX  1
#endif

#define NUM_CL_RAND_GENERATORS 1024  /* number of random number generators running in parallel */
#define PROFILE_OCL_KERNELS 0

typedef struct
{
  cl_ulong min;
  cl_ulong max;
  cl_ulong total;
  cl_ulong count;
} KernelProfileRecord;

static const char *kernelNames[] = {
  "AddNoise",
  "BlurRow",
  "BlurColumn",
  "Composite",
  "ComputeFunction",
  "Contrast",
  "ContrastStretch",
  "Convolve",
  "Equalize",
  "GrayScale",
  "Histogram",
  "HullPass1",
  "HullPass2",
  "LocalContrastBlurRow",
  "LocalContrastBlurApplyColumn",
  "Modulate",
  "MotionBlur",
  "RadialBlur",
  "RandomNumberGenerator",
  "ResizeHorizontal",
  "ResizeVertical",
  "UnsharpMaskBlurColumn",
  "UnsharpMask",
  "NONE" };

KernelProfileRecord
  profileRecords[KERNEL_COUNT];

typedef struct _AccelerateTimer {
  long long _freq;
  long long _clocks;
  long long _start;
} AccelerateTimer;

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
}

void stopAccelerateTimer(AccelerateTimer* timer) {
      long long 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;
}

void resetAccelerateTimer(AccelerateTimer* timer) {
   timer->_clocks = 0;
   timer->_start = 0;
}

void initAccelerateTimer(AccelerateTimer* timer) {
#ifdef _WIN32
    QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq);
#else
    timer->_freq = (long long)1.0E3;
#endif
   resetAccelerateTimer(timer);
}

double readAccelerateTimer(AccelerateTimer* timer) {
  return (double)timer->_clocks/(double)timer->_freq;
};

MagickPrivate void RecordProfileData(MagickCLEnv clEnv, ProfiledKernels kernel, cl_event event)
{
#if PROFILE_OCL_KERNELS
  cl_int status;
  cl_ulong start = 0;
  cl_ulong end = 0;
  cl_ulong elapsed = 0;
  clEnv->library->clWaitForEvents(1, &event);
  status = clEnv->library->clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
  status &= clEnv->library->clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
  if (status == CL_SUCCESS) {
    start /= 1000;      // usecs
    end /= 1000;        // usecs
    elapsed = end - start;
    /* we can use the commandQueuesLock to make the code below thread safe */
    LockSemaphoreInfo(clEnv->commandQueuesLock);
    if ((elapsed < profileRecords[kernel].min) || (profileRecords[kernel].count == 0))
      profileRecords[kernel].min = elapsed;
    if (elapsed > profileRecords[kernel].max)
      profileRecords[kernel].max = elapsed;
    profileRecords[kernel].total += elapsed;
    profileRecords[kernel].count += 1;
    UnlockSemaphoreInfo(clEnv->commandQueuesLock);
  }
#endif
}

void DumpProfileData()
{
#if PROFILE_OCL_KERNELS
  int i;

  OpenCLLog("====================================================");

  // Write out the device info to the profile
  if (0 == 1)
  {
    MagickCLEnv clEnv;
    char buff[2048];
    cl_int status;

    clEnv = GetDefaultOpenCLEnv();

    status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_VENDOR, 2048, buff, NULL);
    OpenCLLog(buff);

    status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, 2048, buff, NULL);
    OpenCLLog(buff);

    status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DRIVER_VERSION, 2048, buff, NULL);
    OpenCLLog(buff);
  }

  OpenCLLog("====================================================");
  OpenCLLog("                        ave\tcalls   \tmin -> max");
  OpenCLLog("                        ---\t-----   \t----------");
  for (i = 0; i < KERNEL_COUNT; ++i) {
    char buf[4096];
    char indent[160];
    strcpy(indent, "                              ");
    strncpy(indent, kernelNames[i], min(strlen(kernelNames[i]), strlen(indent) - 1));
    sprintf(buf, "%s%d\t(%d calls)   \t%d -> %d", indent, profileRecords[i].count > 0 ? (profileRecords[i].total / profileRecords[i].count) : 0, profileRecords[i].count, profileRecords[i].min, profileRecords[i].max);
    //printf("%s%d\t(%d calls)   \t%d -> %d\n", indent, profileRecords[i].count > 0 ? (profileRecords[i].total / profileRecords[i].count) : 0, profileRecords[i].count, profileRecords[i].min, profileRecords[i].max);
    OpenCLLog(buf);
  }
  OpenCLLog("====================================================");
#endif
}

/*
 *
 * Dynamic library loading functions
 *
 */
#ifdef MAGICKCORE_WINDOWS_SUPPORT
#else
#include <dlfcn.h>
#endif

// dynamically load a library.  returns NULL on failure
void *OsLibraryLoad(const char *libraryName)
{
#ifdef MAGICKCORE_WINDOWS_SUPPORT
    return (void *)LoadLibraryA(libraryName);
#else
    return (void *)dlopen(libraryName, RTLD_NOW);
#endif
}

// get a function pointer from a loaded library.  returns NULL on failure.
void *OsLibraryGetFunctionAddress(void *library, const char *functionName)
{
#ifdef MAGICKCORE_WINDOWS_SUPPORT
    if (!library || !functionName)
    {
        return NULL;
    }
    return (void *) GetProcAddress( (HMODULE)library, functionName);
#else
    if (!library || !functionName)
    {
        return NULL;
    }
    return (void *)dlsym(library, functionName);
#endif
}

// unload a library.
void OsLibraryUnload(void *library)
{
#ifdef MAGICKCORE_WINDOWS_SUPPORT
    FreeLibrary( (HMODULE)library);
#else
    dlclose(library);
#endif
}


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

MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
{
  MagickCLEnv clEnv;
  clEnv = (MagickCLEnv) AcquireMagickMemory(sizeof(struct _MagickCLEnv));
  if (clEnv != NULL)
  {
    memset(clEnv, 0, sizeof(struct _MagickCLEnv));
    clEnv->commandQueuesPos=-1;
    ActivateSemaphoreInfo(&clEnv->lock);
    ActivateSemaphoreInfo(&clEnv->commandQueuesLock);
  }
  return clEnv;
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   R e l i n q u i s h M a g i c k O p e n C L E n v                         %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RelinquishMagickOpenCLEnv() destroy the MagickCLEnv structure
%
%  The format of the RelinquishMagickOpenCLEnv method is:
%
%      MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv)
%
%  A description of each parameter follows:
%
%    o clEnv: MagickCLEnv structure to destroy
%
*/

MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv)
{
  if (clEnv != (MagickCLEnv) NULL)
  {
    while (clEnv->commandQueuesPos >= 0)
    {
      clEnv->library->clReleaseCommandQueue(
        clEnv->commandQueues[clEnv->commandQueuesPos--]);
    }
    DestroySemaphoreInfo(&clEnv->lock);
    DestroySemaphoreInfo(&clEnv->commandQueuesLock);
    RelinquishMagickMemory(clEnv);
    return MagickTrue;
  }
  return MagickFalse;
}


/*
* Default OpenCL environment
*/
MagickCLEnv defaultCLEnv;
SemaphoreInfo* defaultCLEnvLock;

/*
* OpenCL library
*/
MagickLibrary * OpenCLLib;
SemaphoreInfo* OpenCLLibLock;


static MagickBooleanType bindOpenCLFunctions(void* library)
{
#ifdef MAGICKCORE_OPENCL_MACOSX
#define BIND(X) OpenCLLib->X= &X;
#else
#define BIND(X)\
  if ((OpenCLLib->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(library,#X)) == NULL)\
  return MagickFalse;
#endif

  BIND(clGetPlatformIDs);
  BIND(clGetPlatformInfo);

  BIND(clGetDeviceIDs);
  BIND(clGetDeviceInfo);

  BIND(clCreateContext);

  BIND(clCreateBuffer);
  BIND(clReleaseMemObject);

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

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

  BIND(clFlush);
  BIND(clFinish);

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

  BIND(clCreateCommandQueue);
  BIND(clReleaseCommandQueue);

  BIND(clGetEventProfilingInfo);
  BIND(clWaitForEvents);
  BIND(clReleaseEvent);

  return MagickTrue;
}

MagickLibrary * GetOpenCLLib()
{
  if (OpenCLLib == NULL)
  {
    if (OpenCLLibLock == NULL)
    {
      ActivateSemaphoreInfo(&OpenCLLibLock);
    }

    LockSemaphoreInfo(OpenCLLibLock);

    OpenCLLib = (MagickLibrary *) AcquireMagickMemory (sizeof (MagickLibrary));

    if (OpenCLLib != NULL)
    {
      MagickBooleanType status = MagickFalse;
      void * library = NULL;

#ifdef MAGICKCORE_OPENCL_MACOSX
      status = bindOpenCLFunctions(library);
#else

      memset(OpenCLLib, 0, sizeof(MagickLibrary));
#ifdef MAGICKCORE_WINDOWS_SUPPORT
      library = OsLibraryLoad("OpenCL.dll");
#else
      library = OsLibraryLoad("libOpenCL.so");
#endif
      if (library)
        status = bindOpenCLFunctions(library);

      if (status==MagickTrue)
        OpenCLLib->base=library;
      else
        OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
#endif
    }

    UnlockSemaphoreInfo(OpenCLLibLock);
  }


  return OpenCLLib;
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   G e t D e f a u l t O p e n C L E n v                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetDefaultOpenCLEnv() returns the default OpenCL env
%
%  The format of the GetDefaultOpenCLEnv method is:
%
%      MagickCLEnv GetDefaultOpenCLEnv()
%
%  A description of each parameter follows:
%
%    o exception: return any errors or warnings.
%
*/

MagickExport MagickCLEnv GetDefaultOpenCLEnv()
{
  if (defaultCLEnv == NULL)
  {
    if (defaultCLEnvLock == NULL)
    {
      ActivateSemaphoreInfo(&defaultCLEnvLock);
    }
    LockSemaphoreInfo(defaultCLEnvLock);
    if (defaultCLEnv == NULL)
      defaultCLEnv = AcquireMagickOpenCLEnv();
    UnlockSemaphoreInfo(defaultCLEnvLock);
  }
  return defaultCLEnv;
}

static void LockDefaultOpenCLEnv() {
  if (defaultCLEnvLock == NULL)
  {
    ActivateSemaphoreInfo(&defaultCLEnvLock);
  }
  LockSemaphoreInfo(defaultCLEnvLock);
}

static void UnlockDefaultOpenCLEnv() {
  if (defaultCLEnvLock == NULL)
  {
    ActivateSemaphoreInfo(&defaultCLEnvLock);
  }
  else
    UnlockSemaphoreInfo(defaultCLEnvLock);
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   S e t D e f a u l t O p e n C L E n v                                     %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  SetDefaultOpenCLEnv() sets the new OpenCL environment as default
%  and returns the old OpenCL environment
%
%  The format of the SetDefaultOpenCLEnv() method is:
%
%      MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv)
%
%  A description of each parameter follows:
%
%    o clEnv: the new default OpenCL environment.
%
*/
MagickExport MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv)
{
  MagickCLEnv oldEnv;
  LockDefaultOpenCLEnv();
  oldEnv = defaultCLEnv;
  defaultCLEnv = clEnv;
  UnlockDefaultOpenCLEnv();
  return oldEnv;
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   S e t M a g i c k O p e n C L E n v P a r a m                             %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  SetMagickOpenCLEnvParam() sets the parameters in the OpenCL environment
%
%  The format of the SetMagickOpenCLEnvParam() method is:
%
%      MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv,
%        MagickOpenCLEnvParam param, size_t dataSize, void* data,
%        ExceptionInfo* exception)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o param: the parameter to be set.
%
%    o dataSize: the data size of the parameter value.
%
%    o data:  the pointer to the new parameter value
%
%    o exception: return any errors or warnings
%
*/

static MagickBooleanType SetMagickOpenCLEnvParamInternal(MagickCLEnv clEnv, MagickOpenCLEnvParam param
                                          , size_t dataSize, void* data, ExceptionInfo* exception)
{
  MagickBooleanType status = MagickFalse;

  if (clEnv == NULL
    || data == NULL)
    goto cleanup;

  switch(param)
  {
  case MAGICK_OPENCL_ENV_PARAM_DEVICE:
    if (dataSize != sizeof(clEnv->device))
      goto cleanup;
    clEnv->device = *((cl_device_id*)data);
    clEnv->OpenCLInitialized = MagickFalse;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED:
    if (dataSize != sizeof(clEnv->OpenCLDisabled))
      goto cleanup;
    clEnv->OpenCLDisabled =  *((MagickBooleanType*)data);
    clEnv->OpenCLInitialized = MagickFalse;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED:
    (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "SetMagickOpenCLEnvParm cannot modify the OpenCL initialization state.", "'%s'", ".");
    break;

  case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
    if (dataSize != sizeof(clEnv->disableProgramCache))
      goto cleanup;
    clEnv->disableProgramCache =  *((MagickBooleanType*)data);
    clEnv->OpenCLInitialized = MagickFalse;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
    if (dataSize != sizeof(clEnv->regenerateProfile))
      goto cleanup;
    clEnv->regenerateProfile =  *((MagickBooleanType*)data);
    clEnv->OpenCLInitialized = MagickFalse;
    status = MagickTrue;
    break;

  default:
    goto cleanup;
  };

cleanup:
  return status;
}

MagickExport
  MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
                                          , size_t dataSize, void* data, ExceptionInfo* exception) {
  MagickBooleanType status = MagickFalse;
  if (clEnv!=NULL) {
    LockSemaphoreInfo(clEnv->lock);
    status = SetMagickOpenCLEnvParamInternal(clEnv,param,dataSize,data,exception);
    UnlockSemaphoreInfo(clEnv->lock);
  }
  return status;
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   G e t M a g i c k O p e n C L E n v P a r a m                             %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetMagickOpenCLEnvParam() gets the parameters in the OpenCL environment
%
%  The format of the GetMagickOpenCLEnvParam() method is:
%
%      MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv,
%        MagickOpenCLEnvParam param, size_t dataSize, void* data,
%        ExceptionInfo* exception)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o param: the parameter to be returned.
%
%    o dataSize: the data size of the parameter value.
%
%    o data:  the location where the returned parameter value will be stored
%
%    o exception: return any errors or warnings
%
*/

MagickExport
  MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
                                          , size_t dataSize, void* data, ExceptionInfo* exception)
{
  MagickBooleanType
   status;

  magick_unreferenced(exception);

  status = MagickFalse;

  if (clEnv == NULL
    || data == NULL)
    goto cleanup;

  switch(param)
  {
  case MAGICK_OPENCL_ENV_PARAM_DEVICE:
    if (dataSize != sizeof(cl_device_id))
      goto cleanup;
    *((cl_device_id*)data) = clEnv->device;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED:
    if (dataSize != sizeof(clEnv->OpenCLDisabled))
      goto cleanup;
    *((MagickBooleanType*)data) = clEnv->OpenCLDisabled;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED:
    if (dataSize != sizeof(clEnv->OpenCLDisabled))
      goto cleanup;
    *((MagickBooleanType*)data) = clEnv->OpenCLInitialized;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
    if (dataSize != sizeof(clEnv->disableProgramCache))
      goto cleanup;
    *((MagickBooleanType*)data) = clEnv->disableProgramCache;
    status = MagickTrue;
    break;

  case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
    if (dataSize != sizeof(clEnv->regenerateProfile))
      goto cleanup;
    *((MagickBooleanType*)data) = clEnv->regenerateProfile;
    status = MagickTrue;
    break;

  default:
    goto cleanup;
  };

cleanup:
  return status;
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   G e t O p e n C L C o n t e x t                                           %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLContext() returns the OpenCL context
%
%  The format of the GetOpenCLContext() method is:
%
%      cl_context GetOpenCLContext(MagickCLEnv clEnv)
%
%  A description of each parameter follows:
%
%    o clEnv: OpenCL environment
%
*/

MagickPrivate
cl_context GetOpenCLContext(MagickCLEnv clEnv) {
  if (clEnv == NULL)
    return NULL;
  else
    return clEnv->context;
}

static char* getBinaryCLProgramName(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
{
  char* name;
  char* ptr;
  char path[MaxTextExtent];
  char deviceName[MaxTextExtent];
  const char* prefix = "magick_opencl";
  clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MaxTextExtent, deviceName, NULL);
  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(path,MaxTextExtent,"%s%s%s_%s_%02d_%08x_%.20g.bin",
         GetOpenCLCachedFilesDirectory(),DirectorySeparator,prefix,deviceName,
         (unsigned int) prog,signature,(double) sizeof(char*)*8);
  name = (char*)AcquireMagickMemory(strlen(path)+1);
  CopyMagickString(name,path,strlen(path)+1);
  return name;
}

static MagickBooleanType saveBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature, ExceptionInfo* exception)
{
  MagickBooleanType saveSuccessful;
  cl_int clStatus;
  size_t binaryProgramSize;
  unsigned char* binaryProgram;
  char* binaryFileName;
  FILE* fileHandle;

#ifdef MAGICKCORE_CLPERFMARKER
  clBeginPerfMarkerAMD(__FUNCTION__,"");
#endif

  binaryProgram = NULL;
  binaryFileName = NULL;
  fileHandle = NULL;
  saveSuccessful = MagickFalse;

  clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binaryProgramSize, NULL);
  if (clStatus != CL_SUCCESS)
  {
    (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
    goto cleanup;
  }

  binaryProgram = (unsigned char*) AcquireMagickMemory(binaryProgramSize);
  clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARIES, sizeof(char*), &binaryProgram, NULL);
  if (clStatus != CL_SUCCESS)
  {
    (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
    goto cleanup;
  }

  binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
  fileHandle = fopen(binaryFileName, "wb");
  if (fileHandle != NULL)
  {
    fwrite(binaryProgram, sizeof(char), binaryProgramSize, fileHandle);
    saveSuccessful = MagickTrue;
  }
  else
  {
    (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
      "Saving binary kernel failed.", "'%s'", ".");
  }

cleanup:
  if (fileHandle != NULL)
    fclose(fileHandle);
  if (binaryProgram != NULL)
    RelinquishMagickMemory(binaryProgram);
  if (binaryFileName != NULL)
    free(binaryFileName);

#ifdef MAGICKCORE_CLPERFMARKER
  clEndPerfMarkerAMD();
#endif

  return saveSuccessful;
}

static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
{
  MagickBooleanType loadSuccessful;
  unsigned char* binaryProgram;
  char* binaryFileName;
  FILE* fileHandle;

#ifdef MAGICKCORE_CLPERFMARKER
  clBeginPerfMarkerAMD(__FUNCTION__,"");
#endif

  binaryProgram = NULL;
  binaryFileName = NULL;
  fileHandle = NULL;
  loadSuccessful = MagickFalse;

  binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
  fileHandle = fopen(binaryFileName, "rb");
  if (fileHandle != NULL)
  {
    int b_error;
    size_t length;
    cl_int clStatus;
    cl_int clBinaryStatus;

    b_error = 0 ;
    length = 0;
    b_error |= fseek( fileHandle, 0, SEEK_END ) < 0;
    b_error |= ( length = ftell( fileHandle ) ) <= 0;
    b_error |= fseek( fileHandle, 0, SEEK_SET ) < 0;
    if( b_error )
      goto cleanup;

    binaryProgram = (unsigned char*)AcquireMagickMemory(length);
    if (binaryProgram == NULL)
      goto cleanup;

    memset(binaryProgram, 0, length);
    b_error |= fread(binaryProgram, 1, length, fileHandle) != length;

    clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
    if (clStatus != CL_SUCCESS
        || clBinaryStatus != CL_SUCCESS)
      goto cleanup;

    loadSuccessful = MagickTrue;
  }

cleanup:
  if (fileHandle != NULL)
    fclose(fileHandle);
  if (binaryFileName != NULL)
    free(binaryFileName);
  if (binaryProgram != NULL)
    RelinquishMagickMemory(binaryProgram);

#ifdef MAGICKCORE_CLPERFMARKER
  clEndPerfMarkerAMD();
#endif

  return loadSuccessful;
}

static unsigned int stringSignature(const char* string)
{
  unsigned int stringLength;
  unsigned int n,i,j;
  unsigned int signature;
  union
  {
    const char* s;
    const unsigned int* u;
  }p;

#ifdef MAGICKCORE_CLPERFMARKER
  clBeginPerfMarkerAMD(__FUNCTION__,"");
#endif

  stringLength = (unsigned int) strlen(string);
  signature = stringLength;
  n = stringLength/sizeof(unsigned int);
  p.s = string;
  for (i = 0; i < n; i++)
  {
    signature^=p.u[i];
  }
  if (n * sizeof(unsigned int) != stringLength)
  {
    char padded[4];
    j = n * sizeof(unsigned int);
    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];
  }

#ifdef MAGICKCORE_CLPERFMARKER
  clEndPerfMarkerAMD();
#endif

  return signature;
}

/* OpenCL kernels for accelerate.c */
extern const char *accelerateKernels, *accelerateKernels2;

static MagickBooleanType CompileOpenCLKernels(MagickCLEnv clEnv, ExceptionInfo* exception)
{
  MagickBooleanType status = MagickFalse;
  cl_int clStatus;
  unsigned int i;
  char* accelerateKernelsBuffer = NULL;

  /* The index of the program strings in this array has to match the value of the enum MagickOpenCLProgram */
  const char* MagickOpenCLProgramStrings[MAGICK_OPENCL_NUM_PROGRAMS];

  char options[MaxTextExtent];
  unsigned int optionsSignature;

#ifdef MAGICKCORE_CLPERFMARKER
  clBeginPerfMarkerAMD(__FUNCTION__,"");
#endif

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

  /*
  if (getenv("MAGICK_OCL_DEF"))
  {
    strcat(options," ");
    strcat(options,getenv("MAGICK_OCL_DEF"));
  }
  */

  /*
  if (getenv("MAGICK_OCL_BUILD"))
    printf("options: %s\n", options);
  */

  optionsSignature = stringSignature(options);

  /* get all the OpenCL program strings here */
  accelerateKernelsBuffer = (char*) AcquireMagickMemory(strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
  sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
  MagickOpenCLProgramStrings[MAGICK_OPENCL_ACCELERATE] = accelerateKernelsBuffer;

  for (i = 0; i < MAGICK_OPENCL_NUM_PROGRAMS; i++)
  {
    MagickBooleanType loadSuccessful = MagickFalse;
    unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature;

    /* try to load the binary first */
    if (clEnv->disableProgramCache != MagickTrue
        && !getenv("MAGICK_OCL_REC"))
      loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature);

    if (loadSuccessful == MagickFalse)
    {
      /* Binary CL program unavailable, compile the program from source */
      size_t programLength = strlen(MagickOpenCLProgramStrings[i]);
      clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
      if (clStatus!=CL_SUCCESS)
      {
        (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
          "clCreateProgramWithSource failed.", "(%d)", (int)clStatus);

        goto cleanup;
      }
    }

    clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
    if (clStatus!=CL_SUCCESS)
    {
      (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
        "clBuildProgram failed.", "(%d)", (int)clStatus);

      if (loadSuccessful == MagickFalse)
      {
        char path[MaxTextExtent];
        FILE* fileHandle;

        /*  dump the source into a file */
        (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
         ,GetOpenCLCachedFilesDirectory()
         ,DirectorySeparator,"magick_badcl.cl");
        fileHandle = fopen(path, "wb"); 
        if (fileHandle != NULL)
        {
          fwrite(MagickOpenCLProgramStrings[i], sizeof(char), strlen(MagickOpenCLProgramStrings[i]), fileHandle);
          fclose(fileHandle);
        }

        /* dump the build log */
        {
          char* log;
          size_t logSize;
          clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
          log = (char*)AcquireMagickMemory(logSize);
          clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);

          (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
           ,GetOpenCLCachedFilesDirectory()
           ,DirectorySeparator,"magick_badcl_build.log");
          fileHandle = fopen(path, "wb");       
          if (fileHandle != NULL)
          {
            const char* buildOptionsTitle = "build options: ";
            fwrite(buildOptionsTitle, sizeof(char), strlen(buildOptionsTitle), fileHandle);
            fwrite(options, sizeof(char), strlen(options), fileHandle);
            fwrite("\n",sizeof(char), 1, fileHandle);
            fwrite(log, sizeof(char), logSize, fileHandle);
            fclose(fileHandle);
          }
          RelinquishMagickMemory(log);
        }
      }
      goto cleanup;
    }

    if (loadSuccessful == MagickFalse)
    {
      /* Save the binary to a file to avoid re-compilation of the kernels in the future */
      saveBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception);
    }

  }
  status = MagickTrue;

cleanup:

  if (accelerateKernelsBuffer!=NULL) RelinquishMagickMemory(accelerateKernelsBuffer);

#ifdef MAGICKCORE_CLPERFMARKER
  clEndPerfMarkerAMD();
#endif

  return status;
}

static MagickBooleanType InitOpenCLPlatformDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
  int i,j;
  cl_int status;
  cl_uint numPlatforms = 0;
  cl_platform_id *platforms = NULL;
  char* MAGICK_OCL_DEVICE = NULL;
  MagickBooleanType OpenCLAvailable = MagickFalse;

#ifdef MAGICKCORE_CLPERFMARKER
  clBeginPerfMarkerAMD(__FUNCTION__,"");
#endif

  /* check if there's an environment variable overriding the device selection */
  MAGICK_OCL_DEVICE = getenv("MAGICK_OCL_DEVICE");
  if (MAGICK_OCL_DEVICE != NULL)
  {
    if (strcmp(MAGICK_OCL_DEVICE, "CPU") == 0)
    {
      clEnv->deviceType = CL_DEVICE_TYPE_CPU;
    }
    else if (strcmp(MAGICK_OCL_DEVICE, "GPU") == 0)
    {
      clEnv->deviceType = CL_DEVICE_TYPE_GPU;
    }
    else if (strcmp(MAGICK_OCL_DEVICE, "OFF") == 0)
    {
      /* OpenCL disabled */
      goto cleanup;
    }
  }
  else if (clEnv->deviceType == 0) {
    clEnv->deviceType = CL_DEVICE_TYPE_ALL;
  }

  if (clEnv->device != NULL)
  {
    status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
    if (status != CL_SUCCESS) {
      (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
          "Failed to get OpenCL platform from the selected device.", "(%d)", status);
    }
    goto cleanup;
  }
  else if (clEnv->platform != NULL)
  {
    numPlatforms = 1;
    platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id));
    if (platforms == (cl_platform_id *) NULL)
    {
      (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
        "AcquireMagickMemory failed.",".");
      goto cleanup;
    }
    platforms[0] = clEnv->platform;
  }
  else
  {
    clEnv->device = NULL;

    /* Get the number of OpenCL platforms available */
    status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms);
    if (status != CL_SUCCESS)
    {
      (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
        "clGetplatformIDs failed.", "(%d)", status);
      goto cleanup;
    }

    /* No OpenCL available, just leave */
    if (numPlatforms == 0) {
      goto cleanup;
    }

    platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id));
    if (platforms == (cl_platform_id *) NULL)
    {
      (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
        "AcquireMagickMemory failed.",".");
      goto cleanup;
    }

    status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL);
    if (status != CL_SUCCESS)
    {
      (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
        "clGetPlatformIDs failed.", "(%d)", status);
      goto cleanup;
    }
  }

  /* Device selection */
  clEnv->device = NULL;
  for (j = 0; j < 2; j++)
  {

    cl_device_type deviceType;
    if (clEnv->deviceType == CL_DEVICE_TYPE_ALL)
    {
      if (j == 0)
        deviceType = CL_DEVICE_TYPE_GPU;
      else
        deviceType = CL_DEVICE_TYPE_CPU;
    }
    else if (j == 1)
    {
      break;
    }
    else
      deviceType = clEnv->deviceType;

    for (i = 0; i < numPlatforms; i++)
    {
      char version[MaxTextExtent];
      cl_uint numDevices;
      status = clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_VERSION, MaxTextExtent, version, NULL);
      if (status != CL_SUCCESS)
      {
        (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
          "clGetPlatformInfo failed.", "(%d)", status);
        goto cleanup;
      }
      if (strncmp(version,"OpenCL 1.0 ",11) == 0)
        continue;
      status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
      if (status != CL_SUCCESS)
      {
        (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
          "clGetDeviceIDs failed.", "(%d)", status);
        goto cleanup;
      }
      if (clEnv->device != NULL)
      {
        clEnv->platform = platforms[i];
  goto cleanup;
      }
    }
  }

cleanup:
  if (platforms!=NULL)
    RelinquishMagickMemory(platforms);

  OpenCLAvailable = (clEnv->platform!=NULL
          && clEnv->device!=NULL)?MagickTrue:MagickFalse;

#ifdef MAGICKCORE_CLPERFMARKER
  clEndPerfMarkerAMD();
#endif

  return OpenCLAvailable;
}

static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) {
  if (clEnv->OpenCLInitialized != MagickFalse
    && clEnv->platform != NULL
    && clEnv->device != NULL) {
      clEnv->OpenCLDisabled = MagickFalse;
      return MagickTrue;
  }
  clEnv->OpenCLDisabled = MagickTrue;
  return MagickFalse;
}


static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception);
/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   I n i t O p e n C L E n v                                                 %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  InitOpenCLEnv() initialize the OpenCL environment
%
%  The format of the RelinquishMagickOpenCLEnv method is:
%
%      MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception)
%
%  A description of each parameter follows:
%
%    o clEnv: OpenCL environment structure
%
%    o exception: return any errors or warnings.
%
*/

MagickExport
MagickBooleanType InitOpenCLEnvInternal(MagickCLEnv clEnv, ExceptionInfo* exception) {
  MagickBooleanType status = MagickTrue;
  cl_int clStatus;
  cl_context_properties cps[3];

#ifdef MAGICKCORE_CLPERFMARKER
  {
    int status = clInitializePerfMarkerAMD();
    if (status == AP_SUCCESS) {
      //printf("PerfMarker successfully initialized\n");
    }
  }
#endif
  clEnv->OpenCLInitialized = MagickTrue;

  /* check and init the global lib */
  OpenCLLib=GetOpenCLLib();
  if (OpenCLLib)
  {
    clEnv->library=OpenCLLib;
  }
  else
  {
    /* turn off opencl */
    MagickBooleanType flag;
    flag = MagickTrue;
    SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
        , sizeof(MagickBooleanType), &flag, exception);
  }

  if (clEnv->OpenCLDisabled != MagickFalse)
    goto cleanup;

  clEnv->OpenCLDisabled = MagickTrue;
  /* setup the OpenCL platform and device */
  status = InitOpenCLPlatformDevice(clEnv, exception);
  if (status == MagickFalse) {
    /* No OpenCL device available */
    goto cleanup;
  }

  /* create an OpenCL context */
  cps[0] = CL_CONTEXT_PLATFORM;
  cps[1] = (cl_context_properties)clEnv->platform;
  cps[2] = 0;
  clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
  if (clStatus != CL_SUCCESS)
  {
    (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
        "clCreateContext failed.", "(%d)", clStatus);
    status = MagickFalse;
    goto cleanup;
  }

  status = CompileOpenCLKernels(clEnv, exception);
  if (status == MagickFalse) {
   (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
        "clCreateCommandQueue failed.", "(%d)", status);

    status = MagickFalse;
    goto cleanup;
  }

  status = EnableOpenCLInternal(clEnv);

cleanup:
  return status;
}


MagickExport
MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) {
  MagickBooleanType status = MagickFalse;

  if (clEnv == NULL)
    return MagickFalse;

#ifdef MAGICKCORE_CLPERFMARKER
  clBeginPerfMarkerAMD(__FUNCTION__,"");
#endif

  LockSemaphoreInfo(clEnv->lock);
  if (clEnv->OpenCLInitialized == MagickFalse) {
    if (clEnv->device==NULL
        && clEnv->OpenCLDisabled == MagickFalse)
      status = autoSelectDevice(clEnv, exception);
    else
      status = InitOpenCLEnvInternal(clEnv, exception);
  }
  UnlockSemaphoreInfo(clEnv->lock);

#ifdef MAGICKCORE_CLPERFMARKER
  clEndPerfMarkerAMD();
#endif
  return status;
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   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(MagickCLEnv clEnv)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
*/

MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
{
  cl_command_queue
    queue;

  cl_command_queue_properties
    properties;

  if (clEnv == (MagickCLEnv) NULL)
    return (cl_command_queue) NULL;
  LockSemaphoreInfo(clEnv->commandQueuesLock);
  if (clEnv->commandQueuesPos >= 0) {
    queue=clEnv->commandQueues[clEnv->commandQueuesPos--];
    UnlockSemaphoreInfo(clEnv->commandQueuesLock);
  }
  else {
    UnlockSemaphoreInfo(clEnv->commandQueuesLock);
    properties=0;
#if PROFILE_OCL_KERNELS
    properties=CL_QUEUE_PROFILING_ENABLE;
#endif
    queue=clEnv->library->clCreateCommandQueue(clEnv->context,clEnv->device,
      properties,NULL);
  }
  return(queue);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   R e l i n q u i s h O p e n C L C o m m a n d Q u e u e                   %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RelinquishOpenCLCommandQueue() releases the OpenCL command queue
%
%  The format of the RelinquishOpenCLCommandQueue method is:
%
%      MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
%        cl_command_queue queue)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o queue: the OpenCL queue to be released.
%
%
*/

MagickPrivate MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
  cl_command_queue queue)
{
  MagickBooleanType
    status;

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

  LockSemaphoreInfo(clEnv->commandQueuesLock);

  if (clEnv->commandQueuesPos >= MAX_COMMAND_QUEUES)
    status=(clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ?
      MagickTrue : MagickFalse;
  else
    {
      clEnv->commandQueues[++clEnv->commandQueuesPos]=queue;
      status=MagickTrue;
    }

  UnlockSemaphoreInfo(clEnv->commandQueuesLock);

  return(status);
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   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(MagickCLEnv clEnv, MagickOpenCLProgram program, const char* kernelName)
{
  cl_int clStatus;
  cl_kernel kernel = NULL;
  if (clEnv != NULL && kernelName!=NULL)
  {
    kernel = clEnv->library->clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
  }
  return kernel;
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   R e l i n q u i s h O p e n C L K e r n e l                               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  RelinquishOpenCLKernel() releases an OpenCL kernel
%
%  The format of the RelinquishOpenCLKernel method is:
%
%    MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv,
%      cl_kernel kernel)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%    o kernel: the OpenCL kernel object to be released.
%
%
*/

MagickPrivate
  MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel)
{
  MagickBooleanType status = MagickFalse;
  if (clEnv != NULL && kernel != NULL)
  {
    status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
  }
  return status;
}

/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   G e t O p e n C L D e v i c e L o c a l M e m o r y S i z e               %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  GetOpenCLDeviceLocalMemorySize() returns local memory size of the device
%
%  The format of the GetOpenCLDeviceLocalMemorySize method is:
%
%    unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
%
%  A description of each parameter follows:
%
%    o clEnv: the OpenCL environment.
%
%
*/

MagickPrivate
 unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
{
  cl_ulong localMemorySize;
  clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
  return (unsigned long)localMemorySize;
}

MagickPrivate
  unsigned long GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv clEnv)
{
  cl_ulong maxMemAllocSize;
  clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
  return (unsigned long)maxMemAllocSize;
}


/*
 Beginning of the OpenCL device selection infrastructure
*/


typedef enum {
  DS_SUCCESS = 0
 ,DS_INVALID_PROFILE = 1000
 ,DS_MEMORY_ERROR
 ,DS_INVALID_PERF_EVALUATOR_TYPE
 ,DS_INVALID_PERF_EVALUATOR
 ,DS_PERF_EVALUATOR_ERROR
 ,DS_FILE_ERROR
 ,DS_UNKNOWN_DEVICE_TYPE
 ,DS_PROFILE_FILE_ERROR
 ,DS_SCORE_SERIALIZER_ERROR
 ,DS_SCORE_DESERIALIZER_ERROR
} ds_status;

/* device type */
typedef enum {
  DS_DEVICE_NATIVE_CPU = 0
 ,DS_DEVICE_OPENCL_DEVICE
} ds_device_type;


typedef struct {
  ds_device_type  type;
  cl_device_type  oclDeviceType;
  cl_device_id    oclDeviceID;
  char*           oclDeviceName;
  char*           oclDriverVersion;
  cl_uint         oclMaxClockFrequency;
  cl_uint         oclMaxComputeUnits;
  void*           score;            /* a pointer to the score data, the content/format is application defined */
} ds_device;

typedef struct {
  unsigned int  numDevices;
  ds_device*    devices;
  const char*   version;
} ds_profile;

/* deallocate memory used by score */
typedef ds_status (*ds_score_release)(void* score);

static ds_status releaseDeviceResource(ds_device* device, ds_score_release sr) {
  ds_status status = DS_SUCCESS;
  if (device) {
    if (device->oclDeviceName)      free(device->oclDeviceName);
    if (device->oclDriverVersion)   free(device->oclDriverVersion);
    if (device->score)              status = sr(device->score);
  }
  return status;
}

static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) {
  ds_status status = DS_SUCCESS;
  if (profile!=NULL) {
    if (profile->devices!=NULL && sr!=NULL) {
      unsigned int i;
      for (i = 0; i < profile->numDevices; i++) {
        status = releaseDeviceResource(profile->devices+i,sr);
        if (status != DS_SUCCESS)
          break;
      }
      free(profile->devices);
    }
    free(profile);
  }
  return status;
}


static ds_status initDSProfile(ds_profile** p, const char* version) {
  int numDevices = 0;
  cl_uint numPlatforms = 0;
  cl_platform_id* platforms = NULL;
  cl_device_id*   devices = NULL;
  ds_status status = DS_SUCCESS;
  ds_profile* profile = NULL;
  unsigned int next = 0;
  unsigned int i;

  if (p == NULL)
    return DS_INVALID_PROFILE;

  profile = (ds_profile*)malloc(sizeof(ds_profile));
  if (profile == NULL)
    return DS_MEMORY_ERROR;

  memset(profile, 0, sizeof(ds_profile));

  OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
  if (numPlatforms > 0) {
    platforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id));
    if (platforms == NULL) {
      status = DS_MEMORY_ERROR;
      goto cleanup;
    }
    OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
    for (i = 0; i < (unsigned int)numPlatforms; i++) {
      cl_uint num;
      if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
        numDevices+=num;
    }
  }

  profile->numDevices = numDevices+1;     /* +1 to numDevices to include the native CPU */

  profile->devices = (ds_device*)malloc(profile->numDevices*sizeof(ds_device));
  if (profile->devices == NULL) {
    profile->numDevices = 0;
    status = DS_MEMORY_ERROR;
    goto cleanup;
  }
  memset(profile->devices, 0, profile->numDevices*sizeof(ds_device));

  if (numDevices > 0) {
    devices = (cl_device_id*)malloc(numDevices*sizeof(cl_device_id));
    if (devices == NULL) {
      status = DS_MEMORY_ERROR;
      goto cleanup;
    }
    for (i = 0; i < (unsigned int)numPlatforms; i++) {
      cl_uint num;

      int d;
      for (d = 0; d < 2; d++) {
        unsigned int j;
        cl_device_type deviceType;
        switch(d) {
        case 0:
          deviceType = CL_DEVICE_TYPE_GPU;
          break;
        case 1:
          deviceType = CL_DEVICE_TYPE_CPU;
          break;
        default:
          continue;
          break;
        }
        if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
          continue;
        for (j = 0; j < num; j++, next++) {
          size_t length;

          profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
          profile->devices[next].oclDeviceID = devices[j];

          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
            , 0, NULL, &length);
          profile->devices[next].oclDeviceName = (char*)malloc(sizeof(char)*length);
          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
            , length, profile->devices[next].oclDeviceName, NULL);

          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
            , 0, NULL, &length);
          profile->devices[next].oclDriverVersion = (char*)malloc(sizeof(char)*length);
          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
            , length, profile->devices[next].oclDriverVersion, NULL);

          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
            , sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);

          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
            , sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);

          OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_TYPE
            , sizeof(cl_device_type), &profile->devices[next].oclDeviceType, NULL);
        }
      }
    }
  }

  profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
  profile->version = version;

cleanup:
  if (platforms)  free(platforms);
  if (devices)    free(devices);
  if (status == DS_SUCCESS) {
    *p = profile;
  }
  else {
    if (profile) {
      if (profile->devices)
        free(profile->devices);
      free(profile);
    }
  }
  return status;
}

/* Pointer to a function that calculates the score of a device (ex: device->score)
 update the data size of score. The encoding and the format of the score data
 is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
 */
typedef ds_status (*ds_perf_evaluator)(ds_device* device, void* data);

typedef enum {
  DS_EVALUATE_ALL
  ,DS_EVALUATE_NEW_ONLY
} ds_evaluation_type;

static ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type
                         ,ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates) {
  ds_status status = DS_SUCCESS;
  unsigned int i;
  unsigned int updates = 0;

  if (profile == NULL) {
    return DS_INVALID_PROFILE;
  }
  if (evaluator == NULL) {
    return DS_INVALID_PERF_EVALUATOR;
  }

  for (i = 0; i < profile->numDevices; i++) {
    ds_status evaluatorStatus;

    switch (type) {
    case DS_EVALUATE_NEW_ONLY:
      if (profile->devices[i].score != NULL)
        break;
      /*  else fall through */
    case DS_EVALUATE_ALL:
      evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
      if (evaluatorStatus != DS_SUCCESS) {
        status = evaluatorStatus;
        return status;
      }
      updates++;
      break;
    default:
      return DS_INVALID_PERF_EVALUATOR_TYPE;
      break;
    };
  }
  if (numUpdates)
    *numUpdates = updates;
  return status;
}


#define DS_TAG_VERSION                      "<version>"
#define DS_TAG_VERSION_END                  "</version>"
#define DS_TAG_DEVICE                       "<device>"
#define DS_TAG_DEVICE_END                   "</device>"
#define DS_TAG_SCORE                        "<score>"
#define DS_TAG_SCORE_END                    "</score>"
#define DS_TAG_DEVICE_TYPE                  "<type>"
#define DS_TAG_DEVICE_TYPE_END              "</type>"
#define DS_TAG_DEVICE_NAME                  "<name>"
#define DS_TAG_DEVICE_NAME_END              "</name>"
#define DS_TAG_DEVICE_DRIVER_VERSION        "<driver>"
#define DS_TAG_DEVICE_DRIVER_VERSION_END    "</driver>"
#define DS_TAG_DEVICE_MAX_COMPUTE_UNITS     "<max cu>"
#define DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END "</max cu>"
#define DS_TAG_DEVICE_MAX_CLOCK_FREQ        "<max clock>"
#define DS_TAG_DEVICE_MAX_CLOCK_FREQ_END    "</max clock>"

#define DS_DEVICE_NATIVE_CPU_STRING  "native_cpu"



typedef ds_status (*ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
static ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file) {
  ds_status status = DS_SUCCESS;
  FILE* profileFile = NULL;


  if (profile == NULL)
    return DS_INVALID_PROFILE;

  profileFile = fopen(file, "wb");
  if (profileFile==NULL) {
    status = DS_FILE_ERROR;
  }
  else {
    unsigned int i;

    /* write version string */
    fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
    fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
    fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
    fwrite("\n", sizeof(char), 1, profileFile);

    for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) {
      void* serializedScore;
      unsigned int serializedScoreSize;

      fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);

      fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
      fwrite(&profile->devices[i].type,sizeof(ds_device_type),1, profileFile);
      fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);

      switch(profile->devices[i].type) {
      case DS_DEVICE_NATIVE_CPU:
        {
          /* There's no need to emit a device name for the native CPU device. */
          /*
          fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
          fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
          fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
          */
        }
        break;
      case DS_DEVICE_OPENCL_DEVICE:
        {
          char tmp[16];

          fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
          fwrite(profile->devices[i].oclDeviceName,sizeof(char),strlen(profile->devices[i].oclDeviceName), profileFile);
          fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);

          fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
          fwrite(profile->devices[i].oclDriverVersion,sizeof(char),strlen(profile->devices[i].oclDriverVersion), profileFile);
          fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);

          fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS), profileFile);
          sprintf(tmp,"%d",profile->devices[i].oclMaxComputeUnits);
          fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
          fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END), profileFile);

          fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ), profileFile);
          sprintf(tmp,"%d",profile->devices[i].oclMaxClockFrequency);
          fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
          fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END), profileFile);
        }
        break;
      default:
        status = DS_UNKNOWN_DEVICE_TYPE;
        break;
      };

      fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
      status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize);
      if (status == DS_SUCCESS && serializedScore!=NULL && serializedScoreSize > 0) {
        fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
        free(serializedScore);
      }
      fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
      fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
      fwrite("\n",sizeof(char),1,profileFile);
    }
    fclose(profileFile);
  }
  return status;
}


static ds_status readProFile(const char* fileName, char** content, size_t* contentSize) {
  ds_status status = DS_SUCCESS;
  FILE * input = NULL;
  size_t size = 0;
  size_t rsize = 0;
  char* binary = NULL;

  *contentSize = 0;
  *content = NULL;

  input = fopen(fileName, "rb");
  if(input == NULL) {
    return DS_FILE_ERROR;
  }

  fseek(input, 0L, SEEK_END);
  size = ftell(input);
  rewind(input);
  binary = (char*)malloc(size);
  if(binary == NULL) {
    status = DS_FILE_ERROR;
    goto cleanup;
  }
  rsize = fread(binary, sizeof(char), size, input);
  if (rsize!=size
      || ferror(input)) {
    status = DS_FILE_ERROR;
    goto cleanup;
  }
  *contentSize = size;
  *content = binary;

cleanup:
  if (input != NULL) fclose(input);
  if (status != DS_SUCCESS
      && binary != NULL) {
      free(binary);
      *content = NULL;
      *contentSize = 0;
  }
  return status;
}


static const char* findString(const char* contentStart, const char* contentEnd, const char* string) {
  size_t stringLength;
  const char* currentPosition;
  const char* found;
  found = NULL;
  stringLength = strlen(string);
  currentPosition = contentStart;
  for(currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) {
    if (*currentPosition == string[0]) {
      if (currentPosition+stringLength < contentEnd) {
        if (strncmp(currentPosition, string, stringLength) == 0) {
          found = currentPosition;
          break;
        }
      }
    }
  }
  return found;
}


typedef ds_status (*ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
static ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file) {

  ds_status status = DS_SUCCESS;
  char* contentStart = NULL;
  const char* contentEnd = NULL;
  size_t contentSize;

  if (profile==NULL)
    return DS_INVALID_PROFILE;

  status = readProFile(file, &contentStart, &contentSize);
  if (status == DS_SUCCESS) {
    const char* currentPosition;
    const char* dataStart;
    const char* dataEnd;
    size_t versionStringLength;

    contentEnd = contentStart + contentSize;
    currentPosition = contentStart;


    /* parse the version string */
    dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
    if (dataStart == NULL) {
      status = DS_PROFILE_FILE_ERROR;
      goto cleanup;
    }
    dataStart += strlen(DS_TAG_VERSION);

    dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
    if (dataEnd==NULL) {
      status = DS_PROFILE_FILE_ERROR;
      goto cleanup;
    }

    versionStringLength = strlen(profile->version);
    if (versionStringLength!=(size_t)(dataEnd-dataStart)
        || strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
      /* version mismatch */
      status = DS_PROFILE_FILE_ERROR;
      goto cleanup;
    }
    currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);

    /* parse the device information */
DisableMSCWarning(4127)
    while (1) {
RestoreMSCWarning
      unsigned int i;

      const char* deviceTypeStart;
      const char* deviceTypeEnd;
      ds_device_type deviceType;

      const char* deviceNameStart;
      const char* deviceNameEnd;

      const char* deviceScoreStart;
      const char* deviceScoreEnd;

      const char* deviceDriverStart;
      const char* deviceDriverEnd;

      const char* tmpStart;
      const char* tmpEnd;
      char tmp[16];

      cl_uint maxClockFrequency;
      cl_uint maxComputeUnits;

      dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
      if (dataStart==NULL) {
        /* nothing useful remain, quit...*/
        break;
      }
      dataStart+=strlen(DS_TAG_DEVICE);
      dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
      if (dataEnd==NULL) {
        status = DS_PROFILE_FILE_ERROR;
        goto cleanup;
      }

      /* parse the device type */
      deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
      if (deviceTypeStart==NULL) {
        status = DS_PROFILE_FILE_ERROR;
        goto cleanup;
      }
      deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE);
      deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
      if (deviceTypeEnd==NULL) {
        status = DS_PROFILE_FILE_ERROR;
        goto cleanup;
      }
      memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));


      /* parse the device name */
      if (deviceType == DS_DEVICE_OPENCL_DEVICE) {

        deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
        if (deviceNameStart==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }
        deviceNameStart+=strlen(DS_TAG_DEVICE_NAME);
        deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
        if (deviceNameEnd==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }


        deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
        if (deviceDriverStart==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }
        deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION);
        deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
        if (deviceDriverEnd ==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }


        tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
        if (tmpStart==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }
        tmpStart+=strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
        tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END);
        if (tmpEnd ==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }
        memcpy(tmp,tmpStart,tmpEnd-tmpStart);
        tmp[tmpEnd-tmpStart] = '\0';
        maxComputeUnits = strtol(tmp,(char **) NULL,10);


        tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ);
        if (tmpStart==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }
        tmpStart+=strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ);
        tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ_END);
        if (tmpEnd ==NULL) {
          status = DS_PROFILE_FILE_ERROR;
          goto cleanup;
        }
        memcpy(tmp,tmpStart,tmpEnd-tmpStart);
        tmp[tmpEnd-tmpStart] = '\0';
        maxClockFrequency = strtol(tmp,(char **) NULL,10);


        /* check if this device is on the system */
        for (i = 0; i < profile->numDevices; i++) {
          if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) {
            size_t actualDeviceNameLength;
            size_t driverVersionLength;

            actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
            driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
            if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
               && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
               && maxComputeUnits == profile->devices[i].oclMaxComputeUnits
               && maxClockFrequency == profile->devices[i].oclMaxClockFrequency
               && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
               && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==(int)0) {

              deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
              if (deviceNameStart==NULL) {
                status = DS_PROFILE_FILE_ERROR;
                goto cleanup;
              }
              deviceScoreStart+=strlen(DS_TAG_SCORE);
              deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
              status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
              if (status != DS_SUCCESS) {
                goto cleanup;
              }
            }
          }
        }

      }
      else if (deviceType == DS_DEVICE_NATIVE_CPU) {
        for (i = 0; i < profile->numDevices; i++) {
          if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) {
            deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
            if (deviceScoreStart==NULL) {
              status = DS_PROFILE_FILE_ERROR;
              goto cleanup;
            }
            deviceScoreStart+=strlen(DS_TAG_SCORE);
            deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
            status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
            if (status != DS_SUCCESS) {
              goto cleanup;
            }
          }
        }
      }

      /* skip over the current one to find the next device */
      currentPosition = dataEnd+strlen(DS_TAG_DEVICE_END);
    }
  }
cleanup:
  if (contentStart!=NULL) free(contentStart);
  return status;
}


#if 0
static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
  unsigned int i;
  if (profile == NULL || num==NULL)
    return DS_MEMORY_ERROR;
  *num=0;
  for (i = 0; i < profile->numDevices; i++) {
    if (profile->devices[i].score == NULL) {
      (*num)++;
    }
  }
  return DS_SUCCESS;
}
#endif

/*
 End of the OpenCL device selection infrastructure
*/


typedef double AccelerateScoreType;

static ds_status AcceleratePerfEvaluator(ds_device *device,
  void *magick_unused(data))
{
#define ACCELERATE_PERF_DIMEN "2048x1536"
#define NUM_ITER  2
#define ReturnStatus(status) \
{ \
  if (clEnv!=NULL) \
    RelinquishMagickOpenCLEnv(clEnv); \
  if (oldClEnv!=NULL) \
    defaultCLEnv = oldClEnv; \
  return status; \
}

  AccelerateTimer
    timer;

  ExceptionInfo
    *exception=NULL;

  MagickCLEnv
    clEnv=NULL,
    oldClEnv=NULL;

  magick_unreferenced(data);

  if (device == NULL)
    ReturnStatus(DS_PERF_EVALUATOR_ERROR);

  clEnv=AcquireMagickOpenCLEnv();
  exception=AcquireExceptionInfo();

  if (device->type == DS_DEVICE_NATIVE_CPU)
    {
      /* CPU device */
      MagickBooleanType flag=MagickTrue;
      SetMagickOpenCLEnvParamInternal(clEnv,
        MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
        &flag,exception);
    }
  else if (device->type == DS_DEVICE_OPENCL_DEVICE)
    {
      /* OpenCL device */
      SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
        sizeof(cl_device_id),&device->oclDeviceID,exception);
    }
  else
    ReturnStatus(DS_PERF_EVALUATOR_ERROR);

  /* recompile the OpenCL kernels if it needs to */
  clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;

  InitOpenCLEnvInternal(clEnv,exception);
  oldClEnv=defaultCLEnv;
  defaultCLEnv=clEnv;

  /* microbenchmark */
  {
    Image
      *inputImage;

    ImageInfo
      *imageInfo;

    int
      i;

    imageInfo=AcquireImageInfo();
    CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
    CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
    inputImage=ReadImage(imageInfo,exception);

    initAccelerateTimer(&timer);

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

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

#ifdef MAGICKCORE_CLPERFMARKER
      clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
#endif

      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,1.0,
        exception);

#ifdef MAGICKCORE_CLPERFMARKER
      clEndPerfMarkerAMD();
#endif

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

      if (bluredImage)
        DestroyImage(bluredImage);
      if (unsharpedImage)
        DestroyImage(unsharpedImage);
      if (resizedImage)
        DestroyImage(resizedImage);
    }
    DestroyImage(inputImage);
  }
  /* end of microbenchmark */

  if (device->score == NULL)
    device->score=malloc(sizeof(AccelerateScoreType));
  *(AccelerateScoreType*)device->score=readAccelerateTimer(&timer);

  ReturnStatus(DS_SUCCESS);
}

ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
  if (device
     && device->score) {
    /* generate a string from the score */
    char* s = (char*)malloc(sizeof(char)*256);
    sprintf(s,"%.4f",*((AccelerateScoreType*)device->score));
    *serializedScore = (void*)s;
    *serializedScoreSize = (unsigned int) strlen(s);
    return DS_SUCCESS;
  }
  else {
    return DS_SCORE_SERIALIZER_ERROR;
  }
}

ds_status AccelerateScoreDeserializer(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize) {
  if (device) {
    /* convert the string back to an int */
    char* s = (char*)malloc(serializedScoreSize+1);
    memcpy(s, serializedScore, serializedScoreSize);
    s[serializedScoreSize] = (char)'\0';
    device->score = malloc(sizeof(AccelerateScoreType));
    *((AccelerateScoreType*)device->score) = (AccelerateScoreType)
      strtod(s, (char **) NULL);
    free(s);
    return DS_SUCCESS;
  }
  else {
    return DS_SCORE_DESERIALIZER_ERROR;
  }
}

ds_status AccelerateScoreRelease(void* score) {
  if (score!=NULL) {
    free(score);
  }
  return DS_SUCCESS;
}

ds_status canWriteProfileToFile(const char *path)
{
  FILE* profileFile = fopen(path, "ab");

  if (profileFile==NULL)
    return DS_FILE_ERROR;

  fclose(profileFile);
  return DS_SUCCESS;
}


#define IMAGEMAGICK_PROFILE_VERSION "ImageMagick Device Selection v0.9"
#define IMAGEMAGICK_PROFILE_FILE    "ImagemagickOpenCLDeviceProfile"
static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {

  MagickBooleanType mStatus = MagickFalse;
  ds_status status;
  ds_profile* profile;
  unsigned int numDeviceProfiled = 0;
  unsigned int i;
  unsigned int bestDeviceIndex;
  AccelerateScoreType bestScore;
  char path[MaxTextExtent];
  MagickBooleanType flag;
  ds_evaluation_type profileType;

  LockDefaultOpenCLEnv();

  /* Initially, just set OpenCL to off */
  flag = MagickTrue;
  SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
    , sizeof(MagickBooleanType), &flag, exception);

  /* check and init the global lib */
  OpenCLLib=GetOpenCLLib();
  if (OpenCLLib==NULL)
  {
    mStatus=InitOpenCLEnvInternal(clEnv, exception);
    goto cleanup;
  }

  status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
  if (status!=DS_SUCCESS) {
    (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
    goto cleanup;
  }

  (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
         ,GetOpenCLCachedFilesDirectory()
         ,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);

  if (canWriteProfileToFile(path) != DS_SUCCESS) {
    /* We can not write out a device profile, so don't run the benchmark */
    /* select the first GPU device */

    bestDeviceIndex = 0;
    for (i = 1; i < profile->numDevices; i++) {
      if ((profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) && (profile->devices[i].oclDeviceType == CL_DEVICE_TYPE_GPU)) {
        bestDeviceIndex = i;
        break;
      }
    }
  }
  else {
    if (clEnv->regenerateProfile != MagickFalse) {
      profileType = DS_EVALUATE_ALL;
    }
    else {
      readProfileFromFile(profile, AccelerateScoreDeserializer, path);
      profileType = DS_EVALUATE_NEW_ONLY;
    }
    status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);

    if (status!=DS_SUCCESS) {
      (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
      goto cleanup;
    }
    if (numDeviceProfiled > 0) {
      status = writeProfileToFile(profile, AccelerateScoreSerializer, path);
      if (status!=DS_SUCCESS) {
        (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when saving the profile into a file", "'%s'", ".");
      }
    }

    /* pick the best device */
    bestDeviceIndex = 0;
    bestScore = *(AccelerateScoreType*)profile->devices[bestDeviceIndex].score;
    for (i = 1; i < profile->numDevices; i++) {
      AccelerateScoreType score = *(AccelerateScoreType*)profile->devices[i].score;
      if (score < bestScore) {
        bestDeviceIndex = i;
        bestScore = score;
      }
    }
  }

  /* set up clEnv with the best device */
  if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
    /* CPU device */
    flag = MagickTrue;
    SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
                                  , sizeof(MagickBooleanType), &flag, exception);
  }
  else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
    /* OpenCL device */
    flag = MagickFalse;
    SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
      , sizeof(MagickBooleanType), &flag, exception);
    SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
      , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
  }
  else {
    status = DS_PERF_EVALUATOR_ERROR;
    goto cleanup;
  }
  mStatus=InitOpenCLEnvInternal(clEnv, exception);

  status = releaseDSProfile(profile, AccelerateScoreRelease);
  if (status!=DS_SUCCESS) {
    (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
  }

cleanup:

  UnlockDefaultOpenCLEnv();
  return mStatus;
}


/*
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%                                                                             %
%                                                                             %
%                                                                             %
+   I n i t I m a g e M a g i c k O p e n C L                                 %
%                                                                             %
%                                                                             %
%                                                                             %
%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
%
%  InitImageMagickOpenCL() provides a simplified interface to initialize
%  the OpenCL environtment in ImageMagick
%
%  The format of the InitImageMagickOpenCL() method is:
%
%      MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
%                                        void* userSelectedDevice,
%                                        void* selectedDevice)
%
%  A description of each parameter follows:
%
%    o mode: OpenCL mode in ImageMagick, could be off,auto,user
%
%    o userSelectedDevice:  when in user mode, a pointer to the selected
%                           cl_device_id
%
%    o selectedDevice: a pointer to cl_device_id where the selected
%                      cl_device_id by ImageMagick could be returned
%
%    o exception: exception
%
*/
MagickExport MagickBooleanType InitImageMagickOpenCL(
  ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
  ExceptionInfo *exception)
{
  MagickBooleanType status = MagickFalse;
  MagickCLEnv clEnv = NULL;
  MagickBooleanType flag;

  clEnv = GetDefaultOpenCLEnv();
  if (clEnv!=NULL) {
    switch(mode) {

    case MAGICK_OPENCL_OFF:
      flag = MagickTrue;
      SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
        , sizeof(MagickBooleanType), &flag, exception);
      status = InitOpenCLEnv(clEnv, exception);

      if (selectedDevice)
        *(cl_device_id*)selectedDevice = NULL;
      break;

    case MAGICK_OPENCL_DEVICE_SELECT_USER:

      if (userSelectedDevice == NULL)
        return MagickFalse;

      flag = MagickFalse;
      SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
        , sizeof(MagickBooleanType), &flag, exception);

      SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
        , sizeof(cl_device_id), userSelectedDevice,exception);

      status = InitOpenCLEnv(clEnv, exception);
      if (selectedDevice) {
        GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
          , sizeof(cl_device_id), selectedDevice, exception);
      }
      break;

    case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
        flag = MagickTrue;
        SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
          , sizeof(MagickBooleanType), &flag, exception);
        flag = MagickTrue;
        SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
          , sizeof(MagickBooleanType), &flag, exception);

    /* fall through here!! */
    case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
    default:
      {
        cl_device_id d = NULL;
        flag = MagickFalse;
        SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
          , sizeof(MagickBooleanType), &flag, exception);
        SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
          , sizeof(cl_device_id), &d,exception);
        status = InitOpenCLEnv(clEnv, exception);
        if (selectedDevice) {
          GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
            , sizeof(cl_device_id),  selectedDevice, exception);
        }
      }
      break;
    };
  }
  return status;
}


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

  MagickCLEnv clEnv;

  status = MagickTrue;

  clEnv = GetDefaultOpenCLEnv();

  assert(exception != (ExceptionInfo *) NULL);
  assert(exception->signature == MagickSignature);

  if (severity!=0) {
    cl_device_type dType;
    clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
    if (dType == CL_DEVICE_TYPE_CPU) {
      char buffer[MaxTextExtent];
      clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);

      /* Workaround for Intel OpenCL CPU runtime bug */
      /* Turn off OpenCL when a problem is detected! */
      if (strncmp(buffer, "Intel",5) == 0) {

        InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
      }
    }
  }

#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);
}

MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
{
  LockSemaphoreInfo(clEnv->lock);
  if (clEnv->seedsLock == NULL)
  {
    ActivateSemaphoreInfo(&clEnv->seedsLock);
  }
  LockSemaphoreInfo(clEnv->seedsLock);

  if (clEnv->seeds == NULL)
  {
    cl_int clStatus;
    clEnv->numGenerators = NUM_CL_RAND_GENERATORS;
    clEnv->seeds = clEnv->library->clCreateBuffer(clEnv->context, CL_MEM_READ_WRITE,
                                  clEnv->numGenerators*4*sizeof(unsigned int),
                                  NULL, &clStatus);
    if (clStatus != CL_SUCCESS)
    {
      clEnv->seeds = NULL;
    }
    else
    {
      unsigned int i;
      cl_command_queue queue = NULL;
      unsigned int *seeds;

      queue = AcquireOpenCLCommandQueue(clEnv);
      seeds = (unsigned int*) clEnv->library->clEnqueueMapBuffer(queue, clEnv->seeds, CL_TRUE,
                                                  CL_MAP_WRITE, 0,
                                                  clEnv->numGenerators*4
                                                  *sizeof(unsigned int),
                                                  0, NULL, NULL, &clStatus);
      if (clStatus!=CL_SUCCESS)
      {
        clEnv->library->clReleaseMemObject(clEnv->seeds);
        goto cleanup;
      }

      for (i = 0; i < clEnv->numGenerators; i++) {
        RandomInfo* randomInfo = AcquireRandomInfo();
        const unsigned long* s = GetRandomInfoSeed(randomInfo);
        if (i == 0)
          clEnv->randNormalize = GetRandomInfoNormalize(randomInfo);

        seeds[i*4]   = (unsigned int) s[0];
        seeds[i*4+1] = (unsigned int) 0x50a7f451;
        seeds[i*4+2] = (unsigned int) 0x5365417e;
        seeds[i*4+3] = (unsigned int) 0xc3a4171a;

        randomInfo = DestroyRandomInfo(randomInfo);
      }
      clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, clEnv->seeds, seeds, 0,
                                          NULL, NULL);
      clEnv->library->clFinish(queue);
cleanup:
      if (queue != NULL)
        RelinquishOpenCLCommandQueue(clEnv, queue);
    }
  }
  UnlockSemaphoreInfo(clEnv->lock);
  return clEnv->seeds;
}

MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv) {
  if (clEnv->seedsLock == NULL)
  {
    ActivateSemaphoreInfo(&clEnv->seedsLock);
  }
  else
    UnlockSemaphoreInfo(clEnv->seedsLock);
}

MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
{
  return clEnv->numGenerators;
}


MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
{
  return clEnv->randNormalize;
}

#else

struct _MagickCLEnv {
  MagickBooleanType OpenCLInitialized;  /* whether OpenCL environment is initialized. */
};

MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
{
  return NULL;
}

MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(
  MagickCLEnv magick_unused(clEnv))
{
  magick_unreferenced(clEnv);

  return MagickFalse;
}

/*
* Return the OpenCL environment
*/
MagickExport MagickCLEnv GetDefaultOpenCLEnv()
{
  return (MagickCLEnv) NULL;
}

MagickExport MagickCLEnv SetDefaultOpenCLEnv(
  MagickCLEnv magick_unused(clEnv))
{
  magick_unreferenced(clEnv);

  return (MagickCLEnv) NULL;
}

MagickExport MagickBooleanType SetMagickOpenCLEnvParam(
  MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
  size_t magick_unused(dataSize),void *magick_unused(data),
  ExceptionInfo *magick_unused(exception))
{
  magick_unreferenced(clEnv);
  magick_unreferenced(param);
  magick_unreferenced(dataSize);
  magick_unreferenced(data);
  magick_unreferenced(exception);

  return MagickFalse;
}

MagickExport MagickBooleanType GetMagickOpenCLEnvParam(
  MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
  size_t magick_unused(dataSize),void *magick_unused(data),
  ExceptionInfo *magick_unused(exception))
{
  magick_unreferenced(clEnv);
  magick_unreferenced(param);
  magick_unreferenced(dataSize);
  magick_unreferenced(data);
  magick_unreferenced(exception);

  return MagickFalse;
}

MagickExport MagickBooleanType InitOpenCLEnv(MagickCLEnv magick_unused(clEnv),
  ExceptionInfo *magick_unused(exception))
{
  magick_unreferenced(clEnv);
  magick_unreferenced(exception);

  return MagickFalse;
}

MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(
  MagickCLEnv magick_unused(clEnv))
{
  magick_unreferenced(clEnv);

  return (cl_command_queue) NULL;
}

MagickExport MagickBooleanType RelinquishCommandQueue(
  MagickCLEnv magick_unused(clEnv),cl_command_queue magick_unused(queue))
{
  magick_unreferenced(clEnv);
  magick_unreferenced(queue);

  return MagickFalse;
}

MagickPrivate cl_kernel AcquireOpenCLKernel(
  MagickCLEnv magick_unused(clEnv),MagickOpenCLProgram magick_unused(program),
  const char *magick_unused(kernelName))
{
  magick_unreferenced(clEnv);
  magick_unreferenced(program);
  magick_unreferenced(kernelName);

  return (cl_kernel) NULL;
}

MagickPrivate MagickBooleanType RelinquishOpenCLKernel(
  MagickCLEnv magick_unused(clEnv),cl_kernel magick_unused(kernel))
{
  magick_unreferenced(clEnv);
  magick_unreferenced(kernel);

  return MagickFalse;
}

MagickPrivate unsigned long GetOpenCLDeviceLocalMemorySize(
  MagickCLEnv magick_unused(clEnv))
{
  magick_unreferenced(clEnv);

  return 0;
}

MagickExport MagickBooleanType InitImageMagickOpenCL(
  ImageMagickOpenCLMode magick_unused(mode),
  void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
  ExceptionInfo *magick_unused(exception))
{
  magick_unreferenced(mode);
  magick_unreferenced(userSelectedDevice);
  magick_unreferenced(selectedDevice);
  magick_unreferenced(exception);
  return MagickFalse;
}


MagickPrivate
MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
  const char *module,const char *function,const size_t line,
  const ExceptionType severity,const char *tag,const char *format,...)
{
  magick_unreferenced(exception);
  magick_unreferenced(module);
  magick_unreferenced(function);
  magick_unreferenced(line);
  magick_unreferenced(severity);
  magick_unreferenced(tag);
  magick_unreferenced(format);
  return(MagickFalse);
}


MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
{
  magick_unreferenced(clEnv);
  return NULL;
}


MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv)
{
  magick_unreferenced(clEnv);
}

MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
{
  magick_unreferenced(clEnv);
  return 0;
}

MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
{
  magick_unreferenced(clEnv);
  return 0.0f;
}

#endif /* MAGICKCORE_OPENCL_SUPPORT */

char* openclCachedFilesDirectory;
SemaphoreInfo* openclCachedFilesDirectoryLock;

MagickPrivate
const char* GetOpenCLCachedFilesDirectory() {
  if (openclCachedFilesDirectory == NULL) {
    if (openclCachedFilesDirectoryLock == NULL)
    {
      ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
    }
    LockSemaphoreInfo(openclCachedFilesDirectoryLock);
    if (openclCachedFilesDirectory == NULL) {
      char path[MaxTextExtent];
      char *home = NULL;
      char *temp = NULL;
      struct stat attributes;
      MagickBooleanType status;
      int mkdirStatus = 0;



      home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
      if (home == (char *) NULL)
      {
        home=GetEnvironmentValue("XDG_CACHE_HOME");
        if (home == (char *) NULL)
          home=GetEnvironmentValue("LOCALAPPDATA");
        if (home == (char *) NULL)
          home=GetEnvironmentValue("APPDATA");
        if (home == (char *) NULL)
          home=GetEnvironmentValue("USERPROFILE");
      }

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

#ifdef MAGICKCORE_WINDOWS_SUPPORT
          mkdirStatus = mkdir(path);
#else
          mkdirStatus = mkdir(path, 0777);
#endif
        }

        /* first check if $HOME/ImageMagick exists */
        if (mkdirStatus==0)
        {
            (void) FormatLocaleString(path,MaxTextExtent,
              "%s%sImageMagick",home,DirectorySeparator);

            status=GetPathAttributes(path,&attributes);
            if (status == MagickFalse)
            {
#ifdef MAGICKCORE_WINDOWS_SUPPORT
              mkdirStatus = mkdir(path);
#else
              mkdirStatus = mkdir(path, 0777);
#endif
            }
        }

        if (mkdirStatus==0)
        {
          temp = (char*)AcquireMagickMemory(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,MaxTextExtent,"%s%s.cache",
            home,DirectorySeparator);
          status=GetPathAttributes(path,&attributes);
          if (status == MagickFalse)
          {

#ifdef   MAGICKCORE_WINDOWS_SUPPORT
            mkdirStatus = mkdir(path);
#else
            mkdirStatus = mkdir(path, 0777);
#endif
          }

          /* first check if $HOME/.cache/ImageMagick exists */
          if (mkdirStatus==0)
          {
              (void) FormatLocaleString(path,MaxTextExtent,
                "%s%s.cache%sImageMagick",home,DirectorySeparator,
                DirectorySeparator);

              status=GetPathAttributes(path,&attributes);
              if (status == MagickFalse)
              {
#ifdef   MAGICKCORE_WINDOWS_SUPPORT
                mkdirStatus = mkdir(path);
#else
                mkdirStatus = mkdir(path, 0777);
#endif
              }
          }

          if (mkdirStatus==0)
          {
            temp = (char*)AcquireMagickMemory(strlen(path)+1);
            CopyMagickString(temp,path,strlen(path)+1);
          }
          home=DestroyString(home);
        }
      }
      openclCachedFilesDirectory = temp;
    }
    UnlockSemaphoreInfo(openclCachedFilesDirectoryLock);
  }
  return openclCachedFilesDirectory;
}

/* create a function for OpenCL log */
MagickPrivate
void OpenCLLog(const char* message) {

#ifdef OPENCLLOG_ENABLED
#define OPENCL_LOG_FILE "ImageMagickOpenCL.log"

  FILE* log;
  if (getenv("MAGICK_OCL_LOG"))
  {
    if (message) {
      char path[MaxTextExtent];
      unsigned long allocSize;

      MagickCLEnv clEnv;

      clEnv = GetDefaultOpenCLEnv();

      /*  dump the source into a file */
      (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
        ,GetOpenCLCachedFilesDirectory()
        ,DirectorySeparator,OPENCL_LOG_FILE);


      log = fopen(path, "ab");
      fwrite(message, sizeof(char), strlen(message), log);
      fwrite("\n", sizeof(char), 1, log);

      if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
      {
        allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
        fprintf(log, "Devic Max Memory Alloc Size: %lu\n", allocSize);
      }

      fclose(log);
    }
  }
#else
  magick_unreferenced(message);
#endif
}

MagickPrivate void OpenCLTerminus()
{
#if MAGICKCORE_OPENCL_SUPPORT
  DumpProfileData();
  if (openclCachedFilesDirectory != (char *) NULL)
    openclCachedFilesDirectory=DestroyString(openclCachedFilesDirectory);
  if (openclCachedFilesDirectoryLock != (SemaphoreInfo*)NULL)
    DestroySemaphoreInfo(&openclCachedFilesDirectoryLock);
  if (defaultCLEnv != (MagickCLEnv) NULL)
    {
       (void) RelinquishMagickOpenCLEnv(defaultCLEnv);
       defaultCLEnv=(MagickCLEnv)NULL;
    }
  if (defaultCLEnvLock != (SemaphoreInfo*) NULL)
    DestroySemaphoreInfo(&defaultCLEnvLock);
  if (OpenCLLib != (MagickLibrary *)NULL)
    OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
  if (OpenCLLibLock != (SemaphoreInfo*)NULL)
    DestroySemaphoreInfo(&OpenCLLibLock);
#endif