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

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