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