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-2015 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     RelinquishSemaphoreInfo(&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 = (unsigned int) 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_type  oclDeviceType;
1474   cl_device_id    oclDeviceID;
1475   char*           oclDeviceName;
1476   char*           oclDriverVersion;
1477   cl_uint         oclMaxClockFrequency;
1478   cl_uint         oclMaxComputeUnits;
1479   void*           score;            /* a pointer to the score data, the content/format is application defined */
1480 } ds_device;
1481
1482 typedef struct {
1483   unsigned int  numDevices;
1484   ds_device*    devices;
1485   const char*   version;
1486 } ds_profile;
1487
1488 /* deallocate memory used by score */
1489 typedef ds_status (*ds_score_release)(void* score);
1490
1491 static ds_status releaseDeviceResource(ds_device* device, ds_score_release sr) {
1492   ds_status status = DS_SUCCESS;
1493   if (device) {
1494     if (device->oclDeviceName)      free(device->oclDeviceName);
1495     if (device->oclDriverVersion)   free(device->oclDriverVersion);
1496     if (device->score)              status = sr(device->score);
1497   }
1498   return status;
1499 }
1500
1501 static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) {
1502   ds_status status = DS_SUCCESS;
1503   if (profile!=NULL) {
1504     if (profile->devices!=NULL && sr!=NULL) {
1505       unsigned int i;
1506       for (i = 0; i < profile->numDevices; i++) {
1507         status = releaseDeviceResource(profile->devices+i,sr);
1508         if (status != DS_SUCCESS)
1509           break;
1510       }
1511       free(profile->devices);
1512     }
1513     free(profile);
1514   }
1515   return status;
1516 }
1517
1518
1519 static ds_status initDSProfile(ds_profile** p, const char* version) {
1520   int numDevices = 0;
1521   cl_uint numPlatforms = 0;
1522   cl_platform_id* platforms = NULL;
1523   cl_device_id*   devices = NULL;
1524   ds_status status = DS_SUCCESS;
1525   ds_profile* profile = NULL;
1526   unsigned int next = 0;
1527   unsigned int i;
1528
1529   if (p == NULL)
1530     return DS_INVALID_PROFILE;
1531
1532   profile = (ds_profile*)malloc(sizeof(ds_profile));
1533   if (profile == NULL)
1534     return DS_MEMORY_ERROR;
1535
1536   memset(profile, 0, sizeof(ds_profile));
1537
1538   OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
1539   if (numPlatforms > 0) {
1540     platforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id));
1541     if (platforms == NULL) {
1542       status = DS_MEMORY_ERROR;
1543       goto cleanup;
1544     }
1545     OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
1546     for (i = 0; i < (unsigned int)numPlatforms; i++) {
1547       cl_uint num;
1548       if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
1549         numDevices+=num;
1550     }
1551   }
1552
1553   profile->numDevices = numDevices+1;     /* +1 to numDevices to include the native CPU */
1554
1555   profile->devices = (ds_device*)malloc(profile->numDevices*sizeof(ds_device));
1556   if (profile->devices == NULL) {
1557     profile->numDevices = 0;
1558     status = DS_MEMORY_ERROR;
1559     goto cleanup;
1560   }
1561   memset(profile->devices, 0, profile->numDevices*sizeof(ds_device));
1562
1563   if (numDevices > 0) {
1564     devices = (cl_device_id*)malloc(numDevices*sizeof(cl_device_id));
1565     if (devices == NULL) {
1566       status = DS_MEMORY_ERROR;
1567       goto cleanup;
1568     }
1569     for (i = 0; i < (unsigned int)numPlatforms; i++) {
1570       cl_uint num;
1571
1572       int d;
1573       for (d = 0; d < 2; d++) {
1574         unsigned int j;
1575         cl_device_type deviceType;
1576         switch(d) {
1577         case 0:
1578           deviceType = CL_DEVICE_TYPE_GPU;
1579           break;
1580         case 1:
1581           deviceType = CL_DEVICE_TYPE_CPU;
1582           break;
1583         default:
1584           continue;
1585           break;
1586         }
1587         if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
1588           continue;
1589         for (j = 0; j < num; j++, next++) {
1590           size_t length;
1591
1592           profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
1593           profile->devices[next].oclDeviceID = devices[j];
1594
1595           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1596             , 0, NULL, &length);
1597           profile->devices[next].oclDeviceName = (char*)malloc(sizeof(char)*length);
1598           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1599             , length, profile->devices[next].oclDeviceName, NULL);
1600
1601           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1602             , 0, NULL, &length);
1603           profile->devices[next].oclDriverVersion = (char*)malloc(sizeof(char)*length);
1604           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1605             , length, profile->devices[next].oclDriverVersion, NULL);
1606
1607           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
1608             , sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);
1609
1610           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
1611             , sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);
1612
1613           OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_TYPE
1614             , sizeof(cl_device_type), &profile->devices[next].oclDeviceType, NULL);
1615         }
1616       }
1617     }
1618   }
1619
1620   profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
1621   profile->version = version;
1622
1623 cleanup:
1624   if (platforms)  free(platforms);
1625   if (devices)    free(devices);
1626   if (status == DS_SUCCESS) {
1627     *p = profile;
1628   }
1629   else {
1630     if (profile) {
1631       if (profile->devices)
1632         free(profile->devices);
1633       free(profile);
1634     }
1635   }
1636   return status;
1637 }
1638
1639 /* Pointer to a function that calculates the score of a device (ex: device->score)
1640  update the data size of score. The encoding and the format of the score data
1641  is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
1642  */
1643 typedef ds_status (*ds_perf_evaluator)(ds_device* device, void* data);
1644
1645 typedef enum {
1646   DS_EVALUATE_ALL
1647   ,DS_EVALUATE_NEW_ONLY
1648 } ds_evaluation_type;
1649
1650 static ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type
1651                          ,ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates) {
1652   ds_status status = DS_SUCCESS;
1653   unsigned int i;
1654   unsigned int updates = 0;
1655
1656   if (profile == NULL) {
1657     return DS_INVALID_PROFILE;
1658   }
1659   if (evaluator == NULL) {
1660     return DS_INVALID_PERF_EVALUATOR;
1661   }
1662
1663   for (i = 0; i < profile->numDevices; i++) {
1664     ds_status evaluatorStatus;
1665
1666     switch (type) {
1667     case DS_EVALUATE_NEW_ONLY:
1668       if (profile->devices[i].score != NULL)
1669         break;
1670       /*  else fall through */
1671     case DS_EVALUATE_ALL:
1672       evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
1673       if (evaluatorStatus != DS_SUCCESS) {
1674         status = evaluatorStatus;
1675         return status;
1676       }
1677       updates++;
1678       break;
1679     default:
1680       return DS_INVALID_PERF_EVALUATOR_TYPE;
1681       break;
1682     };
1683   }
1684   if (numUpdates)
1685     *numUpdates = updates;
1686   return status;
1687 }
1688
1689
1690 #define DS_TAG_VERSION                      "<version>"
1691 #define DS_TAG_VERSION_END                  "</version>"
1692 #define DS_TAG_DEVICE                       "<device>"
1693 #define DS_TAG_DEVICE_END                   "</device>"
1694 #define DS_TAG_SCORE                        "<score>"
1695 #define DS_TAG_SCORE_END                    "</score>"
1696 #define DS_TAG_DEVICE_TYPE                  "<type>"
1697 #define DS_TAG_DEVICE_TYPE_END              "</type>"
1698 #define DS_TAG_DEVICE_NAME                  "<name>"
1699 #define DS_TAG_DEVICE_NAME_END              "</name>"
1700 #define DS_TAG_DEVICE_DRIVER_VERSION        "<driver>"
1701 #define DS_TAG_DEVICE_DRIVER_VERSION_END    "</driver>"
1702 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS     "<max cu>"
1703 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END "</max cu>"
1704 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ        "<max clock>"
1705 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ_END    "</max clock>"
1706
1707 #define DS_DEVICE_NATIVE_CPU_STRING  "native_cpu"
1708
1709
1710
1711 typedef ds_status (*ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
1712 static ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file) {
1713   ds_status status = DS_SUCCESS;
1714   FILE* profileFile = NULL;
1715
1716
1717   if (profile == NULL)
1718     return DS_INVALID_PROFILE;
1719
1720   profileFile = fopen(file, "wb");
1721   if (profileFile==NULL) {
1722     status = DS_FILE_ERROR;
1723   }
1724   else {
1725     unsigned int i;
1726
1727     /* write version string */
1728     fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
1729     fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
1730     fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
1731     fwrite("\n", sizeof(char), 1, profileFile);
1732
1733     for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) {
1734       void* serializedScore;
1735       unsigned int serializedScoreSize;
1736
1737       fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);
1738
1739       fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
1740       fwrite(&profile->devices[i].type,sizeof(ds_device_type),1, profileFile);
1741       fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);
1742
1743       switch(profile->devices[i].type) {
1744       case DS_DEVICE_NATIVE_CPU:
1745         {
1746           /* There's no need to emit a device name for the native CPU device. */
1747           /*
1748           fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
1749           fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
1750           fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
1751           */
1752         }
1753         break;
1754       case DS_DEVICE_OPENCL_DEVICE:
1755         {
1756           char tmp[16];
1757
1758           fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
1759           fwrite(profile->devices[i].oclDeviceName,sizeof(char),strlen(profile->devices[i].oclDeviceName), profileFile);
1760           fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
1761
1762           fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
1763           fwrite(profile->devices[i].oclDriverVersion,sizeof(char),strlen(profile->devices[i].oclDriverVersion), profileFile);
1764           fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);
1765
1766           fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS), profileFile);
1767           sprintf(tmp,"%d",profile->devices[i].oclMaxComputeUnits);
1768           fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
1769           fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END), profileFile);
1770
1771           fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ), profileFile);
1772           sprintf(tmp,"%d",profile->devices[i].oclMaxClockFrequency);
1773           fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
1774           fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END), profileFile);
1775         }
1776         break;
1777       default:
1778         status = DS_UNKNOWN_DEVICE_TYPE;
1779         break;
1780       };
1781
1782       fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
1783       status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize);
1784       if (status == DS_SUCCESS && serializedScore!=NULL && serializedScoreSize > 0) {
1785         fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
1786         free(serializedScore);
1787       }
1788       fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
1789       fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
1790       fwrite("\n",sizeof(char),1,profileFile);
1791     }
1792     fclose(profileFile);
1793   }
1794   return status;
1795 }
1796
1797
1798 static ds_status readProFile(const char* fileName, char** content, size_t* contentSize) {
1799   ds_status status = DS_SUCCESS;
1800   FILE * input = NULL;
1801   size_t size = 0;
1802   size_t rsize = 0;
1803   char* binary = NULL;
1804
1805   *contentSize = 0;
1806   *content = NULL;
1807
1808   input = fopen(fileName, "rb");
1809   if(input == NULL) {
1810     return DS_FILE_ERROR;
1811   }
1812
1813   fseek(input, 0L, SEEK_END);
1814   size = ftell(input);
1815   rewind(input);
1816   binary = (char*)malloc(size);
1817   if(binary == NULL) {
1818     status = DS_FILE_ERROR;
1819     goto cleanup;
1820   }
1821   rsize = fread(binary, sizeof(char), size, input);
1822   if (rsize!=size
1823       || ferror(input)) {
1824     status = DS_FILE_ERROR;
1825     goto cleanup;
1826   }
1827   *contentSize = size;
1828   *content = binary;
1829
1830 cleanup:
1831   if (input != NULL) fclose(input);
1832   if (status != DS_SUCCESS
1833       && binary != NULL) {
1834       free(binary);
1835       *content = NULL;
1836       *contentSize = 0;
1837   }
1838   return status;
1839 }
1840
1841
1842 static const char* findString(const char* contentStart, const char* contentEnd, const char* string) {
1843   size_t stringLength;
1844   const char* currentPosition;
1845   const char* found;
1846   found = NULL;
1847   stringLength = strlen(string);
1848   currentPosition = contentStart;
1849   for(currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) {
1850     if (*currentPosition == string[0]) {
1851       if (currentPosition+stringLength < contentEnd) {
1852         if (strncmp(currentPosition, string, stringLength) == 0) {
1853           found = currentPosition;
1854           break;
1855         }
1856       }
1857     }
1858   }
1859   return found;
1860 }
1861
1862
1863 typedef ds_status (*ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
1864 static ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file) {
1865
1866   ds_status status = DS_SUCCESS;
1867   char* contentStart = NULL;
1868   const char* contentEnd = NULL;
1869   size_t contentSize;
1870
1871   if (profile==NULL)
1872     return DS_INVALID_PROFILE;
1873
1874   status = readProFile(file, &contentStart, &contentSize);
1875   if (status == DS_SUCCESS) {
1876     const char* currentPosition;
1877     const char* dataStart;
1878     const char* dataEnd;
1879     size_t versionStringLength;
1880
1881     contentEnd = contentStart + contentSize;
1882     currentPosition = contentStart;
1883
1884
1885     /* parse the version string */
1886     dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
1887     if (dataStart == NULL) {
1888       status = DS_PROFILE_FILE_ERROR;
1889       goto cleanup;
1890     }
1891     dataStart += strlen(DS_TAG_VERSION);
1892
1893     dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
1894     if (dataEnd==NULL) {
1895       status = DS_PROFILE_FILE_ERROR;
1896       goto cleanup;
1897     }
1898
1899     versionStringLength = strlen(profile->version);
1900     if (versionStringLength!=(size_t)(dataEnd-dataStart)
1901         || strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
1902       /* version mismatch */
1903       status = DS_PROFILE_FILE_ERROR;
1904       goto cleanup;
1905     }
1906     currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);
1907
1908     /* parse the device information */
1909 DisableMSCWarning(4127)
1910     while (1) {
1911 RestoreMSCWarning
1912       unsigned int i;
1913
1914       const char* deviceTypeStart;
1915       const char* deviceTypeEnd;
1916       ds_device_type deviceType;
1917
1918       const char* deviceNameStart;
1919       const char* deviceNameEnd;
1920
1921       const char* deviceScoreStart;
1922       const char* deviceScoreEnd;
1923
1924       const char* deviceDriverStart;
1925       const char* deviceDriverEnd;
1926
1927       const char* tmpStart;
1928       const char* tmpEnd;
1929       char tmp[16];
1930
1931       cl_uint maxClockFrequency;
1932       cl_uint maxComputeUnits;
1933
1934       dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
1935       if (dataStart==NULL) {
1936         /* nothing useful remain, quit...*/
1937         break;
1938       }
1939       dataStart+=strlen(DS_TAG_DEVICE);
1940       dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
1941       if (dataEnd==NULL) {
1942         status = DS_PROFILE_FILE_ERROR;
1943         goto cleanup;
1944       }
1945
1946       /* parse the device type */
1947       deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
1948       if (deviceTypeStart==NULL) {
1949         status = DS_PROFILE_FILE_ERROR;
1950         goto cleanup;
1951       }
1952       deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE);
1953       deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
1954       if (deviceTypeEnd==NULL) {
1955         status = DS_PROFILE_FILE_ERROR;
1956         goto cleanup;
1957       }
1958       memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));
1959
1960
1961       /* parse the device name */
1962       if (deviceType == DS_DEVICE_OPENCL_DEVICE) {
1963
1964         deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
1965         if (deviceNameStart==NULL) {
1966           status = DS_PROFILE_FILE_ERROR;
1967           goto cleanup;
1968         }
1969         deviceNameStart+=strlen(DS_TAG_DEVICE_NAME);
1970         deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
1971         if (deviceNameEnd==NULL) {
1972           status = DS_PROFILE_FILE_ERROR;
1973           goto cleanup;
1974         }
1975
1976
1977         deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
1978         if (deviceDriverStart==NULL) {
1979           status = DS_PROFILE_FILE_ERROR;
1980           goto cleanup;
1981         }
1982         deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION);
1983         deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
1984         if (deviceDriverEnd ==NULL) {
1985           status = DS_PROFILE_FILE_ERROR;
1986           goto cleanup;
1987         }
1988
1989
1990         tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
1991         if (tmpStart==NULL) {
1992           status = DS_PROFILE_FILE_ERROR;
1993           goto cleanup;
1994         }
1995         tmpStart+=strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
1996         tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END);
1997         if (tmpEnd ==NULL) {
1998           status = DS_PROFILE_FILE_ERROR;
1999           goto cleanup;
2000         }
2001         memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2002         tmp[tmpEnd-tmpStart] = '\0';
2003         maxComputeUnits = atoi(tmp);
2004
2005
2006         tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2007         if (tmpStart==NULL) {
2008           status = DS_PROFILE_FILE_ERROR;
2009           goto cleanup;
2010         }
2011         tmpStart+=strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2012         tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ_END);
2013         if (tmpEnd ==NULL) {
2014           status = DS_PROFILE_FILE_ERROR;
2015           goto cleanup;
2016         }
2017         memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2018         tmp[tmpEnd-tmpStart] = '\0';
2019         maxClockFrequency = atoi(tmp);
2020
2021
2022         /* check if this device is on the system */
2023         for (i = 0; i < profile->numDevices; i++) {
2024           if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) {
2025             size_t actualDeviceNameLength;
2026             size_t driverVersionLength;
2027
2028             actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
2029             driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
2030             if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
2031                && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
2032                && maxComputeUnits == profile->devices[i].oclMaxComputeUnits
2033                && maxClockFrequency == profile->devices[i].oclMaxClockFrequency
2034                && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
2035                && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==(int)0) {
2036
2037               deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2038               if (deviceNameStart==NULL) {
2039                 status = DS_PROFILE_FILE_ERROR;
2040                 goto cleanup;
2041               }
2042               deviceScoreStart+=strlen(DS_TAG_SCORE);
2043               deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2044               status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2045               if (status != DS_SUCCESS) {
2046                 goto cleanup;
2047               }
2048             }
2049           }
2050         }
2051
2052       }
2053       else if (deviceType == DS_DEVICE_NATIVE_CPU) {
2054         for (i = 0; i < profile->numDevices; i++) {
2055           if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) {
2056             deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2057             if (deviceScoreStart==NULL) {
2058               status = DS_PROFILE_FILE_ERROR;
2059               goto cleanup;
2060             }
2061             deviceScoreStart+=strlen(DS_TAG_SCORE);
2062             deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2063             status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2064             if (status != DS_SUCCESS) {
2065               goto cleanup;
2066             }
2067           }
2068         }
2069       }
2070
2071       /* skip over the current one to find the next device */
2072       currentPosition = dataEnd+strlen(DS_TAG_DEVICE_END);
2073     }
2074   }
2075 cleanup:
2076   if (contentStart!=NULL) free(contentStart);
2077   return status;
2078 }
2079
2080
2081 #if 0
2082 static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
2083   unsigned int i;
2084   if (profile == NULL || num==NULL)
2085     return DS_MEMORY_ERROR;
2086   *num=0;
2087   for (i = 0; i < profile->numDevices; i++) {
2088     if (profile->devices[i].score == NULL) {
2089       (*num)++;
2090     }
2091   }
2092   return DS_SUCCESS;
2093 }
2094 #endif
2095
2096 /*
2097  End of the OpenCL device selection infrastructure
2098 */
2099
2100
2101 typedef double AccelerateScoreType;
2102
2103 static ds_status AcceleratePerfEvaluator(ds_device *device,
2104   void *magick_unused(data))
2105 {
2106 #define ACCELERATE_PERF_DIMEN "2048x1536"
2107 #define NUM_ITER  2
2108 #define ReturnStatus(status) \
2109 { \
2110   if (clEnv!=NULL) \
2111     RelinquishMagickOpenCLEnv(clEnv); \
2112   if (oldClEnv!=NULL) \
2113     defaultCLEnv = oldClEnv; \
2114   return status; \
2115 }
2116
2117   AccelerateTimer
2118     timer;
2119
2120   ExceptionInfo
2121     *exception=NULL;
2122
2123   MagickCLEnv
2124     clEnv=NULL,
2125     oldClEnv=NULL;
2126
2127   magick_unreferenced(data);
2128
2129   if (device == NULL)
2130     ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2131
2132   clEnv=AcquireMagickOpenCLEnv();
2133   exception=AcquireExceptionInfo();
2134
2135   if (device->type == DS_DEVICE_NATIVE_CPU)
2136     {
2137       /* CPU device */
2138       MagickBooleanType flag=MagickTrue;
2139       SetMagickOpenCLEnvParamInternal(clEnv,
2140         MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
2141         &flag,exception);
2142     }
2143   else if (device->type == DS_DEVICE_OPENCL_DEVICE)
2144     {
2145       /* OpenCL device */
2146       SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
2147         sizeof(cl_device_id),&device->oclDeviceID,exception);
2148     }
2149   else
2150     ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2151
2152   /* recompile the OpenCL kernels if it needs to */
2153   clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;
2154
2155   InitOpenCLEnvInternal(clEnv,exception);
2156   oldClEnv=defaultCLEnv;
2157   defaultCLEnv=clEnv;
2158
2159   /* microbenchmark */
2160   {
2161     Image
2162       *inputImage;
2163
2164     ImageInfo
2165       *imageInfo;
2166
2167     int
2168       i;
2169
2170     imageInfo=AcquireImageInfo();
2171     CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
2172     CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
2173     inputImage=ReadImage(imageInfo,exception);
2174
2175     initAccelerateTimer(&timer);
2176
2177     for (i=0; i<=NUM_ITER; i++)
2178     {
2179       Image
2180         *bluredImage,
2181         *resizedImage,
2182         *unsharpedImage;
2183
2184       if (i > 0)
2185         startAccelerateTimer(&timer);
2186
2187 #ifdef MAGICKCORE_CLPERFMARKER
2188       clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
2189 #endif
2190
2191       bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
2192       unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
2193         exception);
2194       resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
2195         exception);
2196
2197 #ifdef MAGICKCORE_CLPERFMARKER
2198       clEndPerfMarkerAMD();
2199 #endif
2200
2201       if (i > 0)
2202         stopAccelerateTimer(&timer);
2203
2204       if (bluredImage)
2205         DestroyImage(bluredImage);
2206       if (unsharpedImage)
2207         DestroyImage(unsharpedImage);
2208       if (resizedImage)
2209         DestroyImage(resizedImage);
2210     }
2211     DestroyImage(inputImage);
2212   }
2213   /* end of microbenchmark */
2214
2215   if (device->score == NULL)
2216     device->score=malloc(sizeof(AccelerateScoreType));
2217   *(AccelerateScoreType*)device->score=readAccelerateTimer(&timer);
2218
2219   ReturnStatus(DS_SUCCESS);
2220 }
2221
2222 ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
2223   if (device
2224      && device->score) {
2225     /* generate a string from the score */
2226     char* s = (char*)malloc(sizeof(char)*256);
2227     sprintf(s,"%.4f",*((AccelerateScoreType*)device->score));
2228     *serializedScore = (void*)s;
2229     *serializedScoreSize = (unsigned int) strlen(s);
2230     return DS_SUCCESS;
2231   }
2232   else {
2233     return DS_SCORE_SERIALIZER_ERROR;
2234   }
2235 }
2236
2237 ds_status AccelerateScoreDeserializer(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize) {
2238   if (device) {
2239     /* convert the string back to an int */
2240     char* s = (char*)malloc(serializedScoreSize+1);
2241     memcpy(s, serializedScore, serializedScoreSize);
2242     s[serializedScoreSize] = (char)'\0';
2243     device->score = malloc(sizeof(AccelerateScoreType));
2244     *((AccelerateScoreType*)device->score) = (AccelerateScoreType)atof(s);
2245     free(s);
2246     return DS_SUCCESS;
2247   }
2248   else {
2249     return DS_SCORE_DESERIALIZER_ERROR;
2250   }
2251 }
2252
2253 ds_status AccelerateScoreRelease(void* score) {
2254   if (score!=NULL) {
2255     free(score);
2256   }
2257   return DS_SUCCESS;
2258 }
2259
2260 ds_status canWriteProfileToFile(const char *path)
2261 {
2262   FILE* profileFile = fopen(path, "ab");
2263
2264   if (profileFile==NULL)
2265     return DS_FILE_ERROR;
2266
2267   fclose(profileFile);
2268   return DS_SUCCESS;
2269 }
2270
2271 #define IMAGEMAGICK_PROFILE_VERSION "ImageMagick Device Selection v0.9"
2272 #define IMAGEMAGICK_PROFILE_FILE    "ImagemagickOpenCLDeviceProfile"
2273 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
2274
2275   MagickBooleanType mStatus = MagickFalse;
2276   ds_status status;
2277   ds_profile* profile;
2278   unsigned int numDeviceProfiled = 0;
2279   unsigned int i;
2280   unsigned int bestDeviceIndex;
2281   AccelerateScoreType bestScore;
2282   char path[MaxTextExtent];
2283   MagickBooleanType flag;
2284   ds_evaluation_type profileType;
2285
2286   LockDefaultOpenCLEnv();
2287
2288   /* Initially, just set OpenCL to off */
2289   flag = MagickTrue;
2290   SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2291     , sizeof(MagickBooleanType), &flag, exception);
2292
2293   /* check and init the global lib */
2294   OpenCLLib=GetOpenCLLib();
2295   if (OpenCLLib==NULL)
2296   {
2297     mStatus=InitOpenCLEnvInternal(clEnv, exception);
2298     goto cleanup;
2299   }
2300
2301   status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
2302   if (status!=DS_SUCCESS) {
2303     (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2304     goto cleanup;
2305   }
2306
2307   (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2308          ,GetOpenCLCachedFilesDirectory()
2309          ,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
2310
2311   if (canWriteProfileToFile(path) != DS_SUCCESS) {
2312     /* We can not write out a device profile, so don't run the benchmark */
2313     /* select the first GPU device */
2314
2315     bestDeviceIndex = 0;
2316     for (i = 1; i < profile->numDevices; i++) {
2317       if ((profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) && (profile->devices[i].oclDeviceType == CL_DEVICE_TYPE_GPU)) {
2318         bestDeviceIndex = i;
2319         break;
2320       }
2321     }
2322   }
2323   else {
2324     if (clEnv->regenerateProfile != MagickFalse) {
2325       profileType = DS_EVALUATE_ALL;
2326     }
2327     else {
2328       readProfileFromFile(profile, AccelerateScoreDeserializer, path);
2329       profileType = DS_EVALUATE_NEW_ONLY;
2330     }
2331     status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
2332
2333     if (status!=DS_SUCCESS) {
2334       (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2335       goto cleanup;
2336     }
2337     if (numDeviceProfiled > 0) {
2338       status = writeProfileToFile(profile, AccelerateScoreSerializer, path);
2339       if (status!=DS_SUCCESS) {
2340         (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when saving the profile into a file", "'%s'", ".");
2341       }
2342     }
2343
2344     /* pick the best device */
2345     bestDeviceIndex = 0;
2346     bestScore = *(AccelerateScoreType*)profile->devices[bestDeviceIndex].score;
2347     for (i = 1; i < profile->numDevices; i++) {
2348       AccelerateScoreType score = *(AccelerateScoreType*)profile->devices[i].score;
2349       if (score < bestScore) {
2350         bestDeviceIndex = i;
2351         bestScore = score;
2352       }
2353     }
2354   }
2355
2356   /* set up clEnv with the best device */
2357   if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
2358     /* CPU device */
2359     flag = MagickTrue;
2360     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2361                                   , sizeof(MagickBooleanType), &flag, exception);
2362   }
2363   else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
2364     /* OpenCL device */
2365     flag = MagickFalse;
2366     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2367       , sizeof(MagickBooleanType), &flag, exception);
2368     SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2369       , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
2370   }
2371   else {
2372     status = DS_PERF_EVALUATOR_ERROR;
2373     goto cleanup;
2374   }
2375   mStatus=InitOpenCLEnvInternal(clEnv, exception);
2376
2377   status = releaseDSProfile(profile, AccelerateScoreRelease);
2378   if (status!=DS_SUCCESS) {
2379     (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
2380   }
2381
2382 cleanup:
2383
2384   UnlockDefaultOpenCLEnv();
2385   return mStatus;
2386 }
2387
2388
2389 /*
2390 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2391 %                                                                             %
2392 %                                                                             %
2393 %                                                                             %
2394 +   I n i t I m a g e M a g i c k O p e n C L                                 %
2395 %                                                                             %
2396 %                                                                             %
2397 %                                                                             %
2398 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2399 %
2400 %  InitImageMagickOpenCL() provides a simplified interface to initialize
2401 %  the OpenCL environtment in ImageMagick
2402 %
2403 %  The format of the InitImageMagickOpenCL() method is:
2404 %
2405 %      MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
2406 %                                        void* userSelectedDevice,
2407 %                                        void* selectedDevice)
2408 %
2409 %  A description of each parameter follows:
2410 %
2411 %    o mode: OpenCL mode in ImageMagick, could be off,auto,user
2412 %
2413 %    o userSelectedDevice:  when in user mode, a pointer to the selected
2414 %                           cl_device_id
2415 %
2416 %    o selectedDevice: a pointer to cl_device_id where the selected
2417 %                      cl_device_id by ImageMagick could be returned
2418 %
2419 %    o exception: exception
2420 %
2421 */
2422 MagickExport MagickBooleanType InitImageMagickOpenCL(
2423   ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
2424   ExceptionInfo *exception)
2425 {
2426   MagickBooleanType status = MagickFalse;
2427   MagickCLEnv clEnv = NULL;
2428   MagickBooleanType flag;
2429
2430   clEnv = GetDefaultOpenCLEnv();
2431   if (clEnv!=NULL) {
2432     switch(mode) {
2433
2434     case MAGICK_OPENCL_OFF:
2435       flag = MagickTrue;
2436       SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2437         , sizeof(MagickBooleanType), &flag, exception);
2438       status = InitOpenCLEnv(clEnv, exception);
2439
2440       if (selectedDevice)
2441         *(cl_device_id*)selectedDevice = NULL;
2442       break;
2443
2444     case MAGICK_OPENCL_DEVICE_SELECT_USER:
2445
2446       if (userSelectedDevice == NULL)
2447         return MagickFalse;
2448
2449       flag = MagickFalse;
2450       SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2451         , sizeof(MagickBooleanType), &flag, exception);
2452
2453       SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2454         , sizeof(cl_device_id), userSelectedDevice,exception);
2455
2456       status = InitOpenCLEnv(clEnv, exception);
2457       if (selectedDevice) {
2458         GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2459           , sizeof(cl_device_id), selectedDevice, exception);
2460       }
2461       break;
2462
2463     case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
2464         flag = MagickTrue;
2465         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
2466           , sizeof(MagickBooleanType), &flag, exception);
2467         flag = MagickTrue;
2468         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
2469           , sizeof(MagickBooleanType), &flag, exception);
2470
2471     /* fall through here!! */
2472     case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
2473     default:
2474       {
2475         cl_device_id d = NULL;
2476         flag = MagickFalse;
2477         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2478           , sizeof(MagickBooleanType), &flag, exception);
2479         SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2480           , sizeof(cl_device_id), &d,exception);
2481         status = InitOpenCLEnv(clEnv, exception);
2482         if (selectedDevice) {
2483           GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2484             , sizeof(cl_device_id),  selectedDevice, exception);
2485         }
2486       }
2487       break;
2488     };
2489   }
2490   return status;
2491 }
2492
2493
2494 MagickPrivate
2495 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2496   const char *module,const char *function,const size_t line,
2497   const ExceptionType severity,const char *tag,const char *format,...) {
2498   MagickBooleanType
2499     status;
2500
2501   MagickCLEnv clEnv;
2502
2503   status = MagickTrue;
2504
2505   clEnv = GetDefaultOpenCLEnv();
2506
2507   assert(exception != (ExceptionInfo *) NULL);
2508   assert(exception->signature == MagickSignature);
2509
2510   if (severity!=0) {
2511     cl_device_type dType;
2512     clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
2513     if (dType == CL_DEVICE_TYPE_CPU) {
2514       char buffer[MaxTextExtent];
2515       clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);
2516
2517       /* Workaround for Intel OpenCL CPU runtime bug */
2518       /* Turn off OpenCL when a problem is detected! */
2519       if (strncmp(buffer, "Intel",5) == 0) {
2520
2521         InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
2522       }
2523     }
2524   }
2525
2526 #ifdef OPENCLLOG_ENABLED
2527   {
2528     va_list
2529       operands;
2530     va_start(operands,format);
2531     status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands);
2532     va_end(operands);
2533   }
2534 #else
2535   magick_unreferenced(module);
2536   magick_unreferenced(function);
2537   magick_unreferenced(line);
2538   magick_unreferenced(tag);
2539   magick_unreferenced(format);
2540 #endif
2541
2542   return(status);
2543 }
2544
2545 MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
2546 {
2547   LockSemaphoreInfo(clEnv->lock);
2548   if (clEnv->seedsLock == NULL)
2549   {
2550     ActivateSemaphoreInfo(&clEnv->seedsLock);
2551   }
2552   LockSemaphoreInfo(clEnv->seedsLock);
2553
2554   if (clEnv->seeds == NULL)
2555   {
2556     cl_int clStatus;
2557     clEnv->numGenerators = NUM_CL_RAND_GENERATORS;
2558     clEnv->seeds = clEnv->library->clCreateBuffer(clEnv->context, CL_MEM_READ_WRITE,
2559                                   clEnv->numGenerators*4*sizeof(unsigned int),
2560                                   NULL, &clStatus);
2561     if (clStatus != CL_SUCCESS)
2562     {
2563       clEnv->seeds = NULL;
2564     }
2565     else
2566     {
2567       unsigned int i;
2568       cl_command_queue queue = NULL;
2569       unsigned int *seeds;
2570
2571       queue = AcquireOpenCLCommandQueue(clEnv);
2572       seeds = (unsigned int*) clEnv->library->clEnqueueMapBuffer(queue, clEnv->seeds, CL_TRUE,
2573                                                   CL_MAP_WRITE, 0,
2574                                                   clEnv->numGenerators*4
2575                                                   *sizeof(unsigned int),
2576                                                   0, NULL, NULL, &clStatus);
2577       if (clStatus!=CL_SUCCESS)
2578       {
2579         clEnv->library->clReleaseMemObject(clEnv->seeds);
2580         goto cleanup;
2581       }
2582
2583       for (i = 0; i < clEnv->numGenerators; i++) {
2584         RandomInfo* randomInfo = AcquireRandomInfo();
2585         const unsigned long* s = GetRandomInfoSeed(randomInfo);
2586         if (i == 0)
2587           clEnv->randNormalize = GetRandomInfoNormalize(randomInfo);
2588
2589         seeds[i*4]   = (unsigned int) s[0];
2590         seeds[i*4+1] = (unsigned int) 0x50a7f451;
2591         seeds[i*4+2] = (unsigned int) 0x5365417e;
2592         seeds[i*4+3] = (unsigned int) 0xc3a4171a;
2593
2594         randomInfo = DestroyRandomInfo(randomInfo);
2595       }
2596       clStatus = clEnv->library->clEnqueueUnmapMemObject(queue, clEnv->seeds, seeds, 0,
2597                                           NULL, NULL);
2598       clEnv->library->clFinish(queue);
2599 cleanup:
2600       if (queue != NULL)
2601         RelinquishOpenCLCommandQueue(clEnv, queue);
2602     }
2603   }
2604   UnlockSemaphoreInfo(clEnv->lock);
2605   return clEnv->seeds;
2606 }
2607
2608 MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv) {
2609   if (clEnv->seedsLock == NULL)
2610   {
2611     ActivateSemaphoreInfo(&clEnv->seedsLock);
2612   }
2613   else
2614     UnlockSemaphoreInfo(clEnv->seedsLock);
2615 }
2616
2617 MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
2618 {
2619   return clEnv->numGenerators;
2620 }
2621
2622
2623 MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
2624 {
2625   return clEnv->randNormalize;
2626 }
2627
2628 #else
2629
2630 struct _MagickCLEnv {
2631   MagickBooleanType OpenCLInitialized;  /* whether OpenCL environment is initialized. */
2632 };
2633
2634 MagickExport MagickCLEnv AcquireMagickOpenCLEnv()
2635 {
2636   return NULL;
2637 }
2638
2639 MagickExport MagickBooleanType RelinquishMagickOpenCLEnv(
2640   MagickCLEnv magick_unused(clEnv))
2641 {
2642   magick_unreferenced(clEnv);
2643
2644   return MagickFalse;
2645 }
2646
2647 /*
2648 * Return the OpenCL environment
2649 */
2650 MagickExport MagickCLEnv GetDefaultOpenCLEnv(
2651   ExceptionInfo *magick_unused(exception))
2652 {
2653   magick_unreferenced(exception);
2654
2655   return (MagickCLEnv) NULL;
2656 }
2657
2658 MagickExport MagickCLEnv SetDefaultOpenCLEnv(
2659   MagickCLEnv magick_unused(clEnv))
2660 {
2661   magick_unreferenced(clEnv);
2662
2663   return (MagickCLEnv) NULL;
2664 }
2665
2666 MagickExport MagickBooleanType SetMagickOpenCLEnvParam(
2667   MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
2668   size_t magick_unused(dataSize),void *magick_unused(data),
2669   ExceptionInfo *magick_unused(exception))
2670 {
2671   magick_unreferenced(clEnv);
2672   magick_unreferenced(param);
2673   magick_unreferenced(dataSize);
2674   magick_unreferenced(data);
2675   magick_unreferenced(exception);
2676
2677   return MagickFalse;
2678 }
2679
2680 MagickExport MagickBooleanType GetMagickOpenCLEnvParam(
2681   MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
2682   size_t magick_unused(dataSize),void *magick_unused(data),
2683   ExceptionInfo *magick_unused(exception))
2684 {
2685   magick_unreferenced(clEnv);
2686   magick_unreferenced(param);
2687   magick_unreferenced(dataSize);
2688   magick_unreferenced(data);
2689   magick_unreferenced(exception);
2690
2691   return MagickFalse;
2692 }
2693
2694 MagickExport MagickBooleanType InitOpenCLEnv(MagickCLEnv magick_unused(clEnv),
2695   ExceptionInfo *magick_unused(exception))
2696 {
2697   magick_unreferenced(clEnv);
2698   magick_unreferenced(exception);
2699
2700   return MagickFalse;
2701 }
2702
2703 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(
2704   MagickCLEnv magick_unused(clEnv))
2705 {
2706   magick_unreferenced(clEnv);
2707
2708   return (cl_command_queue) NULL;
2709 }
2710
2711 MagickPrivate MagickBooleanType RelinquishCommandQueue(
2712   MagickCLEnv magick_unused(clEnv),cl_command_queue magick_unused(queue))
2713 {
2714   magick_unreferenced(clEnv);
2715   magick_unreferenced(queue);
2716
2717   return MagickFalse;
2718 }
2719
2720 MagickPrivate cl_kernel AcquireOpenCLKernel(
2721   MagickCLEnv magick_unused(clEnv),MagickOpenCLProgram magick_unused(program),
2722   const char *magick_unused(kernelName))
2723 {
2724   magick_unreferenced(clEnv);
2725   magick_unreferenced(program);
2726   magick_unreferenced(kernelName);
2727
2728   return (cl_kernel)NULL;
2729 }
2730
2731 MagickPrivate MagickBooleanType RelinquishOpenCLKernel(
2732   MagickCLEnv magick_unused(clEnv),cl_kernel magick_unused(kernel))
2733 {
2734   magick_unreferenced(clEnv);
2735   magick_unreferenced(kernel);
2736
2737   return MagickFalse;
2738 }
2739
2740 MagickPrivate unsigned long GetOpenCLDeviceLocalMemorySize(
2741   MagickCLEnv magick_unused(clEnv))
2742 {
2743   magick_unreferenced(clEnv);
2744
2745   return 0;
2746 }
2747
2748 MagickExport MagickBooleanType InitImageMagickOpenCL(
2749   ImageMagickOpenCLMode magick_unused(mode),
2750   void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
2751   ExceptionInfo *magick_unused(exception))
2752 {
2753   magick_unreferenced(mode);
2754   magick_unreferenced(userSelectedDevice);
2755   magick_unreferenced(selectedDevice);
2756   magick_unreferenced(exception);
2757   return MagickFalse;
2758 }
2759
2760
2761 MagickPrivate
2762 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2763   const char *module,const char *function,const size_t line,
2764   const ExceptionType severity,const char *tag,const char *format,...)
2765 {
2766   magick_unreferenced(exception);
2767   magick_unreferenced(module);
2768   magick_unreferenced(function);
2769   magick_unreferenced(line);
2770   magick_unreferenced(severity);
2771   magick_unreferenced(tag);
2772   magick_unreferenced(format);
2773   return(MagickFalse);
2774 }
2775
2776
2777 MagickPrivate cl_mem GetAndLockRandSeedBuffer(MagickCLEnv clEnv)
2778 {
2779   magick_unreferenced(clEnv);
2780   return NULL;
2781 }
2782
2783
2784 MagickPrivate void UnlockRandSeedBuffer(MagickCLEnv clEnv)
2785 {
2786   magick_unreferenced(clEnv);
2787 }
2788
2789 MagickPrivate unsigned int GetNumRandGenerators(MagickCLEnv clEnv)
2790 {
2791   magick_unreferenced(clEnv);
2792   return 0;
2793 }
2794
2795 MagickPrivate float GetRandNormalize(MagickCLEnv clEnv)
2796 {
2797   magick_unreferenced(clEnv);
2798   return 0.0f;
2799 }
2800
2801 #endif /* MAGICKCORE_OPENCL_SUPPORT */
2802
2803 char* openclCachedFilesDirectory;
2804 SemaphoreInfo* openclCachedFilesDirectoryLock;
2805
2806 MagickPrivate
2807 const char* GetOpenCLCachedFilesDirectory() {
2808   if (openclCachedFilesDirectory == NULL) {
2809     if (openclCachedFilesDirectoryLock == NULL)
2810     {
2811       ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
2812     }
2813     LockSemaphoreInfo(openclCachedFilesDirectoryLock);
2814     if (openclCachedFilesDirectory == NULL) {
2815       char path[MaxTextExtent];
2816       char *home = NULL;
2817       char *temp = NULL;
2818       struct stat attributes;
2819       MagickBooleanType status;
2820
2821
2822
2823       home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
2824       if (home == (char *) NULL)
2825       {
2826 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2827         home=GetEnvironmentValue("LOCALAPPDATA");
2828         if (home == (char *) NULL)
2829           home=GetEnvironmentValue("APPDATA");
2830         if (home == (char *) NULL)
2831           home=GetEnvironmentValue("USERPROFILE");
2832 #else
2833         home=GetEnvironmentValue("HOME");
2834 #endif
2835       }
2836
2837       if (home != (char *) NULL)
2838       {
2839         int mkdirStatus = 0;
2840         /*
2841         */
2842
2843         /* first check if $HOME/.config exists */
2844         (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config",
2845           home,DirectorySeparator);
2846         status=GetPathAttributes(path,&attributes);
2847         if (status == MagickFalse)
2848         {
2849
2850 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2851           mkdirStatus = mkdir(path);
2852 #else
2853           mkdirStatus = mkdir(path, 0777);
2854 #endif
2855         }
2856
2857         /* first check if $HOME/.config/ImageMagick exists */
2858         if (mkdirStatus==0)
2859         {
2860             (void) FormatLocaleString(path,MaxTextExtent,"%s%s.config%sImageMagick",
2861               home,DirectorySeparator,DirectorySeparator);
2862
2863             status=GetPathAttributes(path,&attributes);
2864             if (status == MagickFalse)
2865             {
2866 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2867               mkdirStatus = mkdir(path);
2868 #else
2869               mkdirStatus = mkdir(path, 0777);
2870 #endif
2871             }
2872         }
2873
2874         if (mkdirStatus==0)
2875         {
2876           temp = (char*)AcquireMagickMemory(strlen(path)+1);
2877           CopyMagickString(temp,path,strlen(path)+1);
2878         }
2879         home=DestroyString(home);
2880       }
2881       openclCachedFilesDirectory = temp;
2882     }
2883     UnlockSemaphoreInfo(openclCachedFilesDirectoryLock);
2884   }
2885   return openclCachedFilesDirectory;
2886 }
2887
2888 void startAccelerateTimer(AccelerateTimer* timer) {
2889 #ifdef _WIN32
2890       QueryPerformanceCounter((LARGE_INTEGER*)&timer->_start);
2891
2892
2893 #else
2894       struct timeval s;
2895       gettimeofday(&s, 0);
2896       timer->_start = (long long)s.tv_sec * (long long)1.0E3 + (long long)s.tv_usec / (long long)1.0E3;
2897 #endif
2898 }
2899
2900 void stopAccelerateTimer(AccelerateTimer* timer) {
2901       long long n=0;
2902 #ifdef _WIN32
2903       QueryPerformanceCounter((LARGE_INTEGER*)&(n));
2904 #else
2905       struct timeval s;
2906       gettimeofday(&s, 0);
2907       n = (long long)s.tv_sec * (long long)1.0E3+ (long long)s.tv_usec / (long long)1.0E3;
2908 #endif
2909       n -= timer->_start;
2910       timer->_start = 0;
2911       timer->_clocks += n;
2912 }
2913
2914 void resetAccelerateTimer(AccelerateTimer* timer) {
2915    timer->_clocks = 0;
2916    timer->_start = 0;
2917 }
2918
2919
2920 void initAccelerateTimer(AccelerateTimer* timer) {
2921 #ifdef _WIN32
2922     QueryPerformanceFrequency((LARGE_INTEGER*)&timer->_freq);
2923 #else
2924     timer->_freq = (long long)1.0E3;
2925 #endif
2926    resetAccelerateTimer(timer);
2927 }
2928
2929 double readAccelerateTimer(AccelerateTimer* timer) {
2930   return (double)timer->_clocks/(double)timer->_freq;
2931 };
2932
2933
2934 /* create a function for OpenCL log */
2935 MagickPrivate
2936 void OpenCLLog(const char* message) {
2937
2938 #ifdef OPENCLLOG_ENABLED
2939 #define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
2940
2941   FILE* log;
2942   if (getenv("MAGICK_OCL_LOG"))
2943   {
2944     if (message) {
2945       char path[MaxTextExtent];
2946       unsigned long allocSize;
2947
2948       MagickCLEnv clEnv;
2949
2950       clEnv = GetDefaultOpenCLEnv();
2951
2952       /*  dump the source into a file */
2953       (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2954         ,GetOpenCLCachedFilesDirectory()
2955         ,DirectorySeparator,OPENCL_LOG_FILE);
2956
2957
2958       log = fopen(path, "ab");
2959       fwrite(message, sizeof(char), strlen(message), log);
2960       fwrite("\n", sizeof(char), 1, log);
2961
2962       if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
2963       {
2964         allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
2965         fprintf(log, "Devic Max Memory Alloc Size: %ld\n", allocSize);
2966       }
2967
2968       fclose(log);
2969     }
2970   }
2971 #else
2972   magick_unreferenced(message);
2973 #endif
2974 }
2975
2976