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-2017 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 %    https://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 \f
  40 /*
  41   Include declarations.
  42 */
  43 #include "MagickCore/studio.h"
  44 #include "MagickCore/artifact.h"
  45 #include "MagickCore/cache.h"
  46 #include "MagickCore/cache-private.h"
  47 #include "MagickCore/color.h"
  48 #include "MagickCore/compare.h"
  49 #include "MagickCore/constitute.h"
  50 #include "MagickCore/configure.h"
  51 #include "MagickCore/distort.h"
  52 #include "MagickCore/draw.h"
  53 #include "MagickCore/effect.h"
  54 #include "MagickCore/exception.h"
  55 #include "MagickCore/exception-private.h"
  56 #include "MagickCore/fx.h"
  57 #include "MagickCore/gem.h"
  58 #include "MagickCore/geometry.h"
  59 #include "MagickCore/image.h"
  60 #include "MagickCore/image-private.h"
  61 #include "MagickCore/layer.h"
  62 #include "MagickCore/mime-private.h"
  63 #include "MagickCore/memory_.h"
  64 #include "MagickCore/monitor.h"
  65 #include "MagickCore/montage.h"
  66 #include "MagickCore/morphology.h"
  67 #include "MagickCore/nt-base.h"
  68 #include "MagickCore/nt-base-private.h"
  69 #include "MagickCore/opencl.h"
  70 #include "MagickCore/opencl-private.h"
  71 #include "MagickCore/option.h"
  72 #include "MagickCore/policy.h"
  73 #include "MagickCore/property.h"
  74 #include "MagickCore/quantize.h"
  75 #include "MagickCore/quantum.h"
  76 #include "MagickCore/random_.h"
  77 #include "MagickCore/random-private.h"
  78 #include "MagickCore/resample.h"
  79 #include "MagickCore/resource_.h"
  80 #include "MagickCore/splay-tree.h"
  81 #include "MagickCore/semaphore.h"
  82 #include "MagickCore/statistic.h"
  83 #include "MagickCore/string_.h"
  84 #include "MagickCore/string-private.h"
  85 #include "MagickCore/token.h"
  86 #include "MagickCore/utility.h"
  87 #include "MagickCore/utility-private.h"
  88
  89 #if defined(MAGICKCORE_OPENCL_SUPPORT)
  90
  91 #ifndef MAGICKCORE_WINDOWS_SUPPORT
  92 #include <dlfcn.h>
  93 #endif
  94
  95 #ifdef MAGICKCORE_HAVE_OPENCL_CL_H
  96 #define MAGICKCORE_OPENCL_MACOSX  1
  97 #endif
  98
  99 /*
 100   Define declarations.
 101 */
 102 #define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"
 103
 104 /*
 105   Typedef declarations.
 106 */
 107 typedef struct
 108 {
 109   long long freq;
 110   long long clocks;
 111   long long start;
 112 } AccelerateTimer;
 113
 114 typedef struct
 115 {
 116   char
 117     *name,
 118     *platform_name,
 119     *vendor_name,
 120     *version;
 121
 122   cl_uint
 123     max_clock_frequency,
 124     max_compute_units;
 125
 126   double
 127     score;
 128 } MagickCLDeviceBenchmark;
 129
 130 /*
 131   Forward declarations.
 132 */
 133
 134 static MagickBooleanType
 135   HasOpenCLDevices(MagickCLEnv,ExceptionInfo *),
 136   LoadOpenCLLibrary(void);
 137
 138 static MagickCLDevice
 139   RelinquishMagickCLDevice(MagickCLDevice);
 140
 141 static MagickCLEnv
 142   RelinquishMagickCLEnv(MagickCLEnv);
 143
 144 static void
 145   BenchmarkOpenCLDevices(MagickCLEnv);
 146
 147 extern const char
 148   *accelerateKernels, *accelerateKernels2;
 149
 150 /* OpenCL library */
 151 MagickLibrary
 152   *openCL_library;
 153
 154 /* Default OpenCL environment */
 155 MagickCLEnv
 156   default_CLEnv;
 157 MagickThreadType
 158   test_thread_id=0;
 159 SemaphoreInfo
 160   *openCL_lock;
 161
 162 /* Cached location of the OpenCL cache files */
 163 char
 164   *cache_directory;
 165 SemaphoreInfo
 166   *cache_directory_lock;
 167
 168 static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
 169   MagickCLDevice b)
 170 {
 171   if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
 172       (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
 173       (LocaleCompare(a->name,b->name) == 0) &&
 174       (LocaleCompare(a->version,b->version) == 0) &&
 175       (a->max_clock_frequency == b->max_clock_frequency) &&
 176       (a->max_compute_units == b->max_compute_units))
 177     return(MagickTrue);
 178
 179   return(MagickFalse);
 180 }
 181
 182 static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
 183   MagickCLDeviceBenchmark *b)
 184 {
 185   if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
 186       (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
 187       (LocaleCompare(a->name,b->name) == 0) &&
 188       (LocaleCompare(a->version,b->version) == 0) &&
 189       (a->max_clock_frequency == b->max_clock_frequency) &&
 190       (a->max_compute_units == b->max_compute_units))
 191     return(MagickTrue);
 192
 193   return(MagickFalse);
 194 }
 195
 196 static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
 197 {
 198   size_t
 199     i;
 200
 201   if (clEnv->devices != (MagickCLDevice *) NULL)
 202     {
 203       for (i = 0; i < clEnv->number_devices; i++)
 204         clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
 205       clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
 206     }
 207   clEnv->number_devices=0;
 208 }
 209
 210 static inline MagickBooleanType MagickCreateDirectory(const char *path)
 211 {
 212   int
 213     status;
 214
 215 #ifdef MAGICKCORE_WINDOWS_SUPPORT
 216   status=mkdir(path);
 217 #else
 218   status=mkdir(path, 0777);
 219 #endif
 220   return(status == 0 ? MagickTrue : MagickFalse);
 221 }
 222
 223 static inline void InitAccelerateTimer(AccelerateTimer *timer)
 224 {
 225 #ifdef _WIN32
 226   QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
 227 #else
 228   timer->freq=(long long)1.0E3;
 229 #endif
 230   timer->clocks=0;
 231   timer->start=0;
 232 }
 233
 234 static inline double ReadAccelerateTimer(AccelerateTimer *timer)
 235 {
 236   return (double)timer->clocks/(double)timer->freq;
 237 }
 238
 239 static inline void StartAccelerateTimer(AccelerateTimer* timer)
 240 {
 241 #ifdef _WIN32
 242   QueryPerformanceCounter((LARGE_INTEGER*)&timer->start);
 243 #else
 244   struct timeval
 245     s;
 246   gettimeofday(&s,0);
 247   timer->start=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
 248     (long long)1.0E3;
 249 #endif
 250 }
 251
 252 static inline void StopAccelerateTimer(AccelerateTimer *timer)
 253 {
 254   long long
 255     n;
 256
 257   n=0;
 258 #ifdef _WIN32
 259   QueryPerformanceCounter((LARGE_INTEGER*)&(n));
 260 #else
 261   struct timeval
 262     s;
 263   gettimeofday(&s,0);
 264   n=(long long)s.tv_sec*(long long)1.0E3+(long long)s.tv_usec/
 265     (long long)1.0E3;
 266 #endif
 267   n-=timer->start;
 268   timer->start=0;
 269   timer->clocks+=n;
 270 }
 271
 272 static const char *GetOpenCLCacheDirectory()
 273 {
 274   if (cache_directory == (char *) NULL)
 275     {
 276       if (cache_directory_lock == (SemaphoreInfo *) NULL)
 277         ActivateSemaphoreInfo(&cache_directory_lock);
 278       LockSemaphoreInfo(cache_directory_lock);
 279       if (cache_directory == (char *) NULL)
 280         {
 281           char
 282             *home,
 283             path[MagickPathExtent],
 284             *temp;
 285
 286           MagickBooleanType
 287             status;
 288
 289           struct stat
 290             attributes;
 291
 292           temp=(char *) NULL;
 293           home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
 294           if (home == (char *) NULL)
 295             {
 296               home=GetEnvironmentValue("XDG_CACHE_HOME");
 297               if (home == (char *) NULL)
 298                 home=GetEnvironmentValue("LOCALAPPDATA");
 299               if (home == (char *) NULL)
 300                 home=GetEnvironmentValue("APPDATA");
 301               if (home == (char *) NULL)
 302                 home=GetEnvironmentValue("USERPROFILE");
 303             }
 304
 305           if (home != (char *) NULL)
 306             {
 307               /* first check if $HOME exists */
 308               (void) FormatLocaleString(path,MagickPathExtent,"%s",home);
 309               status=GetPathAttributes(path,&attributes);
 310               if (status == MagickFalse)
 311                 status=MagickCreateDirectory(path);
 312
 313               /* first check if $HOME/ImageMagick exists */
 314               if (status != MagickFalse)
 315                 {
 316                   (void) FormatLocaleString(path,MagickPathExtent,
 317                     "%s%sImageMagick",home,DirectorySeparator);
 318
 319                   status=GetPathAttributes(path,&attributes);
 320                   if (status == MagickFalse)
 321                     status=MagickCreateDirectory(path);
 322                 }
 323
 324               if (status != MagickFalse)
 325                 {
 326                   temp=(char*) AcquireMagickMemory(strlen(path)+1);
 327                   CopyMagickString(temp,path,strlen(path)+1);
 328                 }
 329               home=DestroyString(home);
 330             }
 331           else
 332             {
 333               home=GetEnvironmentValue("HOME");
 334               if (home != (char *) NULL)
 335                 {
 336                   /* first check if $HOME/.cache exists */
 337                   (void) FormatLocaleString(path,MagickPathExtent,"%s%s.cache",
 338                     home,DirectorySeparator);
 339                   status=GetPathAttributes(path,&attributes);
 340                   if (status == MagickFalse)
 341                     status=MagickCreateDirectory(path);
 342
 343                   /* first check if $HOME/.cache/ImageMagick exists */
 344                   if (status != MagickFalse)
 345                     {
 346                       (void) FormatLocaleString(path,MagickPathExtent,
 347                         "%s%s.cache%sImageMagick",home,DirectorySeparator,
 348                         DirectorySeparator);
 349                       status=GetPathAttributes(path,&attributes);
 350                       if (status == MagickFalse)
 351                         status=MagickCreateDirectory(path);
 352                     }
 353
 354                   if (status != MagickFalse)
 355                     {
 356                       temp=(char*) AcquireMagickMemory(strlen(path)+1);
 357                       CopyMagickString(temp,path,strlen(path)+1);
 358                     }
 359                   home=DestroyString(home);
 360                 }
 361             }
 362           if (temp == (char *) NULL)
 363             temp=AcquireString("?");
 364           cache_directory=temp;
 365         }
 366       UnlockSemaphoreInfo(cache_directory_lock);
 367     }
 368   if (*cache_directory == '?')
 369     return((const char *) NULL);
 370   return(cache_directory);
 371 }
 372
 373 static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
 374 {
 375   MagickCLDevice
 376     device;
 377
 378   size_t
 379     i,
 380     j;
 381
 382   for (i = 0; i < clEnv->number_devices; i++)
 383     clEnv->devices[i]->enabled=MagickFalse;
 384
 385   for (i = 0; i < clEnv->number_devices; i++)
 386   {
 387     device=clEnv->devices[i];
 388     if (device->type != type)
 389       continue;
 390
 391     device->enabled=MagickTrue;
 392     for (j = i+1; j < clEnv->number_devices; j++)
 393     {
 394       MagickCLDevice
 395         other_device;
 396
 397       other_device=clEnv->devices[j];
 398       if (IsSameOpenCLDevice(device,other_device))
 399         other_device->enabled=MagickTrue;
 400     }
 401   }
 402 }
 403
 404 static size_t StringSignature(const char* string)
 405 {
 406   size_t
 407     n,
 408     i,
 409     j,
 410     signature,
 411     stringLength;
 412
 413   union
 414   {
 415     const char* s;
 416     const size_t* u;
 417   } p;
 418
 419   stringLength=(size_t) strlen(string);
 420   signature=stringLength;
 421   n=stringLength/sizeof(size_t);
 422   p.s=string;
 423   for (i = 0; i < n; i++)
 424     signature^=p.u[i];
 425   if (n * sizeof(size_t) != stringLength)
 426     {
 427       char
 428         padded[4];
 429
 430       j=n*sizeof(size_t);
 431       for (i = 0; i < 4; i++, j++)
 432       {
 433         if (j < stringLength)
 434           padded[i]=p.s[j];
 435         else
 436           padded[i]=0;
 437       }
 438       p.s=padded;
 439       signature^=p.u[0];
 440     }
 441   return(signature);
 442 }
 443
 444 /*
 445   Provide call to OpenCL library methods
 446 */
 447
 448 MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
 449   cl_mem_flags flags,size_t size,void *host_ptr)
 450 {
 451   return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr,
 452     (cl_int *) NULL));
 453 }
 454
 455 MagickPrivate void ReleaseOpenCLKernel(cl_kernel kernel)
 456 {
 457   (void) openCL_library->clReleaseKernel(kernel);
 458 }
 459
 460 MagickPrivate void ReleaseOpenCLMemObject(cl_mem memobj)
 461 {
 462   (void) openCL_library->clReleaseMemObject(memobj);
 463 }
 464
 465 MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,size_t arg_index,
 466   size_t arg_size,const void *arg_value)
 467 {
 468   return(openCL_library->clSetKernelArg(kernel,(cl_uint) arg_index,arg_size,
 469     arg_value));
 470 }
 471
 472 /*
 473 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 474 %                                                                             %
 475 %                                                                             %
 476 %                                                                             %
 477 +   A c q u i r e M a g i c k C L C a c h e I n f o                           %
 478 %                                                                             %
 479 %                                                                             %
 480 %                                                                             %
 481 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 482 %
 483 %  AcquireMagickCLCacheInfo() acquires an OpenCL cache info structure.
 484 %
 485 %  The format of the AcquireMagickCLCacheInfo method is:
 486 %
 487 %      MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
 488 %        Quantum *pixels,const MagickSizeType length)
 489 %
 490 %  A description of each parameter follows:
 491 %
 492 %    o device: the OpenCL device.
 493 %
 494 %    o pixels: the pixel buffer of the image.
 495 %
 496 %    o length: the length of the pixel buffer.
 497 %
 498 */
 499
 500 MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
 501   Quantum *pixels,const MagickSizeType length)
 502 {
 503   cl_int
 504     status;
 505
 506   MagickCLCacheInfo
 507     info;
 508
 509   info=(MagickCLCacheInfo) AcquireMagickMemory(sizeof(*info));
 510   if (info == (MagickCLCacheInfo) NULL)
 511     ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
 512   (void) ResetMagickMemory(info,0,sizeof(*info));
 513   LockSemaphoreInfo(openCL_lock);
 514   device->requested++;
 515   UnlockSemaphoreInfo(openCL_lock);
 516   info->device=device;
 517   info->length=length;
 518   info->pixels=pixels;
 519   info->buffer=openCL_library->clCreateBuffer(device->context,
 520     CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(size_t) length,(void *) pixels,
 521     &status);
 522   if (status == CL_SUCCESS)
 523     return(info);
 524   LockSemaphoreInfo(openCL_lock);
 525   device->requested--;
 526   UnlockSemaphoreInfo(openCL_lock);
 527   return((MagickCLCacheInfo) RelinquishMagickMemory(info));
 528 }
 529
 530 /*
 531 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 532 %                                                                             %
 533 %                                                                             %
 534 %                                                                             %
 535 %   A c q u i r e M a g i c k C L D e v i c e                                 %
 536 %                                                                             %
 537 %                                                                             %
 538 %                                                                             %
 539 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 540 %
 541 %  AcquireMagickCLDevice() acquires an OpenCL device
 542 %
 543 %  The format of the AcquireMagickCLDevice method is:
 544 %
 545 %      MagickCLDevice AcquireMagickCLDevice()
 546 %
 547 */
 548
 549 static MagickCLDevice AcquireMagickCLDevice()
 550 {
 551   MagickCLDevice
 552     device;
 553
 554   device=(MagickCLDevice) AcquireMagickMemory(sizeof(*device));
 555   if (device != NULL)
 556   {
 557     (void) ResetMagickMemory(device,0,sizeof(*device));
 558     ActivateSemaphoreInfo(&device->lock);
 559     device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
 560     device->command_queues_index=-1;
 561     device->enabled=MagickTrue;
 562   }
 563   return(device);
 564 }
 565
 566 /*
 567 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 568 %                                                                             %
 569 %                                                                             %
 570 %                                                                             %
 571 %   A c q u i r e M a g i c k C L E n v                                       %
 572 %                                                                             %
 573 %                                                                             %
 574 %                                                                             %
 575 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 576 %
 577 % AcquireMagickCLEnv() allocates the MagickCLEnv structure
 578 %
 579 */
 580
 581 static MagickCLEnv AcquireMagickCLEnv(void)
 582 {
 583   const char
 584     *option;
 585
 586   MagickCLEnv
 587     clEnv;
 588
 589   clEnv=(MagickCLEnv) AcquireMagickMemory(sizeof(*clEnv));
 590   if (clEnv != (MagickCLEnv) NULL)
 591   {
 592     (void) ResetMagickMemory(clEnv,0,sizeof(*clEnv));
 593     ActivateSemaphoreInfo(&clEnv->lock);
 594     clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
 595     clEnv->enabled=MagickTrue;
 596     option=getenv("MAGICK_OCL_DEVICE");
 597     if ((option != (const char *) NULL) && (strcmp(option,"OFF") == 0))
 598       clEnv->enabled=MagickFalse;
 599   }
 600   return clEnv;
 601 }
 602
 603 /*
 604 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 605 %                                                                             %
 606 %                                                                             %
 607 %                                                                             %
 608 +   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                         %
 609 %                                                                             %
 610 %                                                                             %
 611 %                                                                             %
 612 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 613 %
 614 %  AcquireOpenCLCommandQueue() acquires an OpenCL command queue
 615 %
 616 %  The format of the AcquireOpenCLCommandQueue method is:
 617 %
 618 %      cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
 619 %
 620 %  A description of each parameter follows:
 621 %
 622 %    o device: the OpenCL device.
 623 %
 624 */
 625
 626 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
 627 {
 628   cl_command_queue
 629     queue;
 630
 631   cl_command_queue_properties
 632     properties;
 633
 634   assert(device != (MagickCLDevice) NULL);
 635   LockSemaphoreInfo(device->lock);
 636   if ((device->profile_kernels == MagickFalse) &&
 637       (device->command_queues_index >= 0))
 638   {
 639     queue=device->command_queues[device->command_queues_index--];
 640     UnlockSemaphoreInfo(device->lock);
 641   }
 642   else
 643   {
 644     UnlockSemaphoreInfo(device->lock);
 645     properties=(cl_command_queue_properties) NULL;
 646     if (device->profile_kernels != MagickFalse)
 647       properties=CL_QUEUE_PROFILING_ENABLE;
 648     queue=openCL_library->clCreateCommandQueue(device->context,
 649       device->deviceID,properties,(cl_int *) NULL);
 650   }
 651   return(queue);
 652 }
 653
 654 /*
 655 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 656 %                                                                             %
 657 %                                                                             %
 658 %                                                                             %
 659 +   A c q u i r e O p e n C L K e r n e l                                     %
 660 %                                                                             %
 661 %                                                                             %
 662 %                                                                             %
 663 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 664 %
 665 %  AcquireOpenCLKernel() acquires an OpenCL kernel
 666 %
 667 %  The format of the AcquireOpenCLKernel method is:
 668 %
 669 %      cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
 670 %        MagickOpenCLProgram program, const char* kernelName)
 671 %
 672 %  A description of each parameter follows:
 673 %
 674 %    o clEnv: the OpenCL environment.
 675 %
 676 %    o program: the OpenCL program module that the kernel belongs to.
 677 %
 678 %    o kernelName:  the name of the kernel
 679 %
 680 */
 681
 682 MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
 683   const char *kernel_name)
 684 {
 685   cl_kernel
 686     kernel;
 687
 688   assert(device != (MagickCLDevice) NULL);
 689   kernel=openCL_library->clCreateKernel(device->program,kernel_name,
 690     (cl_int *) NULL);
 691   return(kernel);
 692 }
 693
 694 /*
 695 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 696 %                                                                             %
 697 %                                                                             %
 698 %                                                                             %
 699 %   A u t o S e l e c t O p e n C L D e v i c e s                             %
 700 %                                                                             %
 701 %                                                                             %
 702 %                                                                             %
 703 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 704 %
 705 %  AutoSelectOpenCLDevices() determines the best device based on the
 706 %  information from the micro-benchmark.
 707 %
 708 %  The format of the AutoSelectOpenCLDevices method is:
 709 %
 710 %      void AcquireOpenCLKernel(MagickCLEnv clEnv,ExceptionInfo *exception)
 711 %
 712 %  A description of each parameter follows:
 713 %
 714 %    o clEnv: the OpenCL environment.
 715 %
 716 %    o exception: return any errors or warnings in this structure.
 717 %
 718 */
 719
 720 static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,const char *xml)
 721 {
 722   char
 723     keyword[MagickPathExtent],
 724     *token;
 725
 726   const char
 727     *q;
 728
 729   MagickCLDeviceBenchmark
 730     *device_benchmark;
 731
 732   size_t
 733     i,
 734     extent;
 735
 736   if (xml == (char *) NULL)
 737     return;
 738   device_benchmark=(MagickCLDeviceBenchmark *) NULL;
 739   token=AcquireString(xml);
 740   extent=strlen(token)+MagickPathExtent;
 741   for (q=(char *) xml; *q != '\0'; )
 742   {
 743     /*
 744       Interpret XML.
 745     */
 746     GetNextToken(q,&q,extent,token);
 747     if (*token == '\0')
 748       break;
 749     (void) CopyMagickString(keyword,token,MagickPathExtent);
 750     if (LocaleNCompare(keyword,"<!DOCTYPE",9) == 0)
 751       {
 752         /*
 753           Doctype element.
 754         */
 755         while ((LocaleNCompare(q,"]>",2) != 0) && (*q != '\0'))
 756           GetNextToken(q,&q,extent,token);
 757         continue;
 758       }
 759     if (LocaleNCompare(keyword,"<!--",4) == 0)
 760       {
 761         /*
 762           Comment element.
 763         */
 764         while ((LocaleNCompare(q,"->",2) != 0) && (*q != '\0'))
 765           GetNextToken(q,&q,extent,token);
 766         continue;
 767       }
 768     if (LocaleCompare(keyword,"<device") == 0)
 769       {
 770         /*
 771           Device element.
 772         */
 773         device_benchmark=(MagickCLDeviceBenchmark *) AcquireMagickMemory(
 774           sizeof(*device_benchmark));
 775         if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
 776           break;
 777         (void) ResetMagickMemory(device_benchmark,0,sizeof(*device_benchmark));
 778         device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
 779         continue;
 780       }
 781     if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
 782       continue;
 783     if (LocaleCompare(keyword,"/>") == 0)
 784       {
 785         if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
 786           {
 787             if (LocaleCompare(device_benchmark->name, "CPU") == 0)
 788               clEnv->cpu_score=device_benchmark->score;
 789             else
 790               {
 791                 MagickCLDevice
 792                   device;
 793
 794                 /*
 795                   Set the score for all devices that match this device.
 796                 */
 797                 for (i = 0; i < clEnv->number_devices; i++)
 798                 {
 799                   device=clEnv->devices[i];
 800                   if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
 801                     device->score=device_benchmark->score;
 802                 }
 803               }
 804           }
 805
 806         device_benchmark->platform_name=RelinquishMagickMemory(
 807           device_benchmark->platform_name);
 808         device_benchmark->vendor_name=RelinquishMagickMemory(
 809           device_benchmark->vendor_name);
 810         device_benchmark->name=RelinquishMagickMemory(device_benchmark->name);
 811         device_benchmark->version=RelinquishMagickMemory(
 812           device_benchmark->version);
 813         device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
 814           device_benchmark);
 815         continue;
 816       }
 817     GetNextToken(q,(const char **) NULL,extent,token);
 818     if (*token != '=')
 819       continue;
 820     GetNextToken(q,&q,extent,token);
 821     GetNextToken(q,&q,extent,token);
 822     switch (*keyword)
 823     {
 824       case 'M':
 825       case 'm':
 826       {
 827         if (LocaleCompare((char *) keyword,"maxClockFrequency") == 0)
 828           {
 829             device_benchmark->max_clock_frequency=StringToInteger(token);
 830             break;
 831           }
 832         if (LocaleCompare((char *) keyword,"maxComputeUnits") == 0)
 833           {
 834             device_benchmark->max_compute_units=StringToInteger(token);
 835             break;
 836           }
 837         break;
 838       }
 839       case 'N':
 840       case 'n':
 841       {
 842         if (LocaleCompare((char *) keyword,"name") == 0)
 843           device_benchmark->name=ConstantString(token);
 844         break;
 845       }
 846       case 'P':
 847       case 'p':
 848       {
 849         if (LocaleCompare((char *) keyword,"platform") == 0)
 850           device_benchmark->platform_name=ConstantString(token);
 851         break;
 852       }
 853       case 'S':
 854       case 's':
 855       {
 856         if (LocaleCompare((char *) keyword,"score") == 0)
 857           device_benchmark->score=StringToDouble(token,(char **) NULL);
 858         break;
 859       }
 860       case 'V':
 861       case 'v':
 862       {
 863         if (LocaleCompare((char *) keyword,"vendor") == 0)
 864           device_benchmark->vendor_name=ConstantString(token);
 865         if (LocaleCompare((char *) keyword,"version") == 0)
 866           device_benchmark->version=ConstantString(token);
 867         break;
 868       }
 869       default:
 870         break;
 871     }
 872   }
 873   token=(char *) RelinquishMagickMemory(token);
 874   device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
 875     device_benchmark);
 876 }
 877
 878 static MagickBooleanType CanWriteProfileToFile(const char *filename)
 879 {
 880   FILE
 881     *profileFile;
 882
 883   profileFile=fopen(filename,"ab");
 884
 885   if (profileFile == (FILE *)NULL)
 886     return(MagickFalse);
 887
 888   fclose(profileFile);
 889   return(MagickTrue);
 890 }
 891
 892 static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
 893 {
 894   char
 895     filename[MagickPathExtent];
 896
 897   StringInfo
 898     *option;
 899
 900   size_t
 901     i;
 902
 903   (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
 904     GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
 905
 906   /*
 907     We don't run the benchmark when we can not write out a device profile. The
 908     first GPU device will be used.
 909   */
 910 #if !defined(MAGICKCORE_ZERO_CONFIGURATION_SUPPORT)
 911   if (CanWriteProfileToFile(filename) == MagickFalse)
 912 #endif
 913     {
 914       for (i = 0; i < clEnv->number_devices; i++)
 915         clEnv->devices[i]->score=1.0;
 916
 917       SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
 918       return(MagickFalse);
 919     }
 920
 921   option=ConfigureFileToStringInfo(filename);
 922   LoadOpenCLDeviceBenchmark(clEnv,(const char *) GetStringInfoDatum(option));
 923   option=DestroyStringInfo(option);
 924   return(MagickTrue);
 925 }
 926
 927 static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
 928 {
 929   const char
 930     *option;
 931
 932   double
 933     best_score;
 934
 935   MagickBooleanType
 936     benchmark;
 937
 938   size_t
 939     i;
 940
 941   option=getenv("MAGICK_OCL_DEVICE");
 942   if (option != (const char *) NULL)
 943     {
 944       if (strcmp(option,"GPU") == 0)
 945         SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
 946       else if (strcmp(option,"CPU") == 0)
 947         SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
 948       else if (strcmp(option,"OFF") == 0)
 949         {
 950           for (i = 0; i < clEnv->number_devices; i++)
 951             clEnv->devices[i]->enabled=MagickFalse;
 952           clEnv->enabled=MagickFalse;
 953         }
 954     }
 955
 956   if (LoadOpenCLBenchmarks(clEnv) == MagickFalse)
 957     return;
 958
 959   benchmark=MagickFalse;
 960   if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
 961     benchmark=MagickTrue;
 962   else
 963     {
 964       for (i = 0; i < clEnv->number_devices; i++)
 965       {
 966         if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
 967         {
 968           benchmark=MagickTrue;
 969           break;
 970         }
 971       }
 972     }
 973
 974   if (benchmark != MagickFalse)
 975     BenchmarkOpenCLDevices(clEnv);
 976
 977   best_score=clEnv->cpu_score;
 978   for (i = 0; i < clEnv->number_devices; i++)
 979     best_score=MagickMin(clEnv->devices[i]->score,best_score);
 980
 981   for (i = 0; i < clEnv->number_devices; i++)
 982   {
 983     if (clEnv->devices[i]->score != best_score)
 984       clEnv->devices[i]->enabled=MagickFalse;
 985   }
 986 }
 987
 988 /*
 989 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 990 %                                                                             %
 991 %                                                                             %
 992 %                                                                             %
 993 %   B e n c h m a r k O p e n C L D e v i c e s                               %
 994 %                                                                             %
 995 %                                                                             %
 996 %                                                                             %
 997 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
 998 %
 999 %  BenchmarkOpenCLDevices() benchmarks the OpenCL devices and the CPU to help
1000 %  the automatic selection of the best device.
1001 %
1002 %  The format of the BenchmarkOpenCLDevices method is:
1003 %
1004 %    void BenchmarkOpenCLDevices(MagickCLEnv clEnv,ExceptionInfo *exception)
1005 %
1006 %  A description of each parameter follows:
1007 %
1008 %    o clEnv: the OpenCL environment.
1009 %
1010 %    o exception: return any errors or warnings
1011 */
1012
1013 static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
1014 {
1015   AccelerateTimer
1016     timer;
1017
1018   ExceptionInfo
1019     *exception;
1020
1021   Image
1022     *inputImage;
1023
1024   ImageInfo
1025     *imageInfo;
1026
1027   size_t
1028     i;
1029
1030   exception=AcquireExceptionInfo();
1031   imageInfo=AcquireImageInfo();
1032   CloneString(&imageInfo->size,"2048x1536");
1033   CopyMagickString(imageInfo->filename,"xc:none",MagickPathExtent);
1034   inputImage=ReadImage(imageInfo,exception);
1035
1036   InitAccelerateTimer(&timer);
1037
1038   for (i=0; i<=2; i++)
1039   {
1040     Image
1041       *bluredImage,
1042       *resizedImage,
1043       *unsharpedImage;
1044
1045     if (i > 0)
1046       StartAccelerateTimer(&timer);
1047
1048     bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
1049     unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
1050       exception);
1051     resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
1052       exception);
1053
1054     /*
1055       We need this to get a proper performance benchmark, the operations
1056       are executed asynchronous.
1057     */
1058     if (is_cpu == MagickFalse)
1059       {
1060         CacheInfo
1061           *cache_info;
1062
1063         cache_info=(CacheInfo *) resizedImage->cache;
1064         if (cache_info->opencl != (MagickCLCacheInfo) NULL)
1065           openCL_library->clWaitForEvents(cache_info->opencl->event_count,
1066             cache_info->opencl->events);
1067       }
1068
1069     if (i > 0)
1070       StopAccelerateTimer(&timer);
1071
1072     if (bluredImage != (Image *) NULL)
1073       DestroyImage(bluredImage);
1074     if (unsharpedImage != (Image *) NULL)
1075       DestroyImage(unsharpedImage);
1076     if (resizedImage != (Image *) NULL)
1077       DestroyImage(resizedImage);
1078   }
1079   DestroyImage(inputImage);
1080   return(ReadAccelerateTimer(&timer));
1081 }
1082
1083 static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
1084   MagickCLDevice device)
1085 {
1086   testEnv->devices[0]=device;
1087   default_CLEnv=testEnv;
1088   device->score=RunOpenCLBenchmark(MagickFalse);
1089   default_CLEnv=clEnv;
1090   testEnv->devices[0]=(MagickCLDevice) NULL;
1091 }
1092
1093 static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
1094 {
1095   char
1096     filename[MagickPathExtent];
1097
1098   FILE
1099     *cache_file;
1100
1101   MagickCLDevice
1102     device;
1103
1104   size_t
1105     i,
1106     j;
1107
1108   (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1109     GetOpenCLCacheDirectory(),DirectorySeparator,
1110     IMAGEMAGICK_PROFILE_FILE);
1111
1112   cache_file=fopen_utf8(filename,"wb");
1113   if (cache_file == (FILE *) NULL)
1114     return;
1115   fwrite("<devices>\n",sizeof(char),10,cache_file);
1116   fprintf(cache_file,"  <device name=\"CPU\" score=\"%.4g\"/>\n",
1117     clEnv->cpu_score);
1118   for (i = 0; i < clEnv->number_devices; i++)
1119   {
1120     MagickBooleanType
1121       duplicate;
1122
1123     device=clEnv->devices[i];
1124     duplicate=MagickFalse;
1125     for (j = 0; j < i; j++)
1126     {
1127       if (IsSameOpenCLDevice(clEnv->devices[j],device))
1128       {
1129         duplicate=MagickTrue;
1130         break;
1131       }
1132     }
1133
1134     if (duplicate)
1135       continue;
1136
1137     if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1138       fprintf(cache_file,"  <device platform=\"%s\" vendor=\"%s\" name=\"%s\"\
1139  version=\"%s\" maxClockFrequency=\"%d\" maxComputeUnits=\"%d\"\
1140  score=\"%.4g\"/>\n",
1141         device->platform_name,device->vendor_name,device->name,device->version,
1142         (int)device->max_clock_frequency,(int)device->max_compute_units,
1143         device->score);
1144   }
1145   fwrite("</devices>",sizeof(char),10,cache_file);
1146
1147   fclose(cache_file);
1148 }
1149
1150 static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
1151 {
1152   MagickCLDevice
1153     device;
1154
1155   MagickCLEnv
1156     testEnv;
1157
1158   size_t
1159     i,
1160     j;
1161
1162   testEnv=AcquireMagickCLEnv();
1163   testEnv->library=openCL_library;
1164   testEnv->devices=(MagickCLDevice *) AcquireMagickMemory(
1165     sizeof(MagickCLDevice));
1166   testEnv->number_devices=1;
1167   testEnv->benchmark_thread_id=GetMagickThreadId();
1168   testEnv->initialized=MagickTrue;
1169
1170   for (i = 0; i < clEnv->number_devices; i++)
1171     clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
1172
1173   for (i = 0; i < clEnv->number_devices; i++)
1174   {
1175     device=clEnv->devices[i];
1176     if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1177       RunDeviceBenckmark(clEnv,testEnv,device);
1178
1179     /* Set the score on all the other devices that are the same */
1180     for (j = i+1; j < clEnv->number_devices; j++)
1181     {
1182       MagickCLDevice
1183         other_device;
1184
1185       other_device=clEnv->devices[j];
1186       if (IsSameOpenCLDevice(device,other_device))
1187         other_device->score=device->score;
1188     }
1189   }
1190
1191   testEnv->enabled=MagickFalse;
1192   default_CLEnv=testEnv;
1193   clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
1194   default_CLEnv=clEnv;
1195
1196   testEnv=RelinquishMagickCLEnv(testEnv);
1197   CacheOpenCLBenchmarks(clEnv);
1198 }
1199
1200 /*
1201 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1202 %                                                                             %
1203 %                                                                             %
1204 %                                                                             %
1205 %   C o m p i l e O p e n C L K e r n e l                                     %
1206 %                                                                             %
1207 %                                                                             %
1208 %                                                                             %
1209 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1210 %
1211 %  CompileOpenCLKernel() compiles the kernel for the specified device. The
1212 %  kernel will be cached on disk to reduce the compilation time.
1213 %
1214 %  The format of the CompileOpenCLKernel method is:
1215 %
1216 %      MagickBooleanType AcquireOpenCLKernel(MagickCLDevice clEnv,
1217 %        unsigned int signature,const char *kernel,const char *options,
1218 %        ExceptionInfo *exception)
1219 %
1220 %  A description of each parameter follows:
1221 %
1222 %    o device: the OpenCL device.
1223 %
1224 %    o kernel: the source code of the kernel.
1225 %
1226 %    o options: options for the compiler.
1227 %
1228 %    o signature: a number to uniquely identify the kernel
1229 %
1230 %    o exception: return any errors or warnings in this structure.
1231 %
1232 */
1233
1234 static void CacheOpenCLKernel(MagickCLDevice device,char *filename,
1235   ExceptionInfo *exception)
1236 {
1237   cl_uint
1238     status;
1239
1240   size_t
1241     binaryProgramSize;
1242
1243   unsigned char
1244     *binaryProgram;
1245
1246   status=openCL_library->clGetProgramInfo(device->program,
1247     CL_PROGRAM_BINARY_SIZES,sizeof(size_t),&binaryProgramSize,NULL);
1248   if (status != CL_SUCCESS)
1249     return;
1250
1251   binaryProgram=(unsigned char*) AcquireMagickMemory(binaryProgramSize);
1252   status=openCL_library->clGetProgramInfo(device->program,
1253     CL_PROGRAM_BINARIES,sizeof(unsigned char*),&binaryProgram,NULL);
1254   if (status == CL_SUCCESS)
1255     (void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
1256   binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1257 }
1258
1259 static MagickBooleanType LoadCachedOpenCLKernel(MagickCLDevice device,
1260   const char *filename)
1261 {
1262   cl_int
1263     binaryStatus,
1264     status;
1265
1266   ExceptionInfo
1267     *exception;
1268
1269   size_t
1270     length;
1271
1272   unsigned char
1273     *binaryProgram;
1274
1275   exception=AcquireExceptionInfo();
1276   binaryProgram=(unsigned char *) FileToBlob(filename,~0UL,&length,exception);
1277   exception=DestroyExceptionInfo(exception);
1278   if (binaryProgram == (unsigned char *) NULL)
1279     return(MagickFalse);
1280   device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
1281     &device->deviceID,&length,(const unsigned char**)&binaryProgram,
1282     &binaryStatus,&status);
1283   binaryProgram=(unsigned char *) RelinquishMagickMemory(binaryProgram);
1284   return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
1285     MagickTrue);
1286 }
1287
1288 static void LogOpenCLBuildFailure(MagickCLDevice device,const char *kernel,
1289   ExceptionInfo *exception)
1290 {
1291   char
1292     filename[MagickPathExtent],
1293     *log;
1294
1295   size_t
1296     log_size;
1297
1298   (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1299     GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.cl");
1300
1301   (void) remove_utf8(filename);
1302   (void) BlobToFile(filename,kernel,strlen(kernel),exception);
1303
1304   openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1305     CL_PROGRAM_BUILD_LOG,0,NULL,&log_size);
1306   log=(char*)AcquireMagickMemory(log_size);
1307   openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1308     CL_PROGRAM_BUILD_LOG,log_size,log,&log_size);
1309
1310   (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1311     GetOpenCLCacheDirectory(),DirectorySeparator,"magick_badcl.log");
1312
1313   (void) remove_utf8(filename);
1314   (void) BlobToFile(filename,log,log_size,exception);
1315   log=(char*)RelinquishMagickMemory(log);
1316 }
1317
1318 static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
1319   const char *kernel,const char *options,size_t signature,
1320   ExceptionInfo *exception)
1321 {
1322   char
1323     deviceName[MagickPathExtent],
1324     filename[MagickPathExtent],
1325     *ptr;
1326
1327   cl_int
1328     status;
1329
1330   MagickBooleanType
1331     loaded;
1332
1333   size_t
1334     length;
1335
1336   (void) CopyMagickString(deviceName,device->name,MagickPathExtent);
1337   ptr=deviceName;
1338   /* Strip out illegal characters for file names */
1339   while (*ptr != '\0')
1340   {
1341     if ((*ptr == ' ') || (*ptr == '\\') || (*ptr == '/') || (*ptr == ':') ||
1342         (*ptr == '*') || (*ptr == '?') || (*ptr == '"') || (*ptr == '<') ||
1343         (*ptr == '>' || *ptr == '|'))
1344       *ptr = '_';
1345     ptr++;
1346   }
1347   (void) FormatLocaleString(filename,MagickPathExtent,
1348     "%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
1349     DirectorySeparator,"magick_opencl",deviceName,signature,
1350     (double) sizeof(char*)*8);
1351   loaded=LoadCachedOpenCLKernel(device,filename);
1352   if (loaded == MagickFalse)
1353     {
1354       /* Binary CL program unavailable, compile the program from source */
1355       length=strlen(kernel);
1356       device->program=openCL_library->clCreateProgramWithSource(
1357         device->context,1,&kernel,&length,&status);
1358       if (status != CL_SUCCESS)
1359         return(MagickFalse);
1360     }
1361
1362   status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
1363     options,NULL,NULL);
1364   if (status != CL_SUCCESS)
1365   {
1366     (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
1367       "clBuildProgram failed.","(%d)",(int)status);
1368     LogOpenCLBuildFailure(device,kernel,exception);
1369     return(MagickFalse);
1370   }
1371
1372   /* Save the binary to a file to avoid re-compilation of the kernels */
1373   if (loaded == MagickFalse)
1374     CacheOpenCLKernel(device,filename,exception);
1375
1376   return(MagickTrue);
1377 }
1378
1379 /*
1380 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1381 %                                                                             %
1382 %                                                                             %
1383 %                                                                             %
1384 +   C o p y M a g i c k C L C a c h e I n f o                                 %
1385 %                                                                             %
1386 %                                                                             %
1387 %                                                                             %
1388 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1389 %
1390 %  CopyMagickCLCacheInfo() copies the memory from the device into host memory.
1391 %
1392 %  The format of the CopyMagickCLCacheInfo method is:
1393 %
1394 %      void CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1395 %
1396 %  A description of each parameter follows:
1397 %
1398 %    o info: the OpenCL cache info.
1399 %
1400 */
1401 MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1402 {
1403   cl_command_queue
1404     queue;
1405
1406   Quantum
1407     *pixels;
1408
1409   if (info == (MagickCLCacheInfo) NULL)
1410     return((MagickCLCacheInfo) NULL);
1411   if (info->event_count > 0)
1412     {
1413       queue=AcquireOpenCLCommandQueue(info->device);
1414       pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
1415         CL_MAP_READ | CL_MAP_WRITE,0,info->length,info->event_count,
1416         info->events,(cl_event *) NULL,(cl_int *) NULL);
1417       assert(pixels == info->pixels);
1418       ReleaseOpenCLCommandQueue(info->device,queue);
1419     }
1420   return(RelinquishMagickCLCacheInfo(info,MagickFalse));
1421 }
1422
1423 /*
1424 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1425 %                                                                             %
1426 %                                                                             %
1427 %                                                                             %
1428 +   D u m p O p e n C L P r o f i l e D a t a                                 %
1429 %                                                                             %
1430 %                                                                             %
1431 %                                                                             %
1432 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1433 %
1434 %  DumpOpenCLProfileData() dumps the kernel profile data.
1435 %
1436 %  The format of the DumpProfileData method is:
1437 %
1438 %      void DumpProfileData()
1439 %
1440 */
1441
1442 MagickPrivate void DumpOpenCLProfileData()
1443 {
1444 #define OpenCLLog(message) \
1445    fwrite(message,sizeof(char),strlen(message),log); \
1446    fwrite("\n",sizeof(char),1,log);
1447
1448   char
1449     buf[4096],
1450     filename[MagickPathExtent],
1451     indent[160];
1452
1453   FILE
1454     *log;
1455
1456   MagickCLEnv
1457     clEnv;
1458
1459   size_t
1460     i,
1461     j;
1462
1463   clEnv=GetCurrentOpenCLEnv();
1464   if (clEnv == (MagickCLEnv) NULL)
1465     return;
1466
1467   for (i = 0; i < clEnv->number_devices; i++)
1468     if (clEnv->devices[i]->profile_kernels != MagickFalse)
1469       break;
1470   if (i == clEnv->number_devices)
1471     return;
1472
1473   (void) FormatLocaleString(filename,MagickPathExtent,"%s%s%s",
1474     GetOpenCLCacheDirectory(),DirectorySeparator,"ImageMagickOpenCL.log");
1475
1476   log=fopen_utf8(filename,"wb");
1477
1478   for (i = 0; i < clEnv->number_devices; i++)
1479   {
1480     MagickCLDevice
1481       device;
1482
1483     device=clEnv->devices[i];
1484     if ((device->profile_kernels == MagickFalse) ||
1485         (device->profile_records == (KernelProfileRecord *) NULL))
1486       continue;
1487
1488     OpenCLLog("====================================================");
1489     fprintf(log,"Device:  %s\n",device->name);
1490     fprintf(log,"Version: %s\n",device->version);
1491     OpenCLLog("====================================================");
1492     OpenCLLog("                     average   calls     min     max");
1493     OpenCLLog("                     -------   -----     ---     ---");
1494     j=0;
1495     while (device->profile_records[j] != (KernelProfileRecord) NULL)
1496     {
1497       KernelProfileRecord
1498         profile;
1499
1500       profile=device->profile_records[j];
1501       strcpy(indent,"                    ");
1502       strncpy(indent,profile->kernel_name,MagickMin(strlen(
1503         profile->kernel_name),strlen(indent)-1));
1504       sprintf(buf,"%s %7d %7d %7d %7d",indent,(int) (profile->total/
1505         profile->count),(int) profile->count,(int) profile->min,
1506         (int) profile->max);
1507       OpenCLLog(buf);
1508       j++;
1509     }
1510     OpenCLLog("====================================================");
1511     fwrite("\n\n",sizeof(char),2,log);
1512   }
1513   fclose(log);
1514 }
1515 /*
1516 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1517 %                                                                             %
1518 %                                                                             %
1519 %                                                                             %
1520 +   E n q u e u e O p e n C L K e r n e l                                     %
1521 %                                                                             %
1522 %                                                                             %
1523 %                                                                             %
1524 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1525 %
1526 %  EnqueueOpenCLKernel() enques the specified kernel and registers the OpenCL
1527 %  events with the images.
1528 %
1529 %  The format of the EnqueueOpenCLKernel method is:
1530 %
1531 %      MagickBooleanType EnqueueOpenCLKernel(cl_kernel kernel,cl_uint work_dim,
1532 %        const size_t *global_work_offset,const size_t *global_work_size,
1533 %        const size_t *local_work_size,const Image *input_image,
1534 %        const Image *output_image,ExceptionInfo *exception)
1535 %
1536 %  A description of each parameter follows:
1537 %
1538 %    o kernel: the OpenCL kernel.
1539 %
1540 %    o work_dim: the number of dimensions used to specify the global work-items
1541 %                and work-items in the work-group.
1542 %
1543 %    o offset: can be used to specify an array of work_dim unsigned values
1544 %              that describe the offset used to calculate the global ID of a
1545 %              work-item.
1546 %
1547 %    o gsize: points to an array of work_dim unsigned values that describe the
1548 %             number of global work-items in work_dim dimensions that will
1549 %             execute the kernel function.
1550 %
1551 %    o lsize: points to an array of work_dim unsigned values that describe the
1552 %             number of work-items that make up a work-group that will execute
1553 %             the kernel specified by kernel.
1554 %
1555 %    o input_image: the input image of the operation.
1556 %
1557 %    o output_image: the output or secondairy image of the operation.
1558 %
1559 %    o exception: return any errors or warnings in this structure.
1560 %
1561 */
1562
1563 static void RegisterCacheEvent(MagickCLCacheInfo info,cl_event event)
1564 {
1565   assert(info != (MagickCLCacheInfo) NULL);
1566   assert(event != (cl_event) NULL);
1567   if (info->events == (cl_event *) NULL)
1568     {
1569       info->events=AcquireMagickMemory(sizeof(*info->events));
1570       info->event_count=1;
1571     }
1572   else
1573     info->events=ResizeQuantumMemory(info->events,++info->event_count,
1574       sizeof(*info->events));
1575   if (info->events == (cl_event *) NULL)
1576     ThrowFatalException(ResourceLimitFatalError,"MemoryAllocationFailed");
1577   info->events[info->event_count-1]=event;
1578   openCL_library->clRetainEvent(event);
1579 }
1580
1581 MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1582   cl_kernel kernel,cl_uint work_dim,const size_t *offset,const size_t *gsize,
1583   const size_t *lsize,const Image *input_image,const Image *output_image,
1584   MagickBooleanType flush,ExceptionInfo *exception)
1585 {
1586   CacheInfo
1587     *output_info,
1588     *input_info;
1589
1590   cl_event
1591     event,
1592     *events;
1593
1594   cl_int
1595     status;
1596
1597   cl_uint
1598     event_count;
1599
1600   assert(input_image != (const Image *) NULL);
1601   input_info=(CacheInfo *) input_image->cache;
1602   assert(input_info != (CacheInfo *) NULL);
1603   assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1604   event_count=input_info->opencl->event_count;
1605   events=input_info->opencl->events;
1606   output_info=(CacheInfo *) NULL;
1607   if (output_image != (const Image *) NULL)
1608     {
1609       output_info=(CacheInfo *) output_image->cache;
1610       assert(output_info != (CacheInfo *) NULL);
1611       assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1612       if (output_info->opencl->event_count > 0)
1613         {
1614           ssize_t
1615             i;
1616
1617           event_count+=output_info->opencl->event_count;
1618           events=AcquireQuantumMemory(event_count,sizeof(*events));
1619           if (events == (cl_event *) NULL)
1620             return(MagickFalse);
1621           for (i=0; i < (ssize_t) event_count; i++)
1622           {
1623             if (i < (ssize_t) input_info->opencl->event_count)
1624               events[i]=input_info->opencl->events[i];
1625             else
1626               events[i]=output_info->opencl->events[i-
1627                 input_info->opencl->event_count];
1628           }
1629         }
1630     }
1631   status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1632     gsize,lsize,event_count,events,&event);
1633   /* This can fail due to memory issues and calling clFinish might help. */
1634   if ((status != CL_SUCCESS) && (event_count > 0))
1635     {
1636       openCL_library->clFinish(queue);
1637       status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
1638         offset,gsize,lsize,event_count,events,&event);
1639     }
1640   if ((output_info != (CacheInfo *) NULL) &&
1641       (output_info->opencl->event_count > 0))
1642     events=(cl_event *) RelinquishMagickMemory(events);
1643   if (status != CL_SUCCESS)
1644     {
1645       (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1646         GetMagickModule(),ResourceLimitWarning,
1647         "clEnqueueNDRangeKernel failed.","'%s'",".");
1648       return(MagickFalse);
1649     }
1650   if (flush != MagickFalse)
1651     openCL_library->clFlush(queue);
1652   if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1653     {
1654       RegisterCacheEvent(input_info->opencl,event);
1655       if (output_info != (CacheInfo *) NULL)
1656         RegisterCacheEvent(output_info->opencl,event);
1657     }
1658   openCL_library->clReleaseEvent(event);
1659   return(MagickTrue);
1660 }
1661
1662 /*
1663 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1664 %                                                                             %
1665 %                                                                             %
1666 %                                                                             %
1667 +   G e t C u r r u n t O p e n C L E n v                                     %
1668 %                                                                             %
1669 %                                                                             %
1670 %                                                                             %
1671 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1672 %
1673 %  GetCurrentOpenCLEnv() returns the current OpenCL env
1674 %
1675 %  The format of the GetCurrentOpenCLEnv method is:
1676 %
1677 %      MagickCLEnv GetCurrentOpenCLEnv()
1678 %
1679 */
1680
1681 MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(void)
1682 {
1683   if (default_CLEnv != (MagickCLEnv) NULL)
1684   {
1685     if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1686         (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1687       return((MagickCLEnv) NULL);
1688     else
1689       return(default_CLEnv);
1690   }
1691
1692   if (GetOpenCLCacheDirectory() == (char *) NULL)
1693     return((MagickCLEnv) NULL);
1694
1695   if (openCL_lock == (SemaphoreInfo *) NULL)
1696     ActivateSemaphoreInfo(&openCL_lock);
1697
1698   LockSemaphoreInfo(openCL_lock);
1699   if (default_CLEnv == (MagickCLEnv) NULL)
1700     default_CLEnv=AcquireMagickCLEnv();
1701   UnlockSemaphoreInfo(openCL_lock);
1702
1703   return(default_CLEnv);
1704 }
1705
1706 /*
1707 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1708 %                                                                             %
1709 %                                                                             %
1710 %                                                                             %
1711 %   G e t O p e n C L D e v i c e B e n c h m a r k D u r a t i o n           %
1712 %                                                                             %
1713 %                                                                             %
1714 %                                                                             %
1715 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1716 %
1717 %  GetOpenCLDeviceBenchmarkScore() returns the score of the benchmark for the
1718 %  device. The score is determined by the duration of the micro benchmark so
1719 %  that means a lower score is better than a higher score.
1720 %
1721 %  The format of the GetOpenCLDeviceBenchmarkScore method is:
1722 %
1723 %      double GetOpenCLDeviceBenchmarkScore(const MagickCLDevice device)
1724 %
1725 %  A description of each parameter follows:
1726 %
1727 %    o device: the OpenCL device.
1728 */
1729
1730 MagickExport double GetOpenCLDeviceBenchmarkScore(
1731   const MagickCLDevice device)
1732 {
1733   if (device == (MagickCLDevice) NULL)
1734     return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1735   return(device->score);
1736 }
1737
1738 /*
1739 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1740 %                                                                             %
1741 %                                                                             %
1742 %                                                                             %
1743 %   G e t O p e n C L D e v i c e E n a b l e d                               %
1744 %                                                                             %
1745 %                                                                             %
1746 %                                                                             %
1747 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1748 %
1749 %  GetOpenCLDeviceEnabled() returns true if the device is enabled.
1750 %
1751 %  The format of the GetOpenCLDeviceEnabled method is:
1752 %
1753 %      MagickBooleanType GetOpenCLDeviceEnabled(const MagickCLDevice device)
1754 %
1755 %  A description of each parameter follows:
1756 %
1757 %    o device: the OpenCL device.
1758 */
1759
1760 MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
1761   const MagickCLDevice device)
1762 {
1763   if (device == (MagickCLDevice) NULL)
1764     return(MagickFalse);
1765   return(device->enabled);
1766 }
1767
1768 /*
1769 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1770 %                                                                             %
1771 %                                                                             %
1772 %                                                                             %
1773 %   G e t O p e n C L D e v i c e N a m e                                     %
1774 %                                                                             %
1775 %                                                                             %
1776 %                                                                             %
1777 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1778 %
1779 %  GetOpenCLDeviceName() returns the name of the device.
1780 %
1781 %  The format of the GetOpenCLDeviceName method is:
1782 %
1783 %      const char *GetOpenCLDeviceName(const MagickCLDevice device)
1784 %
1785 %  A description of each parameter follows:
1786 %
1787 %    o device: the OpenCL device.
1788 */
1789
1790 MagickExport const char *GetOpenCLDeviceName(const MagickCLDevice device)
1791 {
1792   if (device == (MagickCLDevice) NULL)
1793     return((const char *) NULL);
1794   return(device->name);
1795 }
1796
1797 /*
1798 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1799 %                                                                             %
1800 %                                                                             %
1801 %                                                                             %
1802 %   G e t O p e n C L D e v i c e V e n d o r N a m e                         %
1803 %                                                                             %
1804 %                                                                             %
1805 %                                                                             %
1806 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1807 %
1808 %  GetOpenCLDeviceVendorName() returns the vendor name of the device.
1809 %
1810 %  The format of the GetOpenCLDeviceVendorName method is:
1811 %
1812 %      const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1813 %
1814 %  A description of each parameter follows:
1815 %
1816 %    o device: the OpenCL device.
1817 */
1818
1819 MagickExport const char *GetOpenCLDeviceVendorName(const MagickCLDevice device)
1820 {
1821   if (device == (MagickCLDevice) NULL)
1822     return((const char *) NULL);
1823   return(device->vendor_name);
1824 }
1825
1826 /*
1827 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1828 %                                                                             %
1829 %                                                                             %
1830 %                                                                             %
1831 %   G e t O p e n C L D e v i c e s                                           %
1832 %                                                                             %
1833 %                                                                             %
1834 %                                                                             %
1835 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1836 %
1837 %  GetOpenCLDevices() returns the devices of the OpenCL environment at sets the
1838 %  value of length to the number of devices that are available.
1839 %
1840 %  The format of the GetOpenCLDevices method is:
1841 %
1842 %      const MagickCLDevice *GetOpenCLDevices(size_t *length,
1843 %        ExceptionInfo *exception)
1844 %
1845 %  A description of each parameter follows:
1846 %
1847 %    o length: the number of device.
1848 %
1849 %    o exception: return any errors or warnings in this structure.
1850 %
1851 */
1852
1853 MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
1854   ExceptionInfo *exception)
1855 {
1856   MagickCLEnv
1857     clEnv;
1858
1859   clEnv=GetCurrentOpenCLEnv();
1860   if (clEnv == (MagickCLEnv) NULL)
1861     {
1862       if (length != (size_t *) NULL)
1863         *length=0;
1864       return((MagickCLDevice *) NULL);
1865     }
1866   InitializeOpenCL(clEnv,exception);
1867   if (length != (size_t *) NULL)
1868     *length=clEnv->number_devices;
1869   return(clEnv->devices);
1870 }
1871
1872 /*
1873 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1874 %                                                                             %
1875 %                                                                             %
1876 %                                                                             %
1877 %   G e t O p e n C L D e v i c e T y p e                                     %
1878 %                                                                             %
1879 %                                                                             %
1880 %                                                                             %
1881 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1882 %
1883 %  GetOpenCLDeviceType() returns the type of the device.
1884 %
1885 %  The format of the GetOpenCLDeviceType method is:
1886 %
1887 %      MagickCLDeviceType GetOpenCLDeviceType(const MagickCLDevice device)
1888 %
1889 %  A description of each parameter follows:
1890 %
1891 %    o device: the OpenCL device.
1892 */
1893
1894 MagickExport MagickCLDeviceType GetOpenCLDeviceType(
1895   const MagickCLDevice device)
1896 {
1897   if (device == (MagickCLDevice) NULL)
1898     return(UndefinedCLDeviceType);
1899   if (device->type == CL_DEVICE_TYPE_GPU)
1900     return(GpuCLDeviceType);
1901   if (device->type == CL_DEVICE_TYPE_CPU)
1902     return(CpuCLDeviceType);
1903   return(UndefinedCLDeviceType);
1904 }
1905
1906 /*
1907 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1908 %                                                                             %
1909 %                                                                             %
1910 %                                                                             %
1911 %   G e t O p e n C L D e v i c e V e r s i o n                               %
1912 %                                                                             %
1913 %                                                                             %
1914 %                                                                             %
1915 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1916 %
1917 %  GetOpenCLDeviceVersion() returns the version of the device.
1918 %
1919 %  The format of the GetOpenCLDeviceName method is:
1920 %
1921 %      const char *GetOpenCLDeviceVersion(MagickCLDevice device)
1922 %
1923 %  A description of each parameter follows:
1924 %
1925 %    o device: the OpenCL device.
1926 */
1927
1928 MagickExport const char *GetOpenCLDeviceVersion(const MagickCLDevice device)
1929 {
1930   if (device == (MagickCLDevice) NULL)
1931     return((const char *) NULL);
1932   return(device->version);
1933 }
1934
1935 /*
1936 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1937 %                                                                             %
1938 %                                                                             %
1939 %                                                                             %
1940 %   G e t O p e n C L E n a b l e d                                           %
1941 %                                                                             %
1942 %                                                                             %
1943 %                                                                             %
1944 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1945 %
1946 %  GetOpenCLEnabled() returns true if OpenCL acceleration is enabled.
1947 %
1948 %  The format of the GetOpenCLEnabled method is:
1949 %
1950 %      MagickBooleanType GetOpenCLEnabled()
1951 %
1952 */
1953
1954 MagickExport MagickBooleanType GetOpenCLEnabled(void)
1955 {
1956   MagickCLEnv
1957     clEnv;
1958
1959   clEnv=GetCurrentOpenCLEnv();
1960   if (clEnv == (MagickCLEnv) NULL)
1961     return(MagickFalse);
1962   return(clEnv->enabled);
1963 }
1964
1965 /*
1966 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1967 %                                                                             %
1968 %                                                                             %
1969 %                                                                             %
1970 %   G e t O p e n C L K e r n e l P r o f i l e R e c o r d s                 %
1971 %                                                                             %
1972 %                                                                             %
1973 %                                                                             %
1974 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1975 %
1976 %  GetOpenCLKernelProfileRecords() returns the profile records for the
1977 %  specified device and sets length to the number of profile records.
1978 %
1979 %  The format of the GetOpenCLKernelProfileRecords method is:
1980 %
1981 %      const KernelProfileRecord *GetOpenCLKernelProfileRecords(size *length)
1982 %
1983 %  A description of each parameter follows:
1984 %
1985 %    o length: the number of profiles records.
1986 */
1987
1988 MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
1989   const MagickCLDevice device,size_t *length)
1990 {
1991   if ((device == (const MagickCLDevice) NULL) || (device->profile_records ==
1992       (KernelProfileRecord *) NULL))
1993   {
1994     if (length != (size_t *) NULL)
1995       *length=0;
1996     return((const KernelProfileRecord *) NULL);
1997   }
1998   if (length != (size_t *) NULL)
1999     {
2000       *length=0;
2001       LockSemaphoreInfo(device->lock);
2002       while (device->profile_records[*length] != (KernelProfileRecord) NULL)
2003         *length=*length+1;
2004       UnlockSemaphoreInfo(device->lock);
2005     }
2006   return(device->profile_records);
2007 }
2008
2009 /*
2010 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2011 %                                                                             %
2012 %                                                                             %
2013 %                                                                             %
2014 %   H a s O p e n C L D e v i c e s                                           %
2015 %                                                                             %
2016 %                                                                             %
2017 %                                                                             %
2018 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2019 %
2020 %  HasOpenCLDevices() checks if the OpenCL environment has devices that are
2021 %  enabled and compiles the kernel for the device when necessary. False will be
2022 %  returned if no enabled devices could be found
2023 %
2024 %  The format of the HasOpenCLDevices method is:
2025 %
2026 %    MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2027 %      ExceptionInfo exception)
2028 %
2029 %  A description of each parameter follows:
2030 %
2031 %    o clEnv: the OpenCL environment.
2032 %
2033 %    o exception: return any errors or warnings in this structure.
2034 %
2035 */
2036
2037 static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2038   ExceptionInfo *exception)
2039 {
2040   char
2041     *accelerateKernelsBuffer,
2042     options[MagickPathExtent];
2043
2044   MagickStatusType
2045     status;
2046
2047   size_t
2048     i;
2049
2050   size_t
2051     signature;
2052
2053   /* Check if there are enabled devices */
2054   for (i = 0; i < clEnv->number_devices; i++)
2055   {
2056     if ((clEnv->devices[i]->enabled != MagickFalse))
2057       break;
2058   }
2059   if (i == clEnv->number_devices)
2060     return(MagickFalse);
2061
2062   /* Check if we need to compile a kernel for one of the devices */
2063   status=MagickTrue;
2064   for (i = 0; i < clEnv->number_devices; i++)
2065   {
2066     if ((clEnv->devices[i]->enabled != MagickFalse) &&
2067         (clEnv->devices[i]->program == (cl_program) NULL))
2068     {
2069       status=MagickFalse;
2070       break;
2071     }
2072   }
2073   if (status != MagickFalse)
2074     return(MagickTrue);
2075
2076   /* Get additional options */
2077   (void) FormatLocaleString(options,MaxTextExtent,CLOptions,
2078     (float)QuantumRange,(float)QuantumScale,(float)CLCharQuantumScale,
2079     (float)MagickEpsilon,(float)MagickPI,(unsigned int)MaxMap,
2080     (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2081
2082   signature=StringSignature(options);
2083   accelerateKernelsBuffer=(char*) AcquireMagickMemory(
2084     strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2085   if (accelerateKernelsBuffer == (char*) NULL)
2086     return(MagickFalse);
2087   sprintf(accelerateKernelsBuffer,"%s%s",accelerateKernels,accelerateKernels2);
2088   signature^=StringSignature(accelerateKernelsBuffer);
2089
2090   status=MagickTrue;
2091   for (i = 0; i < clEnv->number_devices; i++)
2092   {
2093     MagickCLDevice
2094       device;
2095
2096     size_t
2097       device_signature;
2098
2099     device=clEnv->devices[i];
2100     if ((device->enabled == MagickFalse) ||
2101         (device->program != (cl_program) NULL))
2102       continue;
2103
2104     LockSemaphoreInfo(device->lock);
2105     if (device->program != (cl_program) NULL)
2106     {
2107       UnlockSemaphoreInfo(device->lock);
2108       continue;
2109     }
2110     device_signature=signature;
2111     device_signature^=StringSignature(device->platform_name);
2112     status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2113       device_signature,exception);
2114     UnlockSemaphoreInfo(device->lock);
2115     if (status == MagickFalse)
2116       break;
2117   }
2118   accelerateKernelsBuffer=RelinquishMagickMemory(accelerateKernelsBuffer);
2119   return(status);
2120 }
2121
2122 /*
2123 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2124 %                                                                             %
2125 %                                                                             %
2126 %                                                                             %
2127 +   I n i t i a l i z e O p e n C L                                           %
2128 %                                                                             %
2129 %                                                                             %
2130 %                                                                             %
2131 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2132 %
2133 %  InitializeOpenCL() is used to initialize the OpenCL environment. This method
2134 %  makes sure the devices are propertly initialized and benchmarked.
2135 %
2136 %  The format of the InitializeOpenCL method is:
2137 %
2138 %    MagickBooleanType InitializeOpenCL(ExceptionInfo exception)
2139 %
2140 %  A description of each parameter follows:
2141 %
2142 %    o exception: return any errors or warnings in this structure.
2143 %
2144 */
2145
2146 static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2147 {
2148   char
2149     version[MagickPathExtent];
2150
2151   cl_uint
2152     num;
2153
2154   if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2155         MagickPathExtent,version,NULL) != CL_SUCCESS)
2156     return(0);
2157   if (strncmp(version, "OpenCL 1.0 ", 11) == 0)
2158     return(0);
2159   if (clEnv->library->clGetDeviceIDs(platform,
2160         CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2161     return(0);
2162   return(num);
2163 }
2164
2165 static void LoadOpenCLDevices(MagickCLEnv clEnv)
2166 {
2167   cl_context_properties
2168     properties[3];
2169
2170   cl_device_id
2171     *devices;
2172
2173   cl_int
2174     status;
2175
2176   cl_platform_id
2177     *platforms;
2178
2179   cl_uint
2180     i,
2181     j,
2182     next,
2183     number_devices,
2184     number_platforms;
2185
2186   size_t
2187     length;
2188
2189   number_platforms=0;
2190   if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2191     return;
2192   if (number_platforms == 0)
2193     return;
2194   platforms=(cl_platform_id *) AcquireMagickMemory(number_platforms*
2195     sizeof(cl_platform_id));
2196   if (platforms == (cl_platform_id *) NULL)
2197     return;
2198   if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2199     {
2200        platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2201        return;
2202     }
2203   for (i = 0; i < number_platforms; i++)
2204   {
2205     number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2206     if (number_devices == 0)
2207       platforms[i]=(cl_platform_id) NULL;
2208     else
2209       clEnv->number_devices+=number_devices;
2210   }
2211   if (clEnv->number_devices == 0)
2212     {
2213       platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2214       return;
2215     }
2216   clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2217     sizeof(MagickCLDevice));
2218   if (clEnv->devices == (MagickCLDevice *) NULL)
2219     {
2220       RelinquishMagickCLDevices(clEnv);
2221       platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2222       return;
2223     }
2224   (void) ResetMagickMemory(clEnv->devices,0,clEnv->number_devices*
2225     sizeof(MagickCLDevice));
2226   devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2227     sizeof(cl_device_id));
2228   if (devices == (cl_device_id *) NULL)
2229     {
2230       platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2231       RelinquishMagickCLDevices(clEnv);
2232       return;
2233     }
2234   clEnv->number_contexts=(size_t) number_platforms;
2235   clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2236     sizeof(cl_context));
2237   if (clEnv->contexts == (cl_context *) NULL)
2238     {
2239       devices=(cl_device_id *) RelinquishMagickMemory(devices);
2240       platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2241       RelinquishMagickCLDevices(clEnv);
2242       return;
2243     }
2244   next=0;
2245   for (i = 0; i < number_platforms; i++)
2246   {
2247     if (platforms[i] == (cl_platform_id) NULL)
2248       continue;
2249
2250     status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2251       CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2252     if (status != CL_SUCCESS)
2253       continue;
2254
2255     properties[0]=CL_CONTEXT_PLATFORM;
2256     properties[1]=(cl_context_properties) platforms[i];
2257     properties[2]=0;
2258     clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2259       devices,NULL,NULL,&status);
2260     if (status != CL_SUCCESS)
2261       continue;
2262
2263     for (j = 0; j < number_devices; j++,next++)
2264     {
2265       MagickCLDevice
2266         device;
2267
2268       device=AcquireMagickCLDevice();
2269       if (device == (MagickCLDevice) NULL)
2270         break;
2271
2272       device->context=clEnv->contexts[i];
2273       device->deviceID=devices[j];
2274
2275       openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,0,NULL,
2276         &length);
2277       device->platform_name=AcquireQuantumMemory(length,
2278         sizeof(*device->platform_name));
2279       openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_NAME,length,
2280         device->platform_name,NULL);
2281
2282       openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,0,NULL,
2283         &length);
2284       device->vendor_name=AcquireQuantumMemory(length,
2285         sizeof(*device->vendor_name));
2286       openCL_library->clGetPlatformInfo(platforms[i],CL_PLATFORM_VENDOR,length,
2287         device->vendor_name,NULL);
2288
2289       openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_NAME,0,NULL,
2290         &length);
2291       device->name=AcquireQuantumMemory(length,sizeof(*device->name));
2292       openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_NAME,length,
2293         device->name,NULL);
2294
2295       openCL_library->clGetDeviceInfo(devices[j],CL_DRIVER_VERSION,0,NULL,
2296         &length);
2297       device->version=AcquireQuantumMemory(length,sizeof(*device->version));
2298       openCL_library->clGetDeviceInfo(devices[j],CL_DRIVER_VERSION,length,
2299         device->version,NULL);
2300
2301       openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2302         sizeof(cl_uint),&device->max_clock_frequency,NULL);
2303
2304       openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2305         sizeof(cl_uint),&device->max_compute_units,NULL);
2306
2307       openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2308         sizeof(cl_device_type),&device->type,NULL);
2309
2310       openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2311         sizeof(cl_ulong),&device->local_memory_size,NULL);
2312
2313       clEnv->devices[next]=device;
2314     }
2315   }
2316   if (next != clEnv->number_devices)
2317     RelinquishMagickCLDevices(clEnv);
2318   platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2319   devices=(cl_device_id *) RelinquishMagickMemory(devices);
2320 }
2321
2322 MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2323   ExceptionInfo *exception)
2324 {
2325   LockSemaphoreInfo(clEnv->lock);
2326   if (clEnv->initialized != MagickFalse)
2327     {
2328       UnlockSemaphoreInfo(clEnv->lock);
2329       return(HasOpenCLDevices(clEnv,exception));
2330     }
2331   if (LoadOpenCLLibrary() != MagickFalse)
2332     {
2333       clEnv->library=openCL_library;
2334       LoadOpenCLDevices(clEnv);
2335       if (clEnv->number_devices > 0)
2336         AutoSelectOpenCLDevices(clEnv);
2337     }
2338   clEnv->initialized=MagickTrue;
2339   UnlockSemaphoreInfo(clEnv->lock);
2340   return(HasOpenCLDevices(clEnv,exception));
2341 }
2342
2343 /*
2344 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2345 %                                                                             %
2346 %                                                                             %
2347 %                                                                             %
2348 %   L o a d O p e n C L L i b r a r y                                         %
2349 %                                                                             %
2350 %                                                                             %
2351 %                                                                             %
2352 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2353 %
2354 %  LoadOpenCLLibrary() load and binds the OpenCL library.
2355 %
2356 %  The format of the LoadOpenCLLibrary method is:
2357 %
2358 %    MagickBooleanType LoadOpenCLLibrary(void)
2359 %
2360 */
2361
2362 void *OsLibraryGetFunctionAddress(void *library,const char *functionName)
2363 {
2364   if ((library == (void *) NULL) || (functionName == (const char *) NULL))
2365     return (void *) NULL;
2366 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2367     return (void *) GetProcAddress((HMODULE)library,functionName);
2368 #else
2369     return (void *) dlsym(library,functionName);
2370 #endif
2371 }
2372
2373 static MagickBooleanType BindOpenCLFunctions()
2374 {
2375 #ifdef MAGICKCORE_OPENCL_MACOSX
2376 #define BIND(X) openCL_library->X= &X;
2377 #else
2378   (void) ResetMagickMemory(openCL_library,0,sizeof(MagickLibrary));
2379 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2380   openCL_library->library=(void *)LoadLibraryA("OpenCL.dll");
2381 #else
2382   openCL_library->library=(void *)dlopen("libOpenCL.so", RTLD_NOW);
2383 #endif
2384 #define BIND(X) \
2385   if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2386     return(MagickFalse);
2387 #endif
2388
2389   if (openCL_library->library == (void*) NULL)
2390     return(MagickFalse);
2391
2392   BIND(clGetPlatformIDs);
2393   BIND(clGetPlatformInfo);
2394
2395   BIND(clGetDeviceIDs);
2396   BIND(clGetDeviceInfo);
2397
2398   BIND(clCreateBuffer);
2399   BIND(clReleaseMemObject);
2400
2401   BIND(clCreateContext);
2402   BIND(clReleaseContext);
2403
2404   BIND(clCreateCommandQueue);
2405   BIND(clReleaseCommandQueue);
2406   BIND(clFlush);
2407   BIND(clFinish);
2408
2409   BIND(clCreateProgramWithSource);
2410   BIND(clCreateProgramWithBinary);
2411   BIND(clReleaseProgram);
2412   BIND(clBuildProgram);
2413   BIND(clGetProgramBuildInfo);
2414   BIND(clGetProgramInfo);
2415
2416   BIND(clCreateKernel);
2417   BIND(clReleaseKernel);
2418   BIND(clSetKernelArg);
2419   BIND(clGetKernelInfo);
2420
2421   BIND(clEnqueueReadBuffer);
2422   BIND(clEnqueueMapBuffer);
2423   BIND(clEnqueueUnmapMemObject);
2424   BIND(clEnqueueNDRangeKernel);
2425
2426   BIND(clGetEventInfo);
2427   BIND(clWaitForEvents);
2428   BIND(clReleaseEvent);
2429   BIND(clRetainEvent);
2430   BIND(clSetEventCallback);
2431
2432   BIND(clGetEventProfilingInfo);
2433
2434   return(MagickTrue);
2435 }
2436
2437 static MagickBooleanType LoadOpenCLLibrary(void)
2438 {
2439   openCL_library=(MagickLibrary *) AcquireMagickMemory(sizeof(MagickLibrary));
2440   if (openCL_library == (MagickLibrary *) NULL)
2441     return(MagickFalse);
2442
2443   if (BindOpenCLFunctions() == MagickFalse)
2444     {
2445       openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2446       return(MagickFalse);
2447     }
2448
2449   return(MagickTrue);
2450 }
2451
2452 /*
2453 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2454 %                                                                             %
2455 %                                                                             %
2456 %                                                                             %
2457 +   O p e n C L T e r m i n u s                                               %
2458 %                                                                             %
2459 %                                                                             %
2460 %                                                                             %
2461 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2462 %
2463 %  OpenCLTerminus() destroys the OpenCL component.
2464 %
2465 %  The format of the OpenCLTerminus method is:
2466 %
2467 %      OpenCLTerminus(void)
2468 %
2469 */
2470
2471 MagickPrivate void OpenCLTerminus()
2472 {
2473   DumpOpenCLProfileData();
2474   if (cache_directory != (char *) NULL)
2475     cache_directory=DestroyString(cache_directory);
2476   if (cache_directory_lock != (SemaphoreInfo *) NULL)
2477     RelinquishSemaphoreInfo(&cache_directory_lock);
2478   if (default_CLEnv != (MagickCLEnv) NULL)
2479     default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2480   if (openCL_lock != (SemaphoreInfo *) NULL)
2481     RelinquishSemaphoreInfo(&openCL_lock);
2482   if (openCL_library != (MagickLibrary *) NULL)
2483     {
2484       if (openCL_library->library != (void *) NULL)
2485         (void) lt_dlclose(openCL_library->library);
2486       openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2487     }
2488 }
2489
2490 /*
2491 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2492 %                                                                             %
2493 %                                                                             %
2494 %                                                                             %
2495 +   O p e n C L T h r o w M a g i c k E x c e p t i o n                       %
2496 %                                                                             %
2497 %                                                                             %
2498 %                                                                             %
2499 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2500 %
2501 %  OpenCLThrowMagickException logs an OpenCL exception as determined by the log
2502 %  configuration file.  If an error occurs, MagickFalse is returned
2503 %  otherwise MagickTrue.
2504 %
2505 %  The format of the OpenCLThrowMagickException method is:
2506 %
2507 %      MagickBooleanType ThrowFileException(ExceptionInfo *exception,
2508 %        const char *module,const char *function,const size_t line,
2509 %        const ExceptionType severity,const char *tag,const char *format,...)
2510 %
2511 %  A description of each parameter follows:
2512 %
2513 %    o exception: the exception info.
2514 %
2515 %    o filename: the source module filename.
2516 %
2517 %    o function: the function name.
2518 %
2519 %    o line: the line number of the source module.
2520 %
2521 %    o severity: Specifies the numeric error category.
2522 %
2523 %    o tag: the locale tag.
2524 %
2525 %    o format: the output format.
2526 %
2527 */
2528
2529 MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2530   MagickCLDevice device,ExceptionInfo *exception,const char *module,
2531   const char *function,const size_t line,const ExceptionType severity,
2532   const char *tag,const char *format,...)
2533 {
2534   MagickBooleanType
2535     status;
2536
2537   assert(device != (MagickCLDevice) NULL);
2538   assert(exception != (ExceptionInfo *) NULL);
2539   assert(exception->signature == MagickCoreSignature);
2540
2541   status=MagickTrue;
2542   if (severity != 0)
2543   {
2544     if (device->type == CL_DEVICE_TYPE_CPU)
2545     {
2546       /* Workaround for Intel OpenCL CPU runtime bug */
2547       /* Turn off OpenCL when a problem is detected! */
2548       if (strncmp(device->platform_name, "Intel",5) == 0)
2549         default_CLEnv->enabled=MagickFalse;
2550     }
2551   }
2552
2553 #ifdef OPENCLLOG_ENABLED
2554   {
2555     va_list
2556       operands;
2557     va_start(operands,format);
2558     status=ThrowMagickExceptionList(exception,module,function,line,severity,tag,
2559       format,operands);
2560     va_end(operands);
2561   }
2562 #else
2563   magick_unreferenced(module);
2564   magick_unreferenced(function);
2565   magick_unreferenced(line);
2566   magick_unreferenced(tag);
2567   magick_unreferenced(format);
2568 #endif
2569
2570   return(status);
2571 }
2572
2573 /*
2574 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2575 %                                                                             %
2576 %                                                                             %
2577 %                                                                             %
2578 +   R e c o r d P r o f i l e D a t a                                         %
2579 %                                                                             %
2580 %                                                                             %
2581 %                                                                             %
2582 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2583 %
2584 %  RecordProfileData() records profile data.
2585 %
2586 %  The format of the RecordProfileData method is:
2587 %
2588 %      void RecordProfileData(MagickCLDevice device,ProfiledKernels kernel,
2589 %        cl_event event)
2590 %
2591 %  A description of each parameter follows:
2592 %
2593 %    o device: the OpenCL device that did the operation.
2594 %
2595 %    o event: the event that contains the profiling data.
2596 %
2597 */
2598
2599 MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2600   cl_kernel kernel,cl_event event)
2601 {
2602   char
2603     *name;
2604
2605   cl_int
2606     status;
2607
2608   cl_ulong
2609     elapsed,
2610     end,
2611     start;
2612
2613   KernelProfileRecord
2614     profile_record;
2615
2616   size_t
2617     i,
2618     length;
2619
2620   if (device->profile_kernels == MagickFalse)
2621     return(MagickFalse);
2622   status=openCL_library->clWaitForEvents(1,&event);
2623   if (status != CL_SUCCESS)
2624     return(MagickFalse);
2625   status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2626     &length);
2627   if (status != CL_SUCCESS)
2628     return(MagickTrue);
2629   name=AcquireQuantumMemory(length,sizeof(*name));
2630   if (name == (char *) NULL)
2631     return(MagickTrue);
2632   start=end=elapsed=0;
2633   status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2634     name,(size_t *) NULL);
2635   status|=openCL_library->clGetEventProfilingInfo(event,
2636     CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&start,NULL);
2637   status|=openCL_library->clGetEventProfilingInfo(event,
2638     CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL);
2639   if (status != CL_SUCCESS)
2640     {
2641       name=DestroyString(name);
2642       return(MagickTrue);
2643     }
2644   start/=1000; // usecs
2645   end/=1000;   // usecs
2646   elapsed=end-start;
2647   LockSemaphoreInfo(device->lock);
2648   i=0;
2649   profile_record=(KernelProfileRecord) NULL;
2650   if (device->profile_records != (KernelProfileRecord *) NULL)
2651     {
2652       while (device->profile_records[i] != (KernelProfileRecord) NULL)
2653       {
2654         if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2655           {
2656             profile_record=device->profile_records[i];
2657             break;
2658           }
2659         i++;
2660       }
2661     }
2662   if (profile_record != (KernelProfileRecord) NULL)
2663     name=DestroyString(name);
2664   else
2665     {
2666       profile_record=AcquireMagickMemory(sizeof(*profile_record));
2667       (void) ResetMagickMemory(profile_record,0,sizeof(*profile_record));
2668       profile_record->kernel_name=name;
2669       device->profile_records=ResizeMagickMemory(device->profile_records,(i+2)*
2670         sizeof(*device->profile_records));
2671       device->profile_records[i]=profile_record;
2672       device->profile_records[i+1]=(KernelProfileRecord) NULL;
2673     }
2674   if ((elapsed < profile_record->min) || (profile_record->count == 0))
2675     profile_record->min=elapsed;
2676   if (elapsed > profile_record->max)
2677     profile_record->max=elapsed;
2678   profile_record->total+=elapsed;
2679   profile_record->count+=1;
2680   UnlockSemaphoreInfo(device->lock);
2681   return(MagickTrue);
2682 }
2683
2684 /*
2685 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2686 %                                                                             %
2687 %                                                                             %
2688 %                                                                             %
2689 +  R e l e a s e O p e n C L C o m m a n d Q u e u e                          %
2690 %                                                                             %
2691 %                                                                             %
2692 %                                                                             %
2693 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2694 %
2695 %  ReleaseOpenCLCommandQueue() releases the OpenCL command queue
2696 %
2697 %  The format of the ReleaseOpenCLCommandQueue method is:
2698 %
2699 %      void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2700 %        cl_command_queue queue)
2701 %
2702 %  A description of each parameter follows:
2703 %
2704 %    o device: the OpenCL device.
2705 %
2706 %    o queue: the OpenCL queue to be released.
2707 */
2708
2709 MagickPrivate void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2710   cl_command_queue queue)
2711 {
2712   if (queue == (cl_command_queue) NULL)
2713     return;
2714
2715   assert(device != (MagickCLDevice) NULL);
2716   LockSemaphoreInfo(device->lock);
2717   if ((device->profile_kernels != MagickFalse) ||
2718       (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2719     {
2720       UnlockSemaphoreInfo(device->lock);
2721       openCL_library->clFinish(queue);
2722       (void) openCL_library->clReleaseCommandQueue(queue);
2723     }
2724   else
2725     {
2726       openCL_library->clFlush(queue);
2727       device->command_queues[++device->command_queues_index]=queue;
2728       UnlockSemaphoreInfo(device->lock);
2729     }
2730 }
2731
2732 /*
2733 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2734 %                                                                             %
2735 %                                                                             %
2736 %                                                                             %
2737 +   R e l e a s e  M a g i c k C L D e v i c e                                %
2738 %                                                                             %
2739 %                                                                             %
2740 %                                                                             %
2741 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2742 %
2743 %  ReleaseOpenCLDevice() returns the OpenCL device to the environment
2744 %
2745 %  The format of the ReleaseOpenCLDevice method is:
2746 %
2747 %      void ReleaseOpenCLDevice(MagickCLDevice device)
2748 %
2749 %  A description of each parameter follows:
2750 %
2751 %    o device: the OpenCL device to be released.
2752 %
2753 */
2754
2755 MagickPrivate void ReleaseOpenCLDevice(MagickCLDevice device)
2756 {
2757   assert(device != (MagickCLDevice) NULL);
2758   LockSemaphoreInfo(openCL_lock);
2759   device->requested--;
2760   UnlockSemaphoreInfo(openCL_lock);
2761 }
2762
2763 /*
2764 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2765 %                                                                             %
2766 %                                                                             %
2767 %                                                                             %
2768 +   R e l i n q u i s h M a g i c k C L C a c h e I n f o                     %
2769 %                                                                             %
2770 %                                                                             %
2771 %                                                                             %
2772 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2773 %
2774 %  RelinquishMagickCLCacheInfo() frees memory acquired with
2775 %  AcquireMagickCLCacheInfo()
2776 %
2777 %  The format of the RelinquishMagickCLCacheInfo method is:
2778 %
2779 %      MagickCLCacheInfo RelinquishMagickCLCacheInfo(MagickCLCacheInfo info,
2780 %        const MagickBooleanType relinquish_pixels)
2781 %
2782 %  A description of each parameter follows:
2783 %
2784 %    o info: the OpenCL cache info.
2785 %
2786 %    o relinquish_pixels: the pixels will be relinquish when set to true.
2787 %
2788 */
2789 static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
2790 {
2791   ssize_t
2792     i;
2793
2794   for (i=0; i < (ssize_t) info->event_count; i++)
2795     openCL_library->clReleaseEvent(info->events[i]);
2796   info->events=(cl_event *) RelinquishMagickMemory(info->events);
2797   if (info->buffer != (cl_mem) NULL)
2798     openCL_library->clReleaseMemObject(info->buffer);
2799   ReleaseOpenCLDevice(info->device);
2800   RelinquishMagickMemory(info);
2801 }
2802
2803 static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2804   cl_event magick_unused(event),
2805   cl_int magick_unused(event_command_exec_status),void *user_data)
2806 {
2807   MagickCLCacheInfo
2808     info;
2809
2810   Quantum
2811     *pixels;
2812
2813   magick_unreferenced(event);
2814   magick_unreferenced(event_command_exec_status);
2815   info=(MagickCLCacheInfo) user_data;
2816   pixels=info->pixels;
2817   RelinquishMagickResource(MemoryResource,info->length);
2818   DestroyMagickCLCacheInfo(info);
2819   (void) RelinquishAlignedMemory(pixels);
2820 }
2821
2822 MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2823   MagickCLCacheInfo info,const MagickBooleanType relinquish_pixels)
2824 {
2825   if (info == (MagickCLCacheInfo) NULL)
2826     return((MagickCLCacheInfo) NULL);
2827   if (relinquish_pixels != MagickFalse)
2828     {
2829       MagickBooleanType
2830         events_completed;
2831
2832       ssize_t
2833         i;
2834
2835       events_completed=MagickTrue;
2836       for (i=0; i < (ssize_t)info->event_count; i++)
2837       {
2838         cl_int
2839           event_status;
2840
2841         cl_uint
2842           status;
2843
2844         status=openCL_library->clGetEventInfo(info->events[i],
2845           CL_EVENT_COMMAND_EXECUTION_STATUS,sizeof(cl_int),&event_status,NULL);
2846         if ((status == CL_SUCCESS) && (event_status != CL_COMPLETE))
2847           {
2848             events_completed=MagickFalse;
2849             break;
2850           }
2851       }
2852       if (events_completed == MagickFalse)
2853         openCL_library->clSetEventCallback(info->events[info->event_count-1],
2854           CL_COMPLETE,&DestroyMagickCLCacheInfoAndPixels,info);
2855       else
2856         DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2857     }
2858   else
2859     DestroyMagickCLCacheInfo(info);
2860   return((MagickCLCacheInfo) NULL);
2861 }
2862
2863 /*
2864 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2865 %                                                                             %
2866 %                                                                             %
2867 %                                                                             %
2868 %   R e l i n q u i s h M a g i c k C L D e v i c e                           %
2869 %                                                                             %
2870 %                                                                             %
2871 %                                                                             %
2872 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2873 %
2874 %  RelinquishMagickCLDevice() releases the OpenCL device
2875 %
2876 %  The format of the RelinquishMagickCLDevice method is:
2877 %
2878 %      MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2879 %
2880 %  A description of each parameter follows:
2881 %
2882 %    o device: the OpenCL device to be released.
2883 %
2884 */
2885
2886 static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2887 {
2888   if (device == (MagickCLDevice) NULL)
2889     return((MagickCLDevice) NULL);
2890
2891   device->platform_name=RelinquishMagickMemory(device->platform_name);
2892   device->vendor_name=RelinquishMagickMemory(device->vendor_name);
2893   device->name=RelinquishMagickMemory(device->name);
2894   device->version=RelinquishMagickMemory(device->version);
2895   if (device->program != (cl_program) NULL)
2896     (void) openCL_library->clReleaseProgram(device->program);
2897   while (device->command_queues_index >= 0)
2898     (void) openCL_library->clReleaseCommandQueue(
2899       device->command_queues[device->command_queues_index--]);
2900   RelinquishSemaphoreInfo(&device->lock);
2901   return((MagickCLDevice) RelinquishMagickMemory(device));
2902 }
2903
2904 /*
2905 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2906 %                                                                             %
2907 %                                                                             %
2908 %                                                                             %
2909 %   R e l i n q u i s h M a g i c k C L E n v                                 %
2910 %                                                                             %
2911 %                                                                             %
2912 %                                                                             %
2913 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2914 %
2915 %  RelinquishMagickCLEnv() releases the OpenCL environment
2916 %
2917 %  The format of the RelinquishMagickCLEnv method is:
2918 %
2919 %      MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv device)
2920 %
2921 %  A description of each parameter follows:
2922 %
2923 %    o clEnv: the OpenCL environment to be released.
2924 %
2925 */
2926
2927 static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
2928 {
2929   if (clEnv == (MagickCLEnv) NULL)
2930     return((MagickCLEnv) NULL);
2931
2932   RelinquishSemaphoreInfo(&clEnv->lock);
2933   RelinquishMagickCLDevices(clEnv);
2934   if (clEnv->contexts != (cl_context *) NULL)
2935     {
2936       ssize_t
2937         i;
2938
2939       for (i=0; i < clEnv->number_contexts; i++)
2940          (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
2941       clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
2942     }
2943   return((MagickCLEnv) RelinquishMagickMemory(clEnv));
2944 }
2945
2946 /*
2947 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2948 %                                                                             %
2949 %                                                                             %
2950 %                                                                             %
2951 +   R e q u e s t O p e n C L D e v i c e                                     %
2952 %                                                                             %
2953 %                                                                             %
2954 %                                                                             %
2955 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2956 %
2957 %  RequestOpenCLDevice() returns one of the enabled OpenCL devices.
2958 %
2959 %  The format of the RequestOpenCLDevice method is:
2960 %
2961 %      MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
2962 %
2963 %  A description of each parameter follows:
2964 %
2965 %    o clEnv: the OpenCL environment.
2966 */
2967
2968 MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
2969 {
2970   MagickCLDevice
2971     device;
2972
2973   double
2974     score,
2975     best_score;
2976
2977   size_t
2978     i;
2979
2980   if (clEnv == (MagickCLEnv) NULL)
2981     return((MagickCLDevice) NULL);
2982
2983   if (clEnv->number_devices == 1)
2984   {
2985     if (clEnv->devices[0]->enabled)
2986       return(clEnv->devices[0]);
2987     else
2988       return((MagickCLDevice) NULL);
2989   }
2990
2991   device=(MagickCLDevice) NULL;
2992   best_score=0.0;
2993   LockSemaphoreInfo(openCL_lock);
2994   for (i = 0; i < clEnv->number_devices; i++)
2995   {
2996     if (clEnv->devices[i]->enabled == MagickFalse)
2997       continue;
2998
2999     score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3000       clEnv->devices[i]->requested);
3001     if ((device == (MagickCLDevice) NULL) || (score < best_score))
3002     {
3003       device=clEnv->devices[i];
3004       best_score=score;
3005     }
3006   }
3007   if (device != (MagickCLDevice)NULL)
3008     device->requested++;
3009   UnlockSemaphoreInfo(openCL_lock);
3010
3011   return(device);
3012 }
3013
3014 /*
3015 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3016 %                                                                             %
3017 %                                                                             %
3018 %                                                                             %
3019 %   S e t O p e n C L D e v i c e E n a b l e d                               %
3020 %                                                                             %
3021 %                                                                             %
3022 %                                                                             %
3023 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3024 %
3025 %  SetOpenCLDeviceEnabled() can be used to enable or disabled the device.
3026 %
3027 %  The format of the SetOpenCLDeviceEnabled method is:
3028 %
3029 %      void SetOpenCLDeviceEnabled(MagickCLDevice device,
3030 %        MagickBooleanType value)
3031 %
3032 %  A description of each parameter follows:
3033 %
3034 %    o device: the OpenCL device.
3035 %
3036 %    o value: determines if the device should be enabled or disabled.
3037 */
3038
3039 MagickExport void SetOpenCLDeviceEnabled(MagickCLDevice device,
3040   const MagickBooleanType value)
3041 {
3042   if (device == (MagickCLDevice) NULL)
3043     return;
3044   device->enabled=value;
3045 }
3046
3047 /*
3048 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3049 %                                                                             %
3050 %                                                                             %
3051 %                                                                             %
3052 %   S e t O p e n C L K e r n e l P r o f i l e E n a b l e d                 %
3053 %                                                                             %
3054 %                                                                             %
3055 %                                                                             %
3056 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3057 %
3058 %  SetOpenCLKernelProfileEnabled() can be used to enable or disabled the
3059 %  kernel profiling of a device.
3060 %
3061 %  The format of the SetOpenCLKernelProfileEnabled method is:
3062 %
3063 %      void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3064 %        MagickBooleanType value)
3065 %
3066 %  A description of each parameter follows:
3067 %
3068 %    o device: the OpenCL device.
3069 %
3070 %    o value: determines if kernel profiling for the device should be enabled
3071 %             or disabled.
3072 */
3073
3074 MagickExport void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3075   const MagickBooleanType value)
3076 {
3077   if (device == (MagickCLDevice) NULL)
3078     return;
3079   device->profile_kernels=value;
3080 }
3081
3082 /*
3083 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3084 %                                                                             %
3085 %                                                                             %
3086 %                                                                             %
3087 %   S e t O p e n C L E n a b l e d                                           %
3088 %                                                                             %
3089 %                                                                             %
3090 %                                                                             %
3091 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
3092 %
3093 %  SetOpenCLEnabled() can be used to enable or disable OpenCL acceleration.
3094 %
3095 %  The format of the SetOpenCLEnabled method is:
3096 %
3097 %      void SetOpenCLEnabled(MagickBooleanType)
3098 %
3099 %  A description of each parameter follows:
3100 %
3101 %    o value: specify true to enable OpenCL acceleration
3102 */
3103
3104 MagickExport MagickBooleanType SetOpenCLEnabled(const MagickBooleanType value)
3105 {
3106   MagickCLEnv
3107     clEnv;
3108
3109   clEnv=GetCurrentOpenCLEnv();
3110   if (clEnv == (MagickCLEnv) NULL)
3111     return(MagickFalse);
3112   clEnv->enabled=value;
3113   return(clEnv->enabled);
3114 }
3115
3116 #else
3117
3118 MagickExport double GetOpenCLDeviceBenchmarkScore(
3119   const MagickCLDevice magick_unused(device))
3120 {
3121   magick_unreferenced(device);
3122   return(0.0);
3123 }
3124
3125 MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3126   const MagickCLDevice magick_unused(device))
3127 {
3128   magick_unreferenced(device);
3129   return(MagickFalse);
3130 }
3131
3132 MagickExport const char *GetOpenCLDeviceName(
3133   const MagickCLDevice magick_unused(device))
3134 {
3135   magick_unreferenced(device);
3136   return((const char *) NULL);
3137 }
3138
3139 MagickExport MagickCLDevice *GetOpenCLDevices(size_t *length,
3140   ExceptionInfo *magick_unused(exception))
3141 {
3142   magick_unreferenced(exception);
3143   if (length != (size_t *) NULL)
3144     *length=0;
3145   return((MagickCLDevice *) NULL);
3146 }
3147
3148 MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3149   const MagickCLDevice magick_unused(device))
3150 {
3151   magick_unreferenced(device);
3152   return(UndefinedCLDeviceType);
3153 }
3154
3155 MagickExport const KernelProfileRecord *GetOpenCLKernelProfileRecords(
3156   const MagickCLDevice magick_unused(device),size_t *length)
3157 {
3158   magick_unreferenced(device);
3159   if (length != (size_t *) NULL)
3160     *length=0;
3161   return((const KernelProfileRecord *) NULL);
3162 }
3163
3164 MagickExport const char *GetOpenCLDeviceVersion(
3165   const MagickCLDevice magick_unused(device))
3166 {
3167   magick_unreferenced(device);
3168   return((const char *) NULL);
3169 }
3170
3171 MagickExport MagickBooleanType GetOpenCLEnabled(void)
3172 {
3173   return(MagickFalse);
3174 }
3175
3176 MagickExport void SetOpenCLDeviceEnabled(
3177   MagickCLDevice magick_unused(device),
3178   const MagickBooleanType magick_unused(value))
3179 {
3180   magick_unreferenced(device);
3181   magick_unreferenced(value);
3182 }
3183
3184 MagickExport MagickBooleanType SetOpenCLEnabled(
3185   const MagickBooleanType magick_unused(value))
3186 {
3187   magick_unreferenced(value);
3188   return(MagickFalse);
3189 }
3190
3191 MagickExport void SetOpenCLKernelProfileEnabled(
3192   MagickCLDevice magick_unused(device),
3193   const MagickBooleanType magick_unused(value))
3194 {
3195   magick_unreferenced(device);
3196   magick_unreferenced(value);
3197 }
3198 #endif