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