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