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