granicus.if.org Git - imagemagick/blob - MagickCore/opencl.c

   1 /*
   2 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
   3 %                                                                             %
   4 %                                                                             %
   5 %                                                                             %
   6 %                   OOO   PPPP   EEEEE  N   N   CCCC  L                       %
   7 %                  O   O  P   P  E      NN  N  C      L                       %
   8 %                  O   O  PPPP   EEE    N N N  C      L                       %
   9 %                  O   O  P      E      N  NN  C      L                       %
  10 %                   OOO   P      EEEEE  N   N   CCCC  LLLLL                   %
  11 %                                                                             %
  12 %                                                                             %
  13 %                         MagickCore OpenCL Methods                           %
  14 %                                                                             %
  15 %                              Software Design                                %
  16 %                                   Cristy                                    %
  17 %                                 March 2000                                  %
  18 %                                                                             %
  19 %                                                                             %
  20 %  Copyright 1999-2014 ImageMagick Studio LLC, a non-profit organization      %
  21 %  dedicated to making software imaging solutions freely available.           %
  22 %                                                                             %
  23 %  You may not use this file except in compliance with the License.  You may  %
  24 %  obtain a copy of the License at                                            %
  25 %                                                                             %
  26 %    http://www.imagemagick.org/script/license.php                            %
  27 %                                                                             %
  28 %  Unless required by applicable law or agreed to in writing, software        %
  29 %  distributed under the License is distributed on an "AS IS" BASIS,          %
  30 %  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.   %
  31 %  See the License for the specific language governing permissions and        %
  32 %  limitations under the License.                                             %
  33 %                                                                             %
  34 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
  35 %
  36 %
  37 %
  38 */
  39
  40 /*
  41 Include declarations.
  42 */
  43 #include "MagickCore/studio.h"
  44 #include "MagickCore/artifact.h"
  45 #include "MagickCore/cache.h"
  46 #include "MagickCore/color.h"
  47 #include "MagickCore/compare.h"
  48 #include "MagickCore/constitute.h"
  49 #include "MagickCore/distort.h"
  50 #include "MagickCore/draw.h"
  51 #include "MagickCore/effect.h"
  52 #include "MagickCore/exception.h"
  53 #include "MagickCore/exception-private.h"
  54 #include "MagickCore/fx.h"
  55 #include "MagickCore/gem.h"
  56 #include "MagickCore/geometry.h"
  57 #include "MagickCore/image.h"
  58 #include "MagickCore/image-private.h"
  59 #include "MagickCore/layer.h"
  60 #include "MagickCore/mime-private.h"
  61 #include "MagickCore/memory_.h"
  62 #include "MagickCore/monitor.h"
  63 #include "MagickCore/montage.h"
  64 #include "MagickCore/morphology.h"
  65 #include "MagickCore/nt-base.h"
  66 #include "MagickCore/nt-base-private.h"
  67 #include "MagickCore/opencl.h"
  68 #include "MagickCore/opencl-private.h"
  69 #include "MagickCore/option.h"
  70 #include "MagickCore/policy.h"
  71 #include "MagickCore/property.h"
  72 #include "MagickCore/quantize.h"
  73 #include "MagickCore/quantum.h"
  74 #include "MagickCore/random_.h"
  75 #include "MagickCore/random-private.h"
  76 #include "MagickCore/resample.h"
  77 #include "MagickCore/resource_.h"
  78 #include "MagickCore/splay-tree.h"
  79 #include "MagickCore/semaphore.h"
  80 #include "MagickCore/statistic.h"
  81 #include "MagickCore/string_.h"
  82 #include "MagickCore/token.h"
  83 #include "MagickCore/utility.h"
  84
  85 #ifdef MAGICKCORE_CLPERFMARKER
  86 #include "CLPerfMarker.h"
  87 #endif
  88
  89
  90 #if defined(MAGICKCORE_OPENCL_SUPPORT)
  91
  92 #ifdef MAGICKCORE_HAVE_OPENCL_CL_H
  93 #define MAGICKCORE_OPENCL_MACOSX  1
  94 #endif
  95
  96
  97 #define NUM_CL_RAND_GENERATORS 1024  /* number of random number generators running in parallel */
  98
  99 /*
 100  *
 101  * Dynamic library loading functions
 102  *
 103  */
 104 #ifdef MAGICKCORE_WINDOWS_SUPPORT
 105 #else
 106 #include <dlfcn.h>
 107 #endif
 108
 109 // dynamically load a library.  returns NULL on failure
 110 void *OsLibraryLoad(const char *libraryName)
 111 {
 112 #ifdef MAGICKCORE_WINDOWS_SUPPORT
 113     return (void *)LoadLibraryA(libraryName);
 114 #else
 115     return (void *)dlopen(libraryName, RTLD_NOW);
 116 #endif
 117 }
 118
 119 // get a function pointer from a loaded library.  returns NULL on failure.
 120 void *OsLibraryGetFunctionAddress(void *library, const char *functionName)
 121 {
 122 #ifdef MAGICKCORE_WINDOWS_SUPPORT
 123     if (!library || !functionName)
 124     {
 125         return NULL;
 126     }
 127     return (void *) GetProcAddress( (HMODULE)library, functionName);
 128 #else
 129     if (!library || !functionName)
 130     {
 131         return NULL;
 132     }
 133     return (void *)dlsym(library, functionName);
 134 #endif
 135 }
 136
 137 // unload a library.
 138 void OsLibraryUnload(void *library)
 139 {
 140 #ifdef MAGICKCORE_WINDOWS_SUPPORT
 141     FreeLibrary( (HMODULE)library);
 142 #else
 143     dlclose(library);
 144 #endif
 145 }
 146
 147
 148 /*
 149 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 150 %                                                                             %
 151 %                                                                             %
 152 %                                                                             %
 153 +   A c q u i r e M a g i c k O p e n C L E n v                               %
 154 %                                                                             %
 155 %                                                                             %
 156 %                                                                             %
 157 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 158 %
 159 % AcquireMagickOpenCLEnv() allocates the MagickCLEnv structure
 160 %
 161 */
 162
 163 MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
 164 {
 165   MagickCLEnv clEnv;
 166   clEnv = (MagickCLEnv) AcquireMagickMemory(sizeof(struct _MagickCLEnv));
 167   if (clEnv != NULL)
 168   {
 169     memset(clEnv, 0, sizeof(struct _MagickCLEnv));
 170     ActivateSemaphoreInfo(&clEnv->lock);
 171   }
 172   return clEnv;
 173 }
 174
 175
 176 /*
 177 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 178 %                                                                             %
 179 %                                                                             %
 180 %                                                                             %
 181 +   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                         %
 182 %                                                                             %
 183 %                                                                             %
 184 %                                                                             %
 185 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 186 %
 187 %  RelinquishMagickOpenCLEnv() destroy the MagickCLEnv structure
 188 %
 189 %  The format of the RelinquishMagickOpenCLEnv method is:
 190 %
 191 %      MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv)
 192 %
 193 %  A description of each parameter follows:
 194 %
 195 %    o clEnv: MagickCLEnv structure to destroy
 196 %
 197 */
 198
 199 MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(MagickCLEnv clEnv)
 200 {
 201   if (clEnv != (MagickCLEnv)NULL)
 202   {
 203     DestroySemaphoreInfo(&clEnv->lock);
 204     RelinquishMagickMemory(clEnv);
 205     return MagickTrue;
 206   }
 207   return MagickFalse;
 208 }
 209
 210
 211 /*
 212 * Default OpenCL environment
 213 */
 214 MagickCLEnv defaultCLEnv;
 215 SemaphoreInfo* defaultCLEnvLock;
 216
 217 /*
 218 * OpenCL library
 219 */
 220 MagickLibrary * OpenCLLib;
 221 SemaphoreInfo* OpenCLLibLock;
 222
 223
 224 static MagickBooleanType bindOpenCLFunctions(void* library)
 225 {
 226 #ifdef MAGICKCORE_OPENCL_MACOSX
 227 #define BIND(X) OpenCLLib->X= &X;
 228 #else
 229 #define BIND(X)\
 230   if ((OpenCLLib->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(library,#X)) == NULL)\
 231   return MagickFalse;
 232 #endif
 233
 234   BIND(clGetPlatformIDs);
 235   BIND(clGetPlatformInfo);
 236
 237   BIND(clGetDeviceIDs);
 238   BIND(clGetDeviceInfo);
 239
 240   BIND(clCreateContext);
 241
 242   BIND(clCreateBuffer);
 243   BIND(clReleaseMemObject);
 244
 245   BIND(clCreateProgramWithSource);
 246   BIND(clCreateProgramWithBinary);
 247   BIND(clBuildProgram);
 248   BIND(clGetProgramInfo);
 249   BIND(clGetProgramBuildInfo);
 250
 251   BIND(clCreateKernel);
 252   BIND(clReleaseKernel);
 253   BIND(clSetKernelArg);
 254
 255   BIND(clFlush);
 256   BIND(clFinish);
 257
 258   BIND(clEnqueueNDRangeKernel);
 259   BIND(clEnqueueReadBuffer);
 260   BIND(clEnqueueMapBuffer);
 261   BIND(clEnqueueUnmapMemObject);
 262
 263   BIND(clCreateCommandQueue);
 264   BIND(clReleaseCommandQueue);
 265
 266   return MagickTrue;
 267 }
 268
 269 MagickLibrary * GetOpenCLLib()
 270 {
 271   if (OpenCLLib == NULL)
 272   {
 273     if (OpenCLLibLock == NULL)
 274     {
 275       ActivateSemaphoreInfo(&OpenCLLibLock);
 276     }
 277
 278     LockSemaphoreInfo(OpenCLLibLock);
 279
 280     OpenCLLib = (MagickLibrary *) AcquireMagickMemory (sizeof (MagickLibrary));
 281
 282     if (OpenCLLib != NULL)
 283     {
 284       MagickBooleanType status = MagickFalse;
 285       void * library = NULL;
 286
 287 #ifdef MAGICKCORE_OPENCL_MACOSX
 288       status = bindOpenCLFunctions(library);
 289 #else
 290
 291       memset(OpenCLLib, 0, sizeof(MagickLibrary));
 292 #ifdef MAGICKCORE_WINDOWS_SUPPORT
 293       library = OsLibraryLoad("OpenCL.dll");
 294 #else
 295       library = OsLibraryLoad("libOpenCL.so");
 296 #endif
 297       if (library)
 298         status = bindOpenCLFunctions(library);
 299
 300       if (status==MagickTrue)
 301         OpenCLLib->base=library;
 302       else
 303         OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
 304 #endif
 305     }
 306
 307     UnlockSemaphoreInfo(OpenCLLibLock);
 308   }
 309
 310
 311   return OpenCLLib;
 312 }
 313
 314
 315 /*
 316 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 317 %                                                                             %
 318 %                                                                             %
 319 %                                                                             %
 320 +   G e t D e f a u l t O p e n C L E n v                                     %
 321 %                                                                             %
 322 %                                                                             %
 323 %                                                                             %
 324 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 325 %
 326 %  GetDefaultOpenCLEnv() returns the default OpenCL env
 327 %
 328 %  The format of the GetDefaultOpenCLEnv method is:
 329 %
 330 %      MagickCLEnv GetDefaultOpenCLEnv()
 331 %
 332 %  A description of each parameter follows:
 333 %
 334 %    o exception: return any errors or warnings.
 335 %
 336 */
 337
 338 MagickExport MagickCLEnv GetDefaultOpenCLEnv()
 339 {
 340   if (defaultCLEnv == NULL)
 341   {
 342     if (defaultCLEnvLock == NULL)
 343     {
 344       ActivateSemaphoreInfo(&defaultCLEnvLock);
 345     }
 346     LockSemaphoreInfo(defaultCLEnvLock);
 347     defaultCLEnv = AcquireMagickOpenCLEnv();
 348     UnlockSemaphoreInfo(defaultCLEnvLock);
 349   }
 350   return defaultCLEnv;
 351 }
 352
 353 static void LockDefaultOpenCLEnv() {
 354   if (defaultCLEnvLock == NULL)
 355   {
 356     ActivateSemaphoreInfo(&defaultCLEnvLock);
 357   }
 358   LockSemaphoreInfo(defaultCLEnvLock);
 359 }
 360
 361 static void UnlockDefaultOpenCLEnv() {
 362   if (defaultCLEnvLock == NULL)
 363   {
 364     ActivateSemaphoreInfo(&defaultCLEnvLock);
 365   }
 366   else
 367     UnlockSemaphoreInfo(defaultCLEnvLock);
 368 }
 369
 370
 371 /*
 372 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 373 %                                                                             %
 374 %                                                                             %
 375 %                                                                             %
 376 +   S e t D e f a u l t O p e n C L E n v                                     %
 377 %                                                                             %
 378 %                                                                             %
 379 %                                                                             %
 380 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 381 %
 382 %  SetDefaultOpenCLEnv() sets the new OpenCL environment as default
 383 %  and returns the old OpenCL environment
 384 %
 385 %  The format of the SetDefaultOpenCLEnv() method is:
 386 %
 387 %      MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv)
 388 %
 389 %  A description of each parameter follows:
 390 %
 391 %    o clEnv: the new default OpenCL environment.
 392 %
 393 */
 394 MagickExport MagickCLEnv SetDefaultOpenCLEnv(MagickCLEnv clEnv)
 395 {
 396   MagickCLEnv oldEnv;
 397   LockDefaultOpenCLEnv();
 398   oldEnv = defaultCLEnv;
 399   defaultCLEnv = clEnv;
 400   UnlockDefaultOpenCLEnv();
 401   return oldEnv;
 402 }
 403
 404
 405
 406 /*
 407 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 408 %                                                                             %
 409 %                                                                             %
 410 %                                                                             %
 411 +   S e t M a g i c k O p e n C L E n v P a r a m                             %
 412 %                                                                             %
 413 %                                                                             %
 414 %                                                                             %
 415 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 416 %
 417 %  SetMagickOpenCLEnvParam() sets the parameters in the OpenCL environment
 418 %
 419 %  The format of the SetMagickOpenCLEnvParam() method is:
 420 %
 421 %      MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv,
 422 %        MagickOpenCLEnvParam param, size_t dataSize, void* data,
 423 %        ExceptionInfo* exception)
 424 %
 425 %  A description of each parameter follows:
 426 %
 427 %    o clEnv: the OpenCL environment.
 428 %
 429 %    o param: the parameter to be set.
 430 %
 431 %    o dataSize: the data size of the parameter value.
 432 %
 433 %    o data:  the pointer to the new parameter value
 434 %
 435 %    o exception: return any errors or warnings
 436 %
 437 */
 438
 439 static MagickBooleanType SetMagickOpenCLEnvParamInternal(MagickCLEnv clEnv, MagickOpenCLEnvParam param
 440                                           , size_t dataSize, void* data, ExceptionInfo* exception)
 441 {
 442   MagickBooleanType status = MagickFalse;
 443
 444   if (clEnv == NULL
 445     || data == NULL)
 446     goto cleanup;
 447
 448   switch(param)
 449   {
 450   case MAGICK_OPENCL_ENV_PARAM_DEVICE:
 451     if (dataSize != sizeof(clEnv->device))
 452       goto cleanup;
 453     clEnv->device = *((cl_device_id*)data);
 454     clEnv->OpenCLInitialized = MagickFalse;
 455     status = MagickTrue;
 456     break;
 457
 458   case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED:
 459     if (dataSize != sizeof(clEnv->OpenCLDisabled))
 460       goto cleanup;
 461     clEnv->OpenCLDisabled =  *((MagickBooleanType*)data);
 462     clEnv->OpenCLInitialized = MagickFalse;
 463     status = MagickTrue;
 464     break;
 465
 466   case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED:
 467     (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "SetMagickOpenCLEnvParm cannot modify the OpenCL initialization state.", "'%s'", ".");
 468     break;
 469
 470   case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
 471     if (dataSize != sizeof(clEnv->disableProgramCache))
 472       goto cleanup;
 473     clEnv->disableProgramCache =  *((MagickBooleanType*)data);
 474     clEnv->OpenCLInitialized = MagickFalse;
 475     status = MagickTrue;
 476     break;
 477
 478   case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
 479     if (dataSize != sizeof(clEnv->regenerateProfile))
 480       goto cleanup;
 481     clEnv->regenerateProfile =  *((MagickBooleanType*)data);
 482     clEnv->OpenCLInitialized = MagickFalse;
 483     status = MagickTrue;
 484     break;
 485
 486   default:
 487     goto cleanup;
 488   };
 489
 490 cleanup:
 491   return status;
 492 }
 493
 494 MagickExport
 495   MagickBooleanType SetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
 496                                           , size_t dataSize, void* data, ExceptionInfo* exception) {
 497   MagickBooleanType status = MagickFalse;
 498   if (clEnv!=NULL) {
 499     LockSemaphoreInfo(clEnv->lock);
 500     status = SetMagickOpenCLEnvParamInternal(clEnv,param,dataSize,data,exception);
 501     UnlockSemaphoreInfo(clEnv->lock);
 502   }
 503   return status;
 504 }
 505
 506 /*
 507 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 508 %                                                                             %
 509 %                                                                             %
 510 %                                                                             %
 511 +   G e t M a g i c k O p e n C L E n v P a r a m                             %
 512 %                                                                             %
 513 %                                                                             %
 514 %                                                                             %
 515 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 516 %
 517 %  GetMagickOpenCLEnvParam() gets the parameters in the OpenCL environment
 518 %
 519 %  The format of the GetMagickOpenCLEnvParam() method is:
 520 %
 521 %      MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv,
 522 %        MagickOpenCLEnvParam param, size_t dataSize, void* data,
 523 %        ExceptionInfo* exception)
 524 %
 525 %  A description of each parameter follows:
 526 %
 527 %    o clEnv: the OpenCL environment.
 528 %
 529 %    o param: the parameter to be returned.
 530 %
 531 %    o dataSize: the data size of the parameter value.
 532 %
 533 %    o data:  the location where the returned parameter value will be stored
 534 %
 535 %    o exception: return any errors or warnings
 536 %
 537 */
 538
 539 MagickExport
 540   MagickBooleanType GetMagickOpenCLEnvParam(MagickCLEnv clEnv, MagickOpenCLEnvParam param
 541                                           , size_t dataSize, void* data, ExceptionInfo* exception)
 542 {
 543   MagickBooleanType
 544    status;
 545
 546   magick_unreferenced(exception);
 547
 548   status = MagickFalse;
 549
 550   if (clEnv == NULL
 551     || data == NULL)
 552     goto cleanup;
 553
 554   switch(param)
 555   {
 556   case MAGICK_OPENCL_ENV_PARAM_DEVICE:
 557     if (dataSize != sizeof(cl_device_id))
 558       goto cleanup;
 559     *((cl_device_id*)data) = clEnv->device;
 560     status = MagickTrue;
 561     break;
 562
 563   case MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED:
 564     if (dataSize != sizeof(clEnv->OpenCLDisabled))
 565       goto cleanup;
 566     *((MagickBooleanType*)data) = clEnv->OpenCLDisabled;
 567     status = MagickTrue;
 568     break;
 569
 570   case MAGICK_OPENCL_ENV_PARAM_OPENCL_INITIALIZED:
 571     if (dataSize != sizeof(clEnv->OpenCLDisabled))
 572       goto cleanup;
 573     *((MagickBooleanType*)data) = clEnv->OpenCLInitialized;
 574     status = MagickTrue;
 575     break;
 576
 577   case MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED:
 578     if (dataSize != sizeof(clEnv->disableProgramCache))
 579       goto cleanup;
 580     *((MagickBooleanType*)data) = clEnv->disableProgramCache;
 581     status = MagickTrue;
 582     break;
 583
 584   case MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE:
 585     if (dataSize != sizeof(clEnv->regenerateProfile))
 586       goto cleanup;
 587     *((MagickBooleanType*)data) = clEnv->regenerateProfile;
 588     status = MagickTrue;
 589     break;
 590
 591   default:
 592     goto cleanup;
 593   };
 594
 595 cleanup:
 596   return status;
 597 }
 598
 599
 600 /*
 601 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 602 %                                                                             %
 603 %                                                                             %
 604 %                                                                             %
 605 +   G e t O p e n C L C o n t e x t                                           %
 606 %                                                                             %
 607 %                                                                             %
 608 %                                                                             %
 609 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 610 %
 611 %  GetOpenCLContext() returns the OpenCL context
 612 %
 613 %  The format of the GetOpenCLContext() method is:
 614 %
 615 %      cl_context GetOpenCLContext(MagickCLEnv clEnv)
 616 %
 617 %  A description of each parameter follows:
 618 %
 619 %    o clEnv: OpenCL environment
 620 %
 621 */
 622
 623 MagickPrivate
 624 cl_context GetOpenCLContext(MagickCLEnv clEnv) {
 625   if (clEnv == NULL)
 626     return NULL;
 627   else
 628     return clEnv->context;
 629 }
 630
 631 static char* getBinaryCLProgramName(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
 632 {
 633   char* name;
 634   char* ptr;
 635   char path[MaxTextExtent];
 636   char deviceName[MaxTextExtent];
 637   const char* prefix = "magick_opencl";
 638   clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_NAME, MaxTextExtent, deviceName, NULL);
 639   ptr=deviceName;
 640   /* strip out illegal characters for file names */
 641   while (*ptr != '\0')
 642   {
 643     if ( *ptr == ' ' || *ptr == '\\' || *ptr == '/' || *ptr == ':' || *ptr == '*'
 644         || *ptr == '?' || *ptr == '"' || *ptr == '<' || *ptr == '>' || *ptr == '|')
 645     {
 646       *ptr = '_';
 647     }
 648     ptr++;
 649   }
 650   (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s_%s_%02d_%08x_%.20g.bin",
 651          GetOpenCLCachedFilesDirectory(),DirectorySeparator,prefix,deviceName,
 652          (unsigned int) prog,signature,(double) sizeof(char*)*8);
 653   name = (char*)AcquireMagickMemory(strlen(path)+1);
 654   CopyMagickString(name,path,strlen(path)+1);
 655   return name;
 656 }
 657
 658 static MagickBooleanType saveBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature, ExceptionInfo* exception)
 659 {
 660   MagickBooleanType saveSuccessful;
 661   cl_int clStatus;
 662   size_t binaryProgramSize;
 663   unsigned char* binaryProgram;
 664   char* binaryFileName;
 665   FILE* fileHandle;
 666
 667 #ifdef MAGICKCORE_CLPERFMARKER
 668   clBeginPerfMarkerAMD(__FUNCTION__,"");
 669 #endif
 670
 671   binaryProgram = NULL;
 672   binaryFileName = NULL;
 673   fileHandle = NULL;
 674   saveSuccessful = MagickFalse;
 675
 676   clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binaryProgramSize, NULL);
 677   if (clStatus != CL_SUCCESS)
 678   {
 679     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
 680     goto cleanup;
 681   }
 682
 683   binaryProgram = (unsigned char*) AcquireMagickMemory(binaryProgramSize);
 684   clStatus = clEnv->library->clGetProgramInfo(clEnv->programs[prog], CL_PROGRAM_BINARIES, sizeof(char*), &binaryProgram, NULL);
 685   if (clStatus != CL_SUCCESS)
 686   {
 687     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "clGetProgramInfo failed.", "'%s'", ".");
 688     goto cleanup;
 689   }
 690
 691   binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
 692   fileHandle = fopen(binaryFileName, "wb");
 693   if (fileHandle != NULL)
 694   {
 695     fwrite(binaryProgram, sizeof(char), binaryProgramSize, fileHandle);
 696     saveSuccessful = MagickTrue;
 697   }
 698   else
 699   {
 700     (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
 701       "Saving binary kernel failed.", "'%s'", ".");
 702   }
 703
 704 cleanup:
 705   if (fileHandle != NULL)
 706     fclose(fileHandle);
 707   if (binaryProgram != NULL)
 708     RelinquishMagickMemory(binaryProgram);
 709   if (binaryFileName != NULL)
 710     free(binaryFileName);
 711
 712 #ifdef MAGICKCORE_CLPERFMARKER
 713   clEndPerfMarkerAMD();
 714 #endif
 715
 716   return saveSuccessful;
 717 }
 718
 719 static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
 720 {
 721   MagickBooleanType loadSuccessful;
 722   unsigned char* binaryProgram;
 723   char* binaryFileName;
 724   FILE* fileHandle;
 725
 726 #ifdef MAGICKCORE_CLPERFMARKER
 727   clBeginPerfMarkerAMD(__FUNCTION__,"");
 728 #endif
 729
 730   binaryProgram = NULL;
 731   binaryFileName = NULL;
 732   fileHandle = NULL;
 733   loadSuccessful = MagickFalse;
 734
 735   binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
 736   fileHandle = fopen(binaryFileName, "rb");
 737   if (fileHandle != NULL)
 738   {
 739     int b_error;
 740     size_t length;
 741     cl_int clStatus;
 742     cl_int clBinaryStatus;
 743
 744     b_error = 0 ;
 745     length = 0;
 746     b_error |= fseek( fileHandle, 0, SEEK_END ) < 0;
 747     b_error |= ( length = ftell( fileHandle ) ) <= 0;
 748     b_error |= fseek( fileHandle, 0, SEEK_SET ) < 0;
 749     if( b_error )
 750       goto cleanup;
 751
 752     binaryProgram = (unsigned char*)AcquireMagickMemory(length);
 753     if (binaryProgram == NULL)
 754       goto cleanup;
 755
 756     memset(binaryProgram, 0, length);
 757     b_error |= fread(binaryProgram, 1, length, fileHandle) != length;
 758
 759     clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
 760     if (clStatus != CL_SUCCESS
 761         || clBinaryStatus != CL_SUCCESS)
 762       goto cleanup;
 763
 764     loadSuccessful = MagickTrue;
 765   }
 766
 767 cleanup:
 768   if (fileHandle != NULL)
 769     fclose(fileHandle);
 770   if (binaryFileName != NULL)
 771     free(binaryFileName);
 772   if (binaryProgram != NULL)
 773     RelinquishMagickMemory(binaryProgram);
 774
 775 #ifdef MAGICKCORE_CLPERFMARKER
 776   clEndPerfMarkerAMD();
 777 #endif
 778
 779   return loadSuccessful;
 780 }
 781
 782 static unsigned int stringSignature(const char* string)
 783 {
 784   unsigned int stringLength;
 785   unsigned int n,i,j;
 786   unsigned int signature;
 787   union
 788   {
 789     const char* s;
 790     const unsigned int* u;
 791   }p;
 792
 793 #ifdef MAGICKCORE_CLPERFMARKER
 794   clBeginPerfMarkerAMD(__FUNCTION__,"");
 795 #endif
 796
 797   stringLength = strlen(string);
 798   signature = stringLength;
 799   n = stringLength/sizeof(unsigned int);
 800   p.s = string;
 801   for (i = 0; i < n; i++)
 802   {
 803     signature^=p.u[i];
 804   }
 805   if (n * sizeof(unsigned int) != stringLength)
 806   {
 807     char padded[4];
 808     j = n * sizeof(unsigned int);
 809     for (i = 0; i < 4; i++,j++)
 810     {
 811       if (j < stringLength)
 812         padded[i] = p.s[j];
 813       else
 814         padded[i] = 0;
 815     }
 816     p.s = padded;
 817     signature^=p.u[0];
 818   }
 819
 820 #ifdef MAGICKCORE_CLPERFMARKER
 821   clEndPerfMarkerAMD();
 822 #endif
 823
 824   return signature;
 825 }
 826
 827 /* OpenCL kernels for accelerate.c */
 828 extern const char *accelerateKernels, *accelerateKernels2;
 829
 830 static MagickBooleanType CompileOpenCLKernels(MagickCLEnv clEnv, ExceptionInfo* exception)
 831 {
 832   MagickBooleanType status = MagickFalse;
 833   cl_int clStatus;
 834   unsigned int i;
 835   char* accelerateKernelsBuffer = NULL;
 836
 837   /* The index of the program strings in this array has to match the value of the enum MagickOpenCLProgram */
 838   const char* MagickOpenCLProgramStrings[MAGICK_OPENCL_NUM_PROGRAMS];
 839
 840   char options[MaxTextExtent];
 841   unsigned int optionsSignature;
 842
 843 #ifdef MAGICKCORE_CLPERFMARKER
 844   clBeginPerfMarkerAMD(__FUNCTION__,"");
 845 #endif
 846
 847   /* Get additional options */
 848   (void) FormatLocaleString(options, MaxTextExtent, CLOptions, (float)QuantumRange,
 849     (float)QuantumScale, (float)CLCharQuantumScale, (float)MagickEpsilon, (float)MagickPI, (unsigned int)MaxMap, (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
 850
 851   /*
 852   if (getenv("MAGICK_OCL_DEF"))
 853   {
 854     strcat(options," ");
 855     strcat(options,getenv("MAGICK_OCL_DEF"));
 856   }
 857   */
 858
 859   /*
 860   if (getenv("MAGICK_OCL_BUILD"))
 861     printf("options: %s\n", options);
 862   */
 863
 864   optionsSignature = stringSignature(options);
 865
 866   /* get all the OpenCL program strings here */
 867   accelerateKernelsBuffer = (char*) AcquireMagickMemory(strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
 868   sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
 869   MagickOpenCLProgramStrings[MAGICK_OPENCL_ACCELERATE] = accelerateKernelsBuffer;
 870
 871   for (i = 0; i < MAGICK_OPENCL_NUM_PROGRAMS; i++)
 872   {
 873     MagickBooleanType loadSuccessful = MagickFalse;
 874     unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature;
 875
 876     /* try to load the binary first */
 877     if (clEnv->disableProgramCache != MagickTrue
 878         && !getenv("MAGICK_OCL_REC"))
 879       loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature);
 880
 881     if (loadSuccessful == MagickFalse)
 882     {
 883       /* Binary CL program unavailable, compile the program from source */
 884       size_t programLength = strlen(MagickOpenCLProgramStrings[i]);
 885       clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
 886       if (clStatus!=CL_SUCCESS)
 887       {
 888         (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
 889           "clCreateProgramWithSource failed.", "(%d)", (int)clStatus);
 890
 891         goto cleanup;
 892       }
 893     }
 894
 895     clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
 896     if (clStatus!=CL_SUCCESS)
 897     {
 898       (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
 899         "clBuildProgram failed.", "(%d)", (int)clStatus);
 900
 901       if (loadSuccessful == MagickFalse)
 902       {
 903         char path[MaxTextExtent];
 904         FILE* fileHandle;
 905
 906         /*  dump the source into a file */
 907         (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
 908          ,GetOpenCLCachedFilesDirectory()
 909          ,DirectorySeparator,"magick_badcl.cl");
 910         fileHandle = fopen(path, "wb");
 911         if (fileHandle != NULL)
 912         {
 913           fwrite(MagickOpenCLProgramStrings[i], sizeof(char), strlen(MagickOpenCLProgramStrings[i]), fileHandle);
 914           fclose(fileHandle);
 915         }
 916
 917         /* dump the build log */
 918         {
 919           char* log;
 920           size_t logSize;
 921           clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
 922           log = (char*)AcquireMagickMemory(logSize);
 923           clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);
 924
 925           (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
 926            ,GetOpenCLCachedFilesDirectory()
 927            ,DirectorySeparator,"magick_badcl_build.log");
 928           fileHandle = fopen(path, "wb");
 929           if (fileHandle != NULL)
 930           {
 931             const char* buildOptionsTitle = "build options: ";
 932             fwrite(buildOptionsTitle, sizeof(char), strlen(buildOptionsTitle), fileHandle);
 933             fwrite(options, sizeof(char), strlen(options), fileHandle);
 934             fwrite("\n",sizeof(char), 1, fileHandle);
 935             fwrite(log, sizeof(char), logSize, fileHandle);
 936             fclose(fileHandle);
 937           }
 938           RelinquishMagickMemory(log);
 939         }
 940       }
 941       goto cleanup;
 942     }
 943
 944     if (loadSuccessful == MagickFalse)
 945     {
 946       /* Save the binary to a file to avoid re-compilation of the kernels in the future */
 947       saveBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception);
 948     }
 949
 950   }
 951   status = MagickTrue;
 952
 953 cleanup:
 954
 955   if (accelerateKernelsBuffer!=NULL) RelinquishMagickMemory(accelerateKernelsBuffer);
 956
 957 #ifdef MAGICKCORE_CLPERFMARKER
 958   clEndPerfMarkerAMD();
 959 #endif
 960
 961   return status;
 962 }
 963
 964 static MagickBooleanType InitOpenCLPlatformDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
 965   int i,j;
 966   cl_int status;
 967   cl_uint numPlatforms = 0;
 968   cl_platform_id *platforms = NULL;
 969   char* MAGICK_OCL_DEVICE = NULL;
 970   MagickBooleanType OpenCLAvailable = MagickFalse;
 971
 972 #ifdef MAGICKCORE_CLPERFMARKER
 973   clBeginPerfMarkerAMD(__FUNCTION__,"");
 974 #endif
 975
 976   /* check if there's an environment variable overriding the device selection */
 977   MAGICK_OCL_DEVICE = getenv("MAGICK_OCL_DEVICE");
 978   if (MAGICK_OCL_DEVICE != NULL)
 979   {
 980     if (strcmp(MAGICK_OCL_DEVICE, "CPU") == 0)
 981     {
 982       clEnv->deviceType = CL_DEVICE_TYPE_CPU;
 983     }
 984     else if (strcmp(MAGICK_OCL_DEVICE, "GPU") == 0)
 985     {
 986       clEnv->deviceType = CL_DEVICE_TYPE_GPU;
 987     }
 988     else if (strcmp(MAGICK_OCL_DEVICE, "OFF") == 0)
 989     {
 990       /* OpenCL disabled */
 991       goto cleanup;
 992     }
 993   }
 994   else if (clEnv->deviceType == 0) {
 995     clEnv->deviceType = CL_DEVICE_TYPE_ALL;
 996   }
 997
 998   if (clEnv->device != NULL)
 999   {
1000     status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
1001     if (status != CL_SUCCESS) {
1002       (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1003           "Failed to get OpenCL platform from the selected device.", "(%d)", status);
1004     }
1005     goto cleanup;
1006   }
1007   else if (clEnv->platform != NULL)
1008   {
1009     numPlatforms = 1;
1010     platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id));
1011     if (platforms == (cl_platform_id *) NULL)
1012     {
1013       (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1014         "AcquireMagickMemory failed.",".");
1015       goto cleanup;
1016     }
1017     platforms[0] = clEnv->platform;
1018   }
1019   else
1020   {
1021     clEnv->device = NULL;
1022
1023     /* Get the number of OpenCL platforms available */
1024     status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms);
1025     if (status != CL_SUCCESS)
1026     {
1027       (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1028         "clGetplatformIDs failed.", "(%d)", status);
1029       goto cleanup;
1030     }
1031
1032     /* No OpenCL available, just leave */
1033     if (numPlatforms == 0) {
1034       goto cleanup;
1035     }
1036
1037     platforms = (cl_platform_id *) AcquireMagickMemory(numPlatforms * sizeof(cl_platform_id));
1038     if (platforms == (cl_platform_id *) NULL)
1039     {
1040       (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1041         "AcquireMagickMemory failed.",".");
1042       goto cleanup;
1043     }
1044
1045     status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL);
1046     if (status != CL_SUCCESS)
1047     {
1048       (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1049         "clGetPlatformIDs failed.", "(%d)", status);
1050       goto cleanup;
1051     }
1052   }
1053
1054   /* Device selection */
1055   clEnv->device = NULL;
1056   for (j = 0; j < 2; j++)
1057   {
1058
1059     cl_device_type deviceType;
1060     if (clEnv->deviceType == CL_DEVICE_TYPE_ALL)
1061     {
1062       if (j == 0)
1063         deviceType = CL_DEVICE_TYPE_GPU;
1064       else
1065         deviceType = CL_DEVICE_TYPE_CPU;
1066     }
1067     else if (j == 1)
1068     {
1069       break;
1070     }
1071     else
1072       deviceType = clEnv->deviceType;
1073
1074     for (i = 0; i < numPlatforms; i++)
1075     {
1076       char version[MaxTextExtent];
1077       cl_uint numDevices;
1078       status = clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_VERSION, MaxTextExtent, version, NULL);
1079       if (status != CL_SUCCESS)
1080       {
1081         (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1082           "clGetPlatformInfo failed.", "(%d)", status);
1083         goto cleanup;
1084       }
1085       if (strncmp(version,"OpenCL 1.0 ",11) == 0)
1086         continue;
1087       status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
1088       if (status != CL_SUCCESS)
1089       {
1090         (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1091           "clGetDeviceIDs failed.", "(%d)", status);
1092         goto cleanup;
1093       }
1094       if (clEnv->device != NULL)
1095       {
1096         clEnv->platform = platforms[i];
1097   goto cleanup;
1098       }
1099     }
1100   }
1101
1102 cleanup:
1103   if (platforms!=NULL)
1104     RelinquishMagickMemory(platforms);
1105
1106   OpenCLAvailable = (clEnv->platform!=NULL
1107           && clEnv->device!=NULL)?MagickTrue:MagickFalse;
1108
1109 #ifdef MAGICKCORE_CLPERFMARKER
1110   clEndPerfMarkerAMD();
1111 #endif
1112
1113   return OpenCLAvailable;
1114 }
1115
1116 static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) {
1117   if (clEnv->OpenCLInitialized != MagickFalse
1118     && clEnv->platform != NULL
1119     && clEnv->device != NULL) {
1120       clEnv->OpenCLDisabled = MagickFalse;
1121       return MagickTrue;
1122   }
1123   clEnv->OpenCLDisabled = MagickTrue;
1124   return MagickFalse;
1125 }
1126
1127
1128 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception);
1129 /*
1130 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1131 %                                                                             %
1132 %                                                                             %
1133 %                                                                             %
1134 +   I n i t O p e n C L E n v                                                 %
1135 %                                                                             %
1136 %                                                                             %
1137 %                                                                             %
1138 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1139 %
1140 %  InitOpenCLEnv() initialize the OpenCL environment
1141 %
1142 %  The format of the RelinquishMagickOpenCLEnv method is:
1143 %
1144 %      MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception)
1145 %
1146 %  A description of each parameter follows:
1147 %
1148 %    o clEnv: OpenCL environment structure
1149 %
1150 %    o exception: return any errors or warnings.
1151 %
1152 */
1153
1154 MagickExport
1155 MagickBooleanType InitOpenCLEnvInternal(MagickCLEnv clEnv, ExceptionInfo* exception) {
1156   MagickBooleanType status = MagickTrue;
1157   cl_int clStatus;
1158   cl_context_properties cps[3];
1159
1160 #ifdef MAGICKCORE_CLPERFMARKER
1161   {
1162     int status = clInitializePerfMarkerAMD();
1163     if (status == AP_SUCCESS) {
1164       //printf("PerfMarker successfully initialized\n");
1165     }
1166   }
1167 #endif
1168   clEnv->OpenCLInitialized = MagickTrue;
1169
1170   /* check and init the global lib */
1171   OpenCLLib=GetOpenCLLib();
1172   if (OpenCLLib)
1173   {
1174     clEnv->library=OpenCLLib;
1175   }
1176   else
1177   {
1178     /* turn off opencl */
1179     MagickBooleanType flag;
1180     flag = MagickTrue;
1181     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
1182         , sizeof(MagickBooleanType), &flag, exception);
1183   }
1184
1185   if (clEnv->OpenCLDisabled != MagickFalse)
1186     goto cleanup;
1187
1188   clEnv->OpenCLDisabled = MagickTrue;
1189   /* setup the OpenCL platform and device */
1190   status = InitOpenCLPlatformDevice(clEnv, exception);
1191   if (status == MagickFalse) {
1192     /* No OpenCL device available */
1193     goto cleanup;
1194   }
1195
1196   /* create an OpenCL context */
1197   cps[0] = CL_CONTEXT_PLATFORM;
1198   cps[1] = (cl_context_properties)clEnv->platform;
1199   cps[2] = 0;
1200   clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
1201   if (clStatus != CL_SUCCESS)
1202   {
1203     (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1204         "clCreateContext failed.", "(%d)", clStatus);
1205     status = MagickFalse;
1206     goto cleanup;
1207   }
1208
1209   status = CompileOpenCLKernels(clEnv, exception);
1210   if (status == MagickFalse) {
1211    (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1212         "clCreateCommandQueue failed.", "(%d)", status);
1213
1214     status = MagickFalse;
1215     goto cleanup;
1216   }
1217
1218   status = EnableOpenCLInternal(clEnv);
1219
1220 cleanup:
1221   return status;
1222 }
1223
1224
1225 MagickExport
1226 MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) {
1227   MagickBooleanType status = MagickFalse;
1228
1229   if (clEnv == NULL)
1230     return MagickFalse;
1231
1232 #ifdef MAGICKCORE_CLPERFMARKER
1233   clBeginPerfMarkerAMD(__FUNCTION__,"");
1234 #endif
1235
1236   LockSemaphoreInfo(clEnv->lock);
1237   if (clEnv->OpenCLInitialized == MagickFalse) {
1238     if (clEnv->device==NULL
1239         && clEnv->OpenCLDisabled == MagickFalse)
1240       status = autoSelectDevice(clEnv, exception);
1241     else
1242       status = InitOpenCLEnvInternal(clEnv, exception);
1243   }
1244   UnlockSemaphoreInfo(clEnv->lock);
1245
1246 #ifdef MAGICKCORE_CLPERFMARKER
1247   clEndPerfMarkerAMD();
1248 #endif
1249   return status;
1250 }
1251
1252
1253 /*
1254 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1255 %                                                                             %
1256 %                                                                             %
1257 %                                                                             %
1258 +   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                         %
1259 %                                                                             %
1260 %                                                                             %
1261 %                                                                             %
1262 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1263 %
1264 %  AcquireOpenCLCommandQueue() acquires an OpenCL command queue
1265 %
1266 %  The format of the AcquireOpenCLCommandQueue method is:
1267 %
1268 %      cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1269 %
1270 %  A description of each parameter follows:
1271 %
1272 %    o clEnv: the OpenCL environment.
1273 %
1274 */
1275
1276 MagickPrivate
1277 cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1278 {
1279   if (clEnv != NULL)
1280     return clEnv->library->clCreateCommandQueue(clEnv->context, clEnv->device, 0, NULL);
1281   else
1282     return NULL;
1283 }
1284
1285
1286 /*
1287 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1288 %                                                                             %
1289 %                                                                             %
1290 %                                                                             %
1291 +   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                   %
1292 %                                                                             %
1293 %                                                                             %
1294 %                                                                             %
1295 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1296 %
1297 %  RelinquishOpenCLCommandQueue() releases the OpenCL command queue
1298 %
1299 %  The format of the RelinquishOpenCLCommandQueue method is:
1300 %
1301 %      MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
1302 %        cl_command_queue queue)
1303 %
1304 %  A description of each parameter follows:
1305 %
1306 %    o clEnv: the OpenCL environment.
1307 %
1308 %    o queue: the OpenCL queue to be released.
1309 %
1310 %
1311 */
1312 MagickPrivate
1313 MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv, cl_command_queue queue)
1314 {
1315   if (clEnv != NULL)
1316   {
1317     return ((clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ? MagickTrue:MagickFalse);
1318   }
1319   else
1320     return MagickFalse;
1321 }
1322
1323
1324
1325 /*
1326 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1327 %                                                                             %
1328 %                                                                             %
1329 %                                                                             %
1330 +   A c q u i r e O p e n C L K e r n e l                                     %
1331 %                                                                             %
1332 %                                                                             %
1333 %                                                                             %
1334 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1335 %
1336 %  AcquireOpenCLKernel() acquires an OpenCL kernel
1337 %
1338 %  The format of the AcquireOpenCLKernel method is:
1339 %
1340 %      cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
1341 %        MagickOpenCLProgram program, const char* kernelName)
1342 %
1343 %  A description of each parameter follows:
1344 %
1345 %    o clEnv: the OpenCL environment.
1346 %
1347 %    o program: the OpenCL program module that the kernel belongs to.
1348 %
1349 %    o kernelName:  the name of the kernel
1350 %
1351 */
1352
1353 MagickPrivate
1354   cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv, MagickOpenCLProgram program, const char* kernelName)
1355 {
1356   cl_int clStatus;
1357   cl_kernel kernel = NULL;
1358   if (clEnv != NULL && kernelName!=NULL)
1359   {
1360     kernel = clEnv->library->clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
1361   }
1362   return kernel;
1363 }
1364
1365
1366 /*
1367 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1368 %                                                                             %
1369 %                                                                             %
1370 %                                                                             %
1371 +   R e l i n q u i s h O p e n C L K e r n e l                               %
1372 %                                                                             %
1373 %                                                                             %
1374 %                                                                             %
1375 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1376 %
1377 %  RelinquishOpenCLKernel() releases an OpenCL kernel
1378 %
1379 %  The format of the RelinquishOpenCLKernel method is:
1380 %
1381 %    MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv,
1382 %      cl_kernel kernel)
1383 %
1384 %  A description of each parameter follows:
1385 %
1386 %    o clEnv: the OpenCL environment.
1387 %
1388 %    o kernel: the OpenCL kernel object to be released.
1389 %
1390 %
1391 */
1392
1393 MagickPrivate
1394   MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel)
1395 {
1396   MagickBooleanType status = MagickFalse;
1397   if (clEnv != NULL && kernel != NULL)
1398   {
1399     status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
1400   }
1401   return status;
1402 }
1403
1404 /*
1405 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1406 %                                                                             %
1407 %                                                                             %
1408 %                                                                             %
1409 +   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               %
1410 %                                                                             %
1411 %                                                                             %
1412 %                                                                             %
1413 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1414 %
1415 %  GetOpenCLDeviceLocalMemorySize() returns local memory size of the device
1416 %
1417 %  The format of the GetOpenCLDeviceLocalMemorySize method is:
1418 %
1419 %    unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1420 %
1421 %  A description of each parameter follows:
1422 %
1423 %    o clEnv: the OpenCL environment.
1424 %
1425 %
1426 */
1427
1428 MagickPrivate
1429  unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1430 {
1431   cl_ulong localMemorySize;
1432   clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
1433   return (unsigned long)localMemorySize;
1434 }
1435
1436 MagickPrivate
1437   unsigned long GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv clEnv)
1438 {
1439   cl_ulong maxMemAllocSize;
1440   clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
1441   return (unsigned long)maxMemAllocSize;
1442 }
1443
1444
1445 /*
1446  Beginning of the OpenCL device selection infrastructure
1447 */
1448
1449
1450 typedef enum {
1451   DS_SUCCESS = 0
1452  ,DS_INVALID_PROFILE = 1000
1453  ,DS_MEMORY_ERROR
1454  ,DS_INVALID_PERF_EVALUATOR_TYPE
1455  ,DS_INVALID_PERF_EVALUATOR
1456  ,DS_PERF_EVALUATOR_ERROR
1457  ,DS_FILE_ERROR
1458  ,DS_UNKNOWN_DEVICE_TYPE
1459  ,DS_PROFILE_FILE_ERROR
1460  ,DS_SCORE_SERIALIZER_ERROR
1461  ,DS_SCORE_DESERIALIZER_ERROR
1462 } ds_status;
1463
1464 /* device type */
1465 typedef enum {
1466   DS_DEVICE_NATIVE_CPU = 0
1467  ,DS_DEVICE_OPENCL_DEVICE
1468 } ds_device_type;
1469
1470
1471 typedef struct {
1472   ds_device_type  type;
1473   cl_device_id    oclDeviceID;
1474   char*           oclDeviceName;
1475   char*           oclDriverVersion;
1476   cl_uint         oclMaxClockFrequency;
1477   cl_uint         oclMaxComputeUnits;
1478   void*           score;            /* a pointer to the score data, the content/format is application defined */
1479 } ds_device;
1480
1481 typedef struct {
1482   unsigned int  numDevices;
1483   ds_device*    devices;
1484   const char*   version;
1485 } ds_profile;
1486
1487 /* deallocate memory used by score */
1488 typedef ds_status (*ds_score_release)(void* score);
1489
1490 static ds_status releaseDeviceResource(ds_device* device, ds_score_release sr) {
1491   ds_status status = DS_SUCCESS;
1492   if (device) {
1493     if (device->oclDeviceName)      free(device->oclDeviceName);
1494     if (device->oclDriverVersion)   free(device->oclDriverVersion);
1495     if (device->score)              status = sr(device->score);
1496   }
1497   return status;
1498 }
1499
1500 static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) {
1501   ds_status status = DS_SUCCESS;
1502   if (profile!=NULL) {
1503     if (profile->devices!=NULL && sr!=NULL) {
1504       unsigned int i;
1505       for (i = 0; i < profile->numDevices; i++) {
1506         status = releaseDeviceResource(profile->devices+i,sr);
1507         if (status != DS_SUCCESS)
1508           break;
1509       }
1510       free(profile->devices);
1511     }
1512     free(profile);
1513   }
1514   return status;
1515 }
1516
1517
1518 static ds_status initDSProfile(ds_profile** p, const char* version) {
1519   int numDevices = 0;
1520   cl_uint numPlatforms = 0;
1521   cl_platform_id* platforms = NULL;
1522   cl_device_id*   devices = NULL;
1523   ds_status status = DS_SUCCESS;
1524   ds_profile* profile = NULL;
1525   unsigned int next = 0;
1526   unsigned int i;
1527
1528   if (p == NULL)
1529     return DS_INVALID_PROFILE;
1530
1531   profile = (ds_profile*)malloc(sizeof(ds_profile));
1532   if (profile == NULL)
1533     return DS_MEMORY_ERROR;
1534
1535   memset(profile, 0, sizeof(ds_profile));
1536
1537   OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
1538   if (numPlatforms > 0) {
1539     platforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id));
1540     if (platforms == NULL) {
1541       status = DS_MEMORY_ERROR;
1542       goto cleanup;
1543     }
1544     OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
1545     for (i = 0; i < (unsigned int)numPlatforms; i++) {
1546       cl_uint num;
1547       if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
1548         numDevices+=num;
1549     }
1550   }
1551
1552   profile->numDevices = numDevices+1;     /* +1 to numDevices to include the native CPU */
1553
1554   profile->devices = (ds_device*)malloc(profile->numDevices*sizeof(ds_device));
1555   if (profile->devices == NULL) {
1556     profile->numDevices = 0;
1557     status = DS_MEMORY_ERROR;
1558     goto cleanup;
1559   }
1560   memset(profile->devices, 0, profile->numDevices*sizeof(ds_device));
1561
1562   if (numDevices > 0) {
1563     devices = (cl_device_id*)malloc(numDevices*sizeof(cl_device_id));
1564     if (devices == NULL) {
1565       status = DS_MEMORY_ERROR;
1566       goto cleanup;
1567     }
1568     for (i = 0; i < (unsigned int)numPlatforms; i++) {
1569       cl_uint num;
1570
1571       int d;
1572       for (d = 0; d < 2; d++) {
1573         unsigned int j;
1574         cl_device_type deviceType;
1575         switch(d) {
1576         case 0:
1577           deviceType = CL_DEVICE_TYPE_GPU;
1578           break;
1579         case 1:
1580           deviceType = CL_DEVICE_TYPE_CPU;
1581           break;
1582         default:
1583           continue;
1584           break;
1585         }
1586         if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
1587           continue;
1588         for (j = 0; j < num; j++, next++) {
1589           size_t length;
1590
1591           profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
1592           profile->devices[next].oclDeviceID = devices[j];
1593
1594           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1595             , 0, NULL, &length);
1596           profile->devices[next].oclDeviceName = (char*)malloc(sizeof(char)*length);
1597           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1598             , length, profile->devices[next].oclDeviceName, NULL);
1599
1600           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1601             , 0, NULL, &length);
1602           profile->devices[next].oclDriverVersion = (char*)malloc(sizeof(char)*length);
1603           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1604             , length, profile->devices[next].oclDriverVersion, NULL);
1605
1606           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
1607             , sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);
1608
1609           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
1610             , sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);
1611         }
1612       }
1613     }
1614   }
1615
1616   profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
1617   profile->version = version;
1618
1619 cleanup:
1620   if (platforms)  free(platforms);
1621   if (devices)    free(devices);
1622   if (status == DS_SUCCESS) {
1623     *p = profile;
1624   }
1625   else {
1626     if (profile) {
1627       if (profile->devices)
1628         free(profile->devices);
1629       free(profile);
1630     }
1631   }
1632   return status;
1633 }
1634
1635 /* Pointer to a function that calculates the score of a device (ex: device->score)
1636  update the data size of score. The encoding and the format of the score data
1637  is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
1638  */
1639 typedef ds_status (*ds_perf_evaluator)(ds_device* device, void* data);
1640
1641 typedef enum {
1642   DS_EVALUATE_ALL
1643   ,DS_EVALUATE_NEW_ONLY
1644 } ds_evaluation_type;
1645
1646 static ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type
1647                          ,ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates) {
1648   ds_status status = DS_SUCCESS;
1649   unsigned int i;
1650   unsigned int updates = 0;
1651
1652   if (profile == NULL) {
1653     return DS_INVALID_PROFILE;
1654   }
1655   if (evaluator == NULL) {
1656     return DS_INVALID_PERF_EVALUATOR;
1657   }
1658
1659   for (i = 0; i < profile->numDevices; i++) {
1660     ds_status evaluatorStatus;
1661
1662     switch (type) {
1663     case DS_EVALUATE_NEW_ONLY:
1664       if (profile->devices[i].score != NULL)
1665         break;
1666       /*  else fall through */
1667     case DS_EVALUATE_ALL:
1668       evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
1669       if (evaluatorStatus != DS_SUCCESS) {
1670         status = evaluatorStatus;
1671         return status;
1672       }
1673       updates++;
1674       break;
1675     default:
1676       return DS_INVALID_PERF_EVALUATOR_TYPE;
1677       break;
1678     };
1679   }
1680   if (numUpdates)
1681     *numUpdates = updates;
1682   return status;
1683 }
1684
1685
1686 #define DS_TAG_VERSION                      "<version>"
1687 #define DS_TAG_VERSION_END                  "</version>"
1688 #define DS_TAG_DEVICE                       "<device>"
1689 #define DS_TAG_DEVICE_END                   "</device>"
1690 #define DS_TAG_SCORE                        "<score>"
1691 #define DS_TAG_SCORE_END                    "</score>"
1692 #define DS_TAG_DEVICE_TYPE                  "<type>"
1693 #define DS_TAG_DEVICE_TYPE_END              "</type>"
1694 #define DS_TAG_DEVICE_NAME                  "<name>"
1695 #define DS_TAG_DEVICE_NAME_END              "</name>"
1696 #define DS_TAG_DEVICE_DRIVER_VERSION        "<driver>"
1697 #define DS_TAG_DEVICE_DRIVER_VERSION_END    "</driver>"
1698 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS     "<max cu>"
1699 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END "</max cu>"
1700 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ        "<max clock>"
1701 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ_END    "</max clock>"
1702
1703 #define DS_DEVICE_NATIVE_CPU_STRING  "native_cpu"
1704
1705
1706
1707 typedef ds_status (*ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
1708 static ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file) {
1709   ds_status status = DS_SUCCESS;
1710   FILE* profileFile = NULL;
1711
1712
1713   if (profile == NULL)
1714     return DS_INVALID_PROFILE;
1715
1716   profileFile = fopen(file, "wb");
1717   if (profileFile==NULL) {
1718     status = DS_FILE_ERROR;
1719   }
1720   else {
1721     unsigned int i;
1722
1723     /* write version string */
1724     fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
1725     fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
1726     fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
1727     fwrite("\n", sizeof(char), 1, profileFile);
1728
1729     for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) {
1730       void* serializedScore;
1731       unsigned int serializedScoreSize;
1732
1733       fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);
1734
1735       fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
1736       fwrite(&profile->devices[i].type,sizeof(ds_device_type),1, profileFile);
1737       fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);
1738
1739       switch(profile->devices[i].type) {
1740       case DS_DEVICE_NATIVE_CPU:
1741         {
1742           /* There's no need to emit a device name for the native CPU device. */
1743           /*
1744           fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
1745           fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
1746           fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
1747           */
1748         }
1749         break;
1750       case DS_DEVICE_OPENCL_DEVICE:
1751         {
1752           char tmp[16];
1753
1754           fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
1755           fwrite(profile->devices[i].oclDeviceName,sizeof(char),strlen(profile->devices[i].oclDeviceName), profileFile);
1756           fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
1757
1758           fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
1759           fwrite(profile->devices[i].oclDriverVersion,sizeof(char),strlen(profile->devices[i].oclDriverVersion), profileFile);
1760           fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);
1761
1762           fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS), profileFile);
1763           sprintf(tmp,"%d",profile->devices[i].oclMaxComputeUnits);
1764           fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
1765           fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END), profileFile);
1766
1767           fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ), profileFile);
1768           sprintf(tmp,"%d",profile->devices[i].oclMaxClockFrequency);
1769           fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
1770           fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END), profileFile);
1771         }
1772         break;
1773       default:
1774         status = DS_UNKNOWN_DEVICE_TYPE;
1775         break;
1776       };
1777
1778       fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
1779       status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize);
1780       if (status == DS_SUCCESS && serializedScore!=NULL && serializedScoreSize > 0) {
1781         fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
1782         free(serializedScore);
1783       }
1784       fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
1785       fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
1786       fwrite("\n",sizeof(char),1,profileFile);
1787     }
1788     fclose(profileFile);
1789   }
1790   return status;
1791 }
1792
1793
1794 static ds_status readProFile(const char* fileName, char** content, size_t* contentSize) {
1795   ds_status status = DS_SUCCESS;
1796   FILE * input = NULL;
1797   size_t size = 0;
1798   size_t rsize = 0;
1799   char* binary = NULL;
1800
1801   *contentSize = 0;
1802   *content = NULL;
1803
1804   input = fopen(fileName, "rb");
1805   if(input == NULL) {
1806     return DS_FILE_ERROR;
1807   }
1808
1809   fseek(input, 0L, SEEK_END);
1810   size = ftell(input);
1811   rewind(input);
1812   binary = (char*)malloc(size);
1813   if(binary == NULL) {
1814     status = DS_FILE_ERROR;
1815     goto cleanup;
1816   }
1817   rsize = fread(binary, sizeof(char), size, input);
1818   if (rsize!=size
1819       || ferror(input)) {
1820     status = DS_FILE_ERROR;
1821     goto cleanup;
1822   }
1823   *contentSize = size;
1824   *content = binary;
1825
1826 cleanup:
1827   if (input != NULL) fclose(input);
1828   if (status != DS_SUCCESS
1829       && binary != NULL) {
1830       free(binary);
1831       *content = NULL;
1832       *contentSize = 0;
1833   }
1834   return status;
1835 }
1836
1837
1838 static const char* findString(const char* contentStart, const char* contentEnd, const char* string) {
1839   size_t stringLength;
1840   const char* currentPosition;
1841   const char* found;
1842   found = NULL;
1843   stringLength = strlen(string);
1844   currentPosition = contentStart;
1845   for(currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) {
1846     if (*currentPosition == string[0]) {
1847       if (currentPosition+stringLength < contentEnd) {
1848         if (strncmp(currentPosition, string, stringLength) == 0) {
1849           found = currentPosition;
1850           break;
1851         }
1852       }
1853     }
1854   }
1855   return found;
1856 }
1857
1858
1859 typedef ds_status (*ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
1860 static ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file) {
1861
1862   ds_status status = DS_SUCCESS;
1863   char* contentStart = NULL;
1864   const char* contentEnd = NULL;
1865   size_t contentSize;
1866
1867   if (profile==NULL)
1868     return DS_INVALID_PROFILE;
1869
1870   status = readProFile(file, &contentStart, &contentSize);
1871   if (status == DS_SUCCESS) {
1872     const char* currentPosition;
1873     const char* dataStart;
1874     const char* dataEnd;
1875     size_t versionStringLength;
1876
1877     contentEnd = contentStart + contentSize;
1878     currentPosition = contentStart;
1879
1880
1881     /* parse the version string */
1882     dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
1883     if (dataStart == NULL) {
1884       status = DS_PROFILE_FILE_ERROR;
1885       goto cleanup;
1886     }
1887     dataStart += strlen(DS_TAG_VERSION);
1888
1889     dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
1890     if (dataEnd==NULL) {
1891       status = DS_PROFILE_FILE_ERROR;
1892       goto cleanup;
1893     }
1894
1895     versionStringLength = strlen(profile->version);
1896     if (versionStringLength!=(size_t)(dataEnd-dataStart)
1897         || strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
1898       /* version mismatch */
1899       status = DS_PROFILE_FILE_ERROR;
1900       goto cleanup;
1901     }
1902     currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);
1903
1904     /* parse the device information */
1905 DisableMSCWarning(4127)
1906     while (1) {
1907 RestoreMSCWarning
1908       unsigned int i;
1909
1910       const char* deviceTypeStart;
1911       const char* deviceTypeEnd;
1912       ds_device_type deviceType;
1913
1914       const char* deviceNameStart;
1915       const char* deviceNameEnd;
1916
1917       const char* deviceScoreStart;
1918       const char* deviceScoreEnd;
1919
1920       const char* deviceDriverStart;
1921       const char* deviceDriverEnd;
1922
1923       const char* tmpStart;
1924       const char* tmpEnd;
1925       char tmp[16];
1926
1927       cl_uint maxClockFrequency;
1928       cl_uint maxComputeUnits;
1929
1930       dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
1931       if (dataStart==NULL) {
1932         /* nothing useful remain, quit...*/
1933         break;
1934       }
1935       dataStart+=strlen(DS_TAG_DEVICE);
1936       dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
1937       if (dataEnd==NULL) {
1938         status = DS_PROFILE_FILE_ERROR;
1939         goto cleanup;
1940       }
1941
1942       /* parse the device type */
1943       deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
1944       if (deviceTypeStart==NULL) {
1945         status = DS_PROFILE_FILE_ERROR;
1946         goto cleanup;
1947       }
1948       deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE);
1949       deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
1950       if (deviceTypeEnd==NULL) {
1951         status = DS_PROFILE_FILE_ERROR;
1952         goto cleanup;
1953       }
1954       memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));
1955
1956
1957       /* parse the device name */
1958       if (deviceType == DS_DEVICE_OPENCL_DEVICE) {
1959
1960         deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
1961         if (deviceNameStart==NULL) {
1962           status = DS_PROFILE_FILE_ERROR;
1963           goto cleanup;
1964         }
1965         deviceNameStart+=strlen(DS_TAG_DEVICE_NAME);
1966         deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
1967         if (deviceNameEnd==NULL) {
1968           status = DS_PROFILE_FILE_ERROR;
1969           goto cleanup;
1970         }
1971
1972
1973         deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
1974         if (deviceDriverStart==NULL) {
1975           status = DS_PROFILE_FILE_ERROR;
1976           goto cleanup;
1977         }
1978         deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION);
1979         deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
1980         if (deviceDriverEnd ==NULL) {
1981           status = DS_PROFILE_FILE_ERROR;
1982           goto cleanup;
1983         }
1984
1985
1986         tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
1987         if (tmpStart==NULL) {
1988           status = DS_PROFILE_FILE_ERROR;
1989           goto cleanup;
1990         }
1991         tmpStart+=strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
1992         tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END);
1993         if (tmpEnd ==NULL) {
1994           status = DS_PROFILE_FILE_ERROR;
1995           goto cleanup;
1996         }
1997         memcpy(tmp,tmpStart,tmpEnd-tmpStart);
1998         tmp[tmpEnd-tmpStart] = '\0';
1999         maxComputeUnits = atoi(tmp);
2000
2001
2002         tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2003         if (tmpStart==NULL) {
2004           status = DS_PROFILE_FILE_ERROR;
2005           goto cleanup;
2006         }
2007         tmpStart+=strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2008         tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ_END);
2009         if (tmpEnd ==NULL) {
2010           status = DS_PROFILE_FILE_ERROR;
2011           goto cleanup;
2012         }
2013         memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2014         tmp[tmpEnd-tmpStart] = '\0';
2015         maxClockFrequency = atoi(tmp);
2016
2017
2018         /* check if this device is on the system */
2019         for (i = 0; i < profile->numDevices; i++) {
2020           if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) {
2021             size_t actualDeviceNameLength;
2022             size_t driverVersionLength;
2023
2024             actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
2025             driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
2026             if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
2027                && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
2028                && maxComputeUnits == profile->devices[i].oclMaxComputeUnits
2029                && maxClockFrequency == profile->devices[i].oclMaxClockFrequency
2030                && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
2031                && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==(int)0) {
2032
2033               deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2034               if (deviceNameStart==NULL) {
2035                 status = DS_PROFILE_FILE_ERROR;
2036                 goto cleanup;
2037               }
2038               deviceScoreStart+=strlen(DS_TAG_SCORE);
2039               deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2040               status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2041               if (status != DS_SUCCESS) {
2042                 goto cleanup;
2043               }
2044             }
2045           }
2046         }
2047
2048       }
2049       else if (deviceType == DS_DEVICE_NATIVE_CPU) {
2050         for (i = 0; i < profile->numDevices; i++) {
2051           if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) {
2052             deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2053             if (deviceScoreStart==NULL) {
2054               status = DS_PROFILE_FILE_ERROR;
2055               goto cleanup;
2056             }
2057             deviceScoreStart+=strlen(DS_TAG_SCORE);
2058             deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2059             status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2060             if (status != DS_SUCCESS) {
2061               goto cleanup;
2062             }
2063           }
2064         }
2065       }
2066
2067       /* skip over the current one to find the next device */
2068       currentPosition = dataEnd+strlen(DS_TAG_DEVICE_END);
2069     }
2070   }
2071 cleanup:
2072   if (contentStart!=NULL) free(contentStart);
2073   return status;
2074 }
2075
2076
2077 #if 0
2078 static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
2079   unsigned int i;
2080   if (profile == NULL || num==NULL)
2081     return DS_MEMORY_ERROR;
2082   *num=0;
2083   for (i = 0; i < profile->numDevices; i++) {
2084     if (profile->devices[i].score == NULL) {
2085       (*num)++;
2086     }
2087   }
2088   return DS_SUCCESS;
2089 }
2090 #endif
2091
2092 /*
2093  End of the OpenCL device selection infrastructure
2094 */
2095
2096
2097 typedef double AccelerateScoreType;
2098
2099 static ds_status AcceleratePerfEvaluator(ds_device *device,
2100   void *magick_unused(data))
2101 {
2102 #define ACCELERATE_PERF_DIMEN "2048x1536"
2103 #define NUM_ITER  2
2104 #define ReturnStatus(status) \
2105 { \
2106   if (clEnv!=NULL) \
2107     RelinquishMagickOpenCLEnv(clEnv); \
2108   if (oldClEnv!=NULL) \
2109     defaultCLEnv = oldClEnv; \
2110   return status; \
2111 }
2112
2113   AccelerateTimer
2114     timer;
2115
2116   ExceptionInfo
2117     *exception=NULL;
2118
2119   MagickCLEnv
2120     clEnv=NULL,
2121     oldClEnv=NULL;
2122
2123   magick_unreferenced(data);
2124
2125   if (device == NULL)
2126     ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2127
2128   clEnv=AcquireMagickOpenCLEnv();
2129   exception=AcquireExceptionInfo();
2130
2131   if (device->type == DS_DEVICE_NATIVE_CPU)
2132     {
2133       /* CPU device */
2134       MagickBooleanType flag=MagickTrue;
2135       SetMagickOpenCLEnvParamInternal(clEnv,
2136         MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
2137         &flag,exception);
2138     }
2139   else if (device->type == DS_DEVICE_OPENCL_DEVICE)
2140     {
2141       /* OpenCL device */
2142       SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
2143         sizeof(cl_device_id),&device->oclDeviceID,exception);
2144     }
2145   else
2146     ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2147
2148   /* recompile the OpenCL kernels if it needs to */
2149   clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;
2150
2151   InitOpenCLEnvInternal(clEnv,exception);
2152   oldClEnv=defaultCLEnv;
2153   defaultCLEnv=clEnv;
2154
2155   /* microbenchmark */
2156   {
2157     Image
2158       *inputImage;
2159
2160     ImageInfo
2161       *imageInfo;
2162
2163     int
2164       i;
2165
2166     imageInfo=AcquireImageInfo();
2167     CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
2168     CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
2169     inputImage=ReadImage(imageInfo,exception);
2170
2171     initAccelerateTimer(&timer);
2172
2173     for (i=0; i<=NUM_ITER; i++)
2174     {
2175       Image
2176         *bluredImage,
2177         *resizedImage,
2178         *unsharpedImage;
2179
2180       if (i > 0)
2181         startAccelerateTimer(&timer);
2182
2183 #ifdef MAGICKCORE_CLPERFMARKER
2184       clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
2185 #endif
2186
2187       bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
2188       unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
2189         exception);
2190       resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,1.0,
2191         exception);
2192
2193 #ifdef MAGICKCORE_CLPERFMARKER
2194       clEndPerfMarkerAMD();
2195 #endif
2196
2197       if (i > 0)
2198         stopAccelerateTimer(&timer);
2199
2200       if (bluredImage)
2201         DestroyImage(bluredImage);
2202       if (unsharpedImage)
2203         DestroyImage(unsharpedImage);
2204       if (resizedImage)
2205         DestroyImage(resizedImage);
2206     }
2207     DestroyImage(inputImage);
2208   }
2209   /* end of microbenchmark */
2210
2211   if (device->score == NULL)
2212     device->score=malloc(sizeof(AccelerateScoreType));
2213   *(AccelerateScoreType*)device->score=readAccelerateTimer(&timer);
2214
2215   ReturnStatus(DS_SUCCESS);
2216 }
2217
2218 ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
2219   if (device
2220      && device->score) {
2221     /* generate a string from the score */
2222     char* s = (char*)malloc(sizeof(char)*256);
2223     sprintf(s,"%.4f",*((AccelerateScoreType*)device->score));
2224     *serializedScore = (void*)s;
2225     *serializedScoreSize = strlen(s);
2226     return DS_SUCCESS;
2227   }
2228   else {
2229     return DS_SCORE_SERIALIZER_ERROR;
2230   }
2231 }
2232
2233 ds_status AccelerateScoreDeserializer(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize) {
2234   if (device) {
2235     /* convert the string back to an int */
2236     char* s = (char*)malloc(serializedScoreSize+1);
2237     memcpy(s, serializedScore, serializedScoreSize);
2238     s[serializedScoreSize] = (char)'\0';
2239     device->score = malloc(sizeof(AccelerateScoreType));
2240     *((AccelerateScoreType*)device->score) = (AccelerateScoreType)atof(s);
2241     free(s);
2242     return DS_SUCCESS;
2243   }
2244   else {
2245     return DS_SCORE_DESERIALIZER_ERROR;
2246   }
2247 }
2248
2249 ds_status AccelerateScoreRelease(void* score) {
2250   if (score!=NULL) {
2251     free(score);
2252   }
2253   return DS_SUCCESS;
2254 }
2255
2256
2257 #define IMAGEMAGICK_PROFILE_VERSION "ImageMagick Device Selection v0.9"
2258 #define IMAGEMAGICK_PROFILE_FILE    "ImagemagickOpenCLDeviceProfile"
2259 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
2260
2261   MagickBooleanType mStatus = MagickFalse;
2262   ds_status status;
2263   ds_profile* profile;
2264   unsigned int numDeviceProfiled = 0;
2265   unsigned int i;
2266   unsigned int bestDeviceIndex;
2267   AccelerateScoreType bestScore;
2268   char path[MaxTextExtent];
2269   MagickBooleanType flag;
2270   ds_evaluation_type profileType;
2271
2272   LockDefaultOpenCLEnv();
2273
2274   /* Initially, just set OpenCL to off */
2275   flag = MagickTrue;
2276   SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2277     , sizeof(MagickBooleanType), &flag, exception);
2278
2279   /* check and init the global lib */
2280   OpenCLLib=GetOpenCLLib();
2281   if (OpenCLLib==NULL)
2282   {
2283     mStatus=InitOpenCLEnvInternal(clEnv, exception);
2284     goto cleanup;
2285   }
2286
2287   status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
2288   if (status!=DS_SUCCESS) {
2289     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2290     goto cleanup;
2291   }
2292
2293   (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2294          ,GetOpenCLCachedFilesDirectory()
2295          ,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
2296
2297   if (clEnv->regenerateProfile != MagickFalse) {
2298     profileType = DS_EVALUATE_ALL;
2299   }
2300   else {
2301     readProfileFromFile(profile, AccelerateScoreDeserializer, path);
2302     profileType = DS_EVALUATE_NEW_ONLY;
2303   }
2304   status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
2305
2306   if (status!=DS_SUCCESS) {
2307     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2308     goto cleanup;
2309   }
2310   if (numDeviceProfiled > 0) {
2311     status = writeProfileToFile(profile, AccelerateScoreSerializer, path);
2312     if (status!=DS_SUCCESS) {
2313       (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when saving the profile into a file", "'%s'", ".");
2314     }
2315   }
2316
2317   /* pick the best device */
2318   bestDeviceIndex = 0;
2319   bestScore = *(AccelerateScoreType*)profile->devices[bestDeviceIndex].score;
2320   for (i = 1; i < profile->numDevices; i++) {
2321     AccelerateScoreType score = *(AccelerateScoreType*)profile->devices[i].score;
2322     if (score < bestScore) {
2323       bestDeviceIndex = i;
2324       bestScore = score;
2325     }
2326   }
2327
2328   /* set up clEnv with the best device */
2329   if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
2330     /* CPU device */
2331     flag = MagickTrue;
2332     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2333                                   , sizeof(MagickBooleanType), &flag, exception);
2334   }
2335   else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
2336     /* OpenCL device */
2337     flag = MagickFalse;
2338     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2339       , sizeof(MagickBooleanType), &flag, exception);
2340     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2341       , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
2342   }
2343   else {
2344     status = DS_PERF_EVALUATOR_ERROR;
2345     goto cleanup;
2346   }
2347   mStatus=InitOpenCLEnvInternal(clEnv, exception);
2348
2349   status = releaseDSProfile(profile, AccelerateScoreRelease);
2350   if (status!=DS_SUCCESS) {
2351     (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
2352   }
2353
2354 cleanup:
2355
2356   UnlockDefaultOpenCLEnv();
2357   return mStatus;
2358 }
2359
2360
2361 /*
2362 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2363 %                                                                             %
2364 %                                                                             %
2365 %                                                                             %
2366 +   I n i t I m a g e M a g i c k O p e n C L                                 %
2367 %                                                                             %
2368 %                                                                             %
2369 %                                                                             %
2370 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2371 %
2372 %  InitImageMagickOpenCL() provides a simplified interface to initialize
2373 %  the OpenCL environtment in ImageMagick
2374 %
2375 %  The format of the InitImageMagickOpenCL() method is:
2376 %
2377 %      MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
2378 %                                        void* userSelectedDevice,
2379 %                                        void* selectedDevice)
2380 %
2381 %  A description of each parameter follows:
2382 %
2383 %    o mode: OpenCL mode in ImageMagick, could be off,auto,user
2384 %
2385 %    o userSelectedDevice:  when in user mode, a pointer to the selected
2386 %                           cl_device_id
2387 %
2388 %    o selectedDevice: a pointer to cl_device_id where the selected
2389 %                      cl_device_id by ImageMagick could be returned
2390 %
2391 %    o exception: exception
2392 %
2393 */
2394 MagickExport MagickBooleanType InitImageMagickOpenCL(
2395   ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
2396   ExceptionInfo *exception)
2397 {
2398   MagickBooleanType status = MagickFalse;
2399   MagickCLEnv clEnv = NULL;
2400   MagickBooleanType flag;
2401
2402   clEnv = GetDefaultOpenCLEnv();
2403   if (clEnv!=NULL) {
2404     switch(mode) {
2405
2406     case MAGICK_OPENCL_OFF:
2407       flag = MagickTrue;
2408       SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2409         , sizeof(MagickBooleanType), &flag, exception);
2410       status = InitOpenCLEnv(clEnv, exception);
2411
2412       if (selectedDevice)
2413         *(cl_device_id*)selectedDevice = NULL;
2414       break;
2415
2416     case MAGICK_OPENCL_DEVICE_SELECT_USER:
2417
2418       if (userSelectedDevice == NULL)
2419         return MagickFalse;
2420
2421       flag = MagickFalse;
2422       SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2423         , sizeof(MagickBooleanType), &flag, exception);
2424
2425       SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2426         , sizeof(cl_device_id), userSelectedDevice,exception);
2427
2428       status = InitOpenCLEnv(clEnv, exception);
2429       if (selectedDevice) {
2430         GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2431           , sizeof(cl_device_id), selectedDevice, exception);
2432       }
2433       break;
2434
2435     case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
2436         flag = MagickTrue;
2437         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
2438           , sizeof(MagickBooleanType), &flag, exception);
2439         flag = MagickTrue;
2440         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
2441           , sizeof(MagickBooleanType), &flag, exception);
2442
2443     /* fall through here!! */
2444     case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
2445     default:
2446       {
2447         cl_device_id d = NULL;
2448         flag = MagickFalse;
2449         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2450           , sizeof(MagickBooleanType), &flag, exception);
2451         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2452           , sizeof(cl_device_id), &d,exception);
2453         status = InitOpenCLEnv(clEnv, exception);
2454         if (selectedDevice) {
2455           GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2456             , sizeof(cl_device_id),  selectedDevice, exception);
2457         }
2458       }
2459       break;
2460     };
2461   }
2462   return status;
2463 }
2464
2465
2466 MagickPrivate
2467 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2468   const char *module,const char *function,const size_t line,
2469   const ExceptionType severity,const char *tag,const char *format,...) {
2470   MagickBooleanType
2471     status;
2472
2473   MagickCLEnv clEnv;
2474
2475   status = MagickTrue;
2476
2477   clEnv = GetDefaultOpenCLEnv();
2478
2479   assert(exception != (ExceptionInfo *) NULL);
2480   assert(exception->signature == MagickSignature);
2481
2482   if (severity!=0) {
2483     cl_device_type dType;
2484     clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
2485     if (dType == CL_DEVICE_TYPE_CPU) {
2486       char buffer[MaxTextExtent];
2487       clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);
2488
2489       /* Workaround for Intel OpenCL CPU runtime bug */
2490       /* Turn off OpenCL when a problem is detected! */
2491       if (strncmp(buffer, "Intel",5) == 0) {
2492
2493         InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
2494       }
2495     }
2496   }
2497
2498 #ifdef OPENCLLOG_ENABLED
2499   {
2500     va_list
2501       operands;
2502     va_start(operands,format);
2503     status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands);
2504     va_end(operands);
2505   }
2506 #else
2507   magick_unreferenced(module);
2508   magick_unreferenced(function);
2509   magick_unreferenced(line);
2510   magick_unreferenced(tag);
2511   magick_unreferenced(format);
2512 #endif
2513
2514   return(status);
2515 }
2516
2517 MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
2518 {
2519   LockSemaphoreInfo(clEnv->lock);
2520   if (clEnv->seedsLock == NULL)
2521   {
2522     ActivateSemaphoreInfo(&clEnv->seedsLock);
2523   }
2524   LockSemaphoreInfo(clEnv->seedsLock);
2525
2526   if (clEnv->seeds == NULL)
2527   {
2528     cl_int clStatus;
2529     clEnv->numGenerators = NUM_CL_RAND_GENERATORS;
2530     clEnv->seeds = clEnv->library->clCreateBuffer(clEnv->context, CL_MEM_READ_WRITE,
2531                                   clEnv->numGenerators*4*sizeof(unsigned int),
2532                                   NULL, &clStatus);
2533     if (clStatus != CL_SUCCESS)
2534     {
2535       clEnv->seeds = NULL;
2536     }
2537     else
2538     {
2539       unsigned int i;
2540       cl_command_queue queue = NULL;
2541       unsigned int *seeds;
2542
2543       queue = AcquireOpenCLCommandQueue(clEnv);
2544       seeds = (unsigned int*) clEnv->library->clEnqueueMapBuffer(queue, clEnv->seeds, CL_TRUE,
2545                                                   CL_MAP_WRITE, 0,
2546                                                   clEnv->numGenerators*4
2547                                                   *sizeof(unsigned int),
2548                                                   0, NULL, NULL, &clStatus);
2549       if (clStatus!=CL_SUCCESS)
2550       {
2551         clEnv->library->clReleaseMemObject(clEnv->seeds);
2552         goto cleanup;
2553       }
2554
2555       for (i = 0; i < clEnv->numGenerators; i++) {
2556         RandomInfo* randomInfo = AcquireRandomInfo();
2557         const unsigned long* s = GetRandomInfoSeed(randomInfo);
2558         if (i == 0)
2559           clEnv->randNormalize = GetRandomInfoNormalize(randomInfo);
2560
2561         seeds[i*4]   = (unsigned int) s[0];
2562         seeds[i*4+1] = (unsigned int) 0x50a7f451;
2563         seeds[i*4+2] = (unsigned int) 0x5365417e;
2564         seeds[i*4+3] = (unsigned int) 0xc3a4171a;
2565
2566         randomInfo = DestroyRandomInfo(randomInfo);
2567       }
2568       clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, clEnv->seeds, seeds, 0,
2569                                           NULL, NULL);
2570       clEnv->library->clFinish(queue);
2571 cleanup:
2572       if (queue != NULL)
2573         RelinquishOpenCLCommandQueue(clEnv, queue);
2574     }
2575   }
2576   UnlockSemaphoreInfo(clEnv->lock);
2577   return clEnv->seeds;
2578 }
2579
2580 MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv) {
2581   if (clEnv->seedsLock == NULL)
2582   {
2583     ActivateSemaphoreInfo(&clEnv->seedsLock);
2584   }
2585   else
2586     UnlockSemaphoreInfo(clEnv->seedsLock);
2587 }
2588
2589 MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
2590 {
2591   return clEnv->numGenerators;
2592 }
2593
2594
2595 MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
2596 {
2597   return clEnv->randNormalize;
2598 }
2599
2600 #else
2601
2602 struct _MagickCLEnv {
2603   MagickBooleanType OpenCLInitialized;  /* whether OpenCL environment is initialized. */
2604 };
2605
2606 MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
2607 {
2608   return NULL;
2609 }
2610
2611 MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(
2612   MagickCLEnv magick_unused(clEnv))
2613 {
2614   magick_unreferenced(clEnv);
2615
2616   return MagickFalse;
2617 }
2618
2619 /*
2620 * Return the OpenCL environment
2621 */
2622 MagickExport MagickCLEnv GetDefaultOpenCLEnv(
2623   ExceptionInfo *magick_unused(exception))
2624 {
2625   magick_unreferenced(exception);
2626
2627   return (MagickCLEnv) NULL;
2628 }
2629
2630 MagickExport MagickCLEnv SetDefaultOpenCLEnv(
2631   MagickCLEnv magick_unused(clEnv))
2632 {
2633   magick_unreferenced(clEnv);
2634
2635   return (MagickCLEnv) NULL;
2636 }
2637
2638 MagickExport MagickBooleanType SetMagickOpenCLEnvParam(
2639   MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
2640   size_t magick_unused(dataSize),void *magick_unused(data),
2641   ExceptionInfo *magick_unused(exception))
2642 {
2643   magick_unreferenced(clEnv);
2644   magick_unreferenced(param);
2645   magick_unreferenced(dataSize);
2646   magick_unreferenced(data);
2647   magick_unreferenced(exception);
2648
2649   return MagickFalse;
2650 }
2651
2652 MagickExport MagickBooleanType GetMagickOpenCLEnvParam(
2653   MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
2654   size_t magick_unused(dataSize),void *magick_unused(data),
2655   ExceptionInfo *magick_unused(exception))
2656 {
2657   magick_unreferenced(clEnv);
2658   magick_unreferenced(param);
2659   magick_unreferenced(dataSize);
2660   magick_unreferenced(data);
2661   magick_unreferenced(exception);
2662
2663   return MagickFalse;
2664 }
2665
2666 MagickExport MagickBooleanType InitOpenCLEnv(MagickCLEnv magick_unused(clEnv),
2667   ExceptionInfo *magick_unused(exception))
2668 {
2669   magick_unreferenced(clEnv);
2670   magick_unreferenced(exception);
2671
2672   return MagickFalse;
2673 }
2674
2675 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(
2676   MagickCLEnv magick_unused(clEnv))
2677 {
2678   magick_unreferenced(clEnv);
2679
2680   return (cl_command_queue) NULL;
2681 }
2682
2683 MagickPrivate MagickBooleanType RelinquishCommandQueue(
2684   MagickCLEnv magick_unused(clEnv),cl_command_queue magick_unused(queue))
2685 {
2686   magick_unreferenced(clEnv);
2687   magick_unreferenced(queue);
2688
2689   return MagickFalse;
2690 }
2691
2692 MagickPrivate cl_kernel AcquireOpenCLKernel(
2693   MagickCLEnv magick_unused(clEnv),MagickOpenCLProgram magick_unused(program),
2694   const char *magick_unused(kernelName))
2695 {
2696   magick_unreferenced(clEnv);
2697   magick_unreferenced(program);
2698   magick_unreferenced(kernelName);
2699
2700   return (cl_kernel)NULL;
2701 }
2702
2703 MagickPrivate MagickBooleanType RelinquishOpenCLKernel(
2704   MagickCLEnv magick_unused(clEnv),cl_kernel magick_unused(kernel))
2705 {
2706   magick_unreferenced(clEnv);
2707   magick_unreferenced(kernel);
2708
2709   return MagickFalse;
2710 }
2711
2712 MagickPrivate unsigned long GetOpenCLDeviceLocalMemorySize(
2713   MagickCLEnv magick_unused(clEnv))
2714 {
2715   magick_unreferenced(clEnv);
2716
2717   return 0;
2718 }
2719
2720 MagickExport MagickBooleanType InitImageMagickOpenCL(
2721   ImageMagickOpenCLMode magick_unused(mode),
2722   void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
2723   ExceptionInfo *magick_unused(exception))
2724 {
2725   magick_unreferenced(mode);
2726   magick_unreferenced(userSelectedDevice);
2727   magick_unreferenced(selectedDevice);
2728   magick_unreferenced(exception);
2729   return MagickFalse;
2730 }
2731
2732
2733 MagickPrivate
2734 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2735   const char *module,const char *function,const size_t line,
2736   const ExceptionType severity,const char *tag,const char *format,...)
2737 {
2738   magick_unreferenced(exception);
2739   magick_unreferenced(module);
2740   magick_unreferenced(function);
2741   magick_unreferenced(line);
2742   magick_unreferenced(severity);
2743   magick_unreferenced(tag);
2744   magick_unreferenced(format);
2745   return(MagickFalse);
2746 }
2747
2748
2749 MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
2750 {
2751   magick_unreferenced(clEnv);
2752   return NULL;
2753 }
2754
2755
2756 MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv)
2757 {
2758   magick_unreferenced(clEnv);
2759 }
2760
2761 MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
2762 {
2763   magick_unreferenced(clEnv);
2764   return 0;
2765 }
2766
2767 MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
2768 {
2769   magick_unreferenced(clEnv);
2770   return 0.0f;
2771 }
2772
2773 #endif /* MAGICKCORE_OPENCL_SUPPORT */
2774
2775 char* openclCachedFilesDirectory;
2776 SemaphoreInfo* openclCachedFilesDirectoryLock;
2777
2778 MagickPrivate
2779 const char* GetOpenCLCachedFilesDirectory() {
2780   if (openclCachedFilesDirectory == NULL) {
2781     if (openclCachedFilesDirectoryLock == NULL)
2782     {
2783       ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
2784     }
2785     LockSemaphoreInfo(openclCachedFilesDirectoryLock);
2786     if (openclCachedFilesDirectory == NULL) {
2787       char path[MaxTextExtent];
2788       char *home = NULL;
2789       char *temp = NULL;
2790       struct stat attributes;
2791       MagickBooleanType status;
2792
2793
2794
2795       home=GetEnvironmentValue("IMAGEMAGICK_OPENCL_CACHE_DIR");
2796       if (home == (char *) NULL)
2797       {
2798 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2799         home=GetEnvironmentValue("LOCALAPPDATA");
2800         if (home == (char *) NULL)
2801           home=GetEnvironmentValue("APPDATA");
2802         if (home == (char *) NULL)
2803           home=GetEnvironmentValue("USERPROFILE");
2804 #else
2805         home=GetEnvironmentValue("HOME");
2806 #endif
2807       }
2808
2809       if (home != (char *) NULL)
2810       {
2811         int mkdirStatus = 0;
2812         /*
2813         */
2814
2815         /* first check if $HOME/.config exists */
2816         (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config",
2817           home,DirectorySeparator);
2818         status=GetPathAttributes(path,&attributes);
2819         if (status == MagickFalse)
2820         {
2821
2822 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2823           mkdirStatus = mkdir(path);
2824 #else
2825           mkdirStatus = mkdir(path, 0777);
2826 #endif
2827         }
2828
2829         /* first check if $HOME/.config/ImageMagick exists */
2830         if (mkdirStatus==0)
2831         {
2832             (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config%sImageMagick",
2833               home,DirectorySeparator,DirectorySeparator);
2834
2835             status=GetPathAttributes(path,&attributes);
2836             if (status == MagickFalse)
2837             {
2838 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2839               mkdirStatus = mkdir(path);
2840 #else
2841               mkdirStatus = mkdir(path, 0777);
2842 #endif
2843             }
2844         }
2845
2846         if (mkdirStatus==0)
2847         {
2848           temp = (char*)AcquireMagickMemory(strlen(path)+1);
2849           CopyMagickString(temp,path,strlen(path)+1);
2850         }
2851         home=DestroyString(home);
2852       }
2853       openclCachedFilesDirectory = temp;
2854     }
2855     UnlockSemaphoreInfo(openclCachedFilesDirectoryLock);
2856   }
2857   return openclCachedFilesDirectory;
2858 }
2859
2860 void startAccelerateTimer(AccelerateTimer* timer) {
2861 #ifdef _WIN32
2862       QueryPerformanceCounter((LARGE_INTEGER*)&timer->_start);
2863
2864
2865 #else
2866       struct timeval s;
2867       gettimeofday(&s, 0);
2868       timer->_start = (long long)s.tv_sec * (long long)1.0E3 + (long long)s.tv_usec / (long long)1.0E3;
2869 #endif
2870 }
2871
2872 void stopAccelerateTimer(AccelerateTimer* timer) {
2873       long long n=0;
2874 #ifdef _WIN32
2875       QueryPerformanceCounter((LARGE_INTEGER*)&(n));
2876 #else
2877       struct timeval s;
2878       gettimeofday(&s, 0);
2879       n = (long long)s.tv_sec * (long long)1.0E3+ (long long)s.tv_usec / (long long)1.0E3;
2880 #endif
2881       n -= timer->_start;
2882       timer->_start = 0;
2883       timer->_clocks += n;
2884 }
2885
2886 void resetAccelerateTimer(AccelerateTimer* timer) {
2887    timer->_clocks = 0;
2888    timer->_start = 0;
2889 }
2890
2891
2892 void initAccelerateTimer(AccelerateTimer* timer) {
2893 #ifdef _WIN32
2894     QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq);
2895 #else
2896     timer->_freq = (long long)1.0E3;
2897 #endif
2898    resetAccelerateTimer(timer);
2899 }
2900
2901 double readAccelerateTimer(AccelerateTimer* timer) {
2902   return (double)timer->_clocks/(double)timer->_freq;
2903 };
2904
2905
2906 /* create a function for OpenCL log */
2907 MagickPrivate
2908 void OpenCLLog(const char* message) {
2909
2910 #ifdef OPENCLLOG_ENABLED
2911 #define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
2912
2913   FILE* log;
2914   if (getenv("MAGICK_OCL_LOG"))
2915   {
2916     if (message) {
2917       char path[MaxTextExtent];
2918       unsigned long allocSize;
2919
2920       MagickCLEnv clEnv;
2921
2922       clEnv = GetDefaultOpenCLEnv();
2923
2924       /*  dump the source into a file */
2925       (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2926         ,GetOpenCLCachedFilesDirectory()
2927         ,DirectorySeparator,OPENCL_LOG_FILE);
2928
2929
2930       log = fopen(path, "ab");
2931       fwrite(message, sizeof(char), strlen(message), log);
2932       fwrite("\n", sizeof(char), 1, log);
2933
2934       if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
2935       {
2936         allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
2937         fprintf(log, "Devic Max Memory Alloc Size: %ld\n", allocSize);
2938       }
2939
2940       fclose(log);
2941     }
2942   }
2943 #else
2944   magick_unreferenced(message);
2945 #endif
2946 }
2947
2948