43 #include "MagickCore/studio.h"
44 #include "MagickCore/artifact.h"
45 #include "MagickCore/cache.h"
46 #include "MagickCore/cache-private.h"
47 #include "MagickCore/color.h"
48 #include "MagickCore/compare.h"
49 #include "MagickCore/constitute.h"
50 #include "MagickCore/configure.h"
51 #include "MagickCore/distort.h"
52 #include "MagickCore/draw.h"
53 #include "MagickCore/effect.h"
54 #include "MagickCore/exception.h"
55 #include "MagickCore/exception-private.h"
56 #include "MagickCore/fx.h"
57 #include "MagickCore/gem.h"
58 #include "MagickCore/geometry.h"
59 #include "MagickCore/image.h"
60 #include "MagickCore/image-private.h"
61 #include "MagickCore/layer.h"
62 #include "MagickCore/mime-private.h"
63 #include "MagickCore/memory_.h"
64 #include "MagickCore/memory-private.h"
65 #include "MagickCore/monitor.h"
66 #include "MagickCore/montage.h"
67 #include "MagickCore/morphology.h"
68 #include "MagickCore/nt-base.h"
69 #include "MagickCore/nt-base-private.h"
70 #include "MagickCore/opencl.h"
71 #include "MagickCore/opencl-private.h"
72 #include "MagickCore/option.h"
73 #include "MagickCore/policy.h"
74 #include "MagickCore/property.h"
75 #include "MagickCore/quantize.h"
76 #include "MagickCore/quantum.h"
77 #include "MagickCore/random_.h"
78 #include "MagickCore/random-private.h"
79 #include "MagickCore/resample.h"
80 #include "MagickCore/resource_.h"
81 #include "MagickCore/splay-tree.h"
82 #include "MagickCore/semaphore.h"
83 #include "MagickCore/statistic.h"
84 #include "MagickCore/string_.h"
85 #include "MagickCore/string-private.h"
86 #include "MagickCore/token.h"
87 #include "MagickCore/utility.h"
88 #include "MagickCore/utility-private.h"
90 #if defined(MAGICKCORE_OPENCL_SUPPORT)
91 #if defined(MAGICKCORE_LTDL_DELEGATE)
95 #ifndef MAGICKCORE_WINDOWS_SUPPORT
102 #define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile.xml"
128 } MagickCLDeviceBenchmark;
134 static MagickBooleanType
136 LoadOpenCLLibrary(
void);
138 static MagickCLDevice
139 RelinquishMagickCLDevice(MagickCLDevice);
142 RelinquishMagickCLEnv(MagickCLEnv);
145 BenchmarkOpenCLDevices(MagickCLEnv);
148 *accelerateKernels, *accelerateKernels2;
166 *cache_directory_lock;
168 static inline MagickBooleanType IsSameOpenCLDevice(MagickCLDevice a,
171 if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
172 (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
173 (LocaleCompare(a->name,b->name) == 0) &&
174 (LocaleCompare(a->version,b->version) == 0) &&
175 (a->max_clock_frequency == b->max_clock_frequency) &&
176 (a->max_compute_units == b->max_compute_units))
182 static inline MagickBooleanType IsBenchmarkedOpenCLDevice(MagickCLDevice a,
183 MagickCLDeviceBenchmark *b)
185 if ((LocaleCompare(a->platform_name,b->platform_name) == 0) &&
186 (LocaleCompare(a->vendor_name,b->vendor_name) == 0) &&
187 (LocaleCompare(a->name,b->name) == 0) &&
188 (LocaleCompare(a->version,b->version) == 0) &&
189 (a->max_clock_frequency == b->max_clock_frequency) &&
190 (a->max_compute_units == b->max_compute_units))
196 static inline void RelinquishMagickCLDevices(MagickCLEnv clEnv)
201 if (clEnv->devices != (MagickCLDevice *) NULL)
203 for (i = 0; i < clEnv->number_devices; i++)
204 clEnv->devices[i]=RelinquishMagickCLDevice(clEnv->devices[i]);
205 clEnv->devices=(MagickCLDevice *) RelinquishMagickMemory(clEnv->devices);
207 clEnv->number_devices=0;
210 static inline MagickBooleanType MagickCreateDirectory(
const char *path)
215 #ifdef MAGICKCORE_WINDOWS_SUPPORT
218 status=mkdir(path,0777);
220 return(status == 0 ? MagickTrue : MagickFalse);
223 static inline void InitAccelerateTimer(AccelerateTimer *timer)
226 QueryPerformanceFrequency((LARGE_INTEGER*)&timer->freq);
228 timer->freq=(
long long)1.0E3;
234 static inline double ReadAccelerateTimer(AccelerateTimer *timer)
236 return (
double)timer->clocks/(double)timer->freq;
239 static inline void StartAccelerateTimer(AccelerateTimer* timer)
242 QueryPerformanceCounter((LARGE_INTEGER*)&timer->start);
247 timer->start=(
long long)s.tv_sec*(
long long)1.0E3+(
long long)s.tv_usec/
252 static inline void StopAccelerateTimer(AccelerateTimer *timer)
259 QueryPerformanceCounter((LARGE_INTEGER*)&(n));
264 n=(
long long)s.tv_sec*(
long long)1.0E3+(
long long)s.tv_usec/
272 static const char *GetOpenCLCacheDirectory()
274 if (cache_directory == (
char *) NULL)
277 ActivateSemaphoreInfo(&cache_directory_lock);
278 LockSemaphoreInfo(cache_directory_lock);
279 if (cache_directory == (
char *) NULL)
283 path[MagickPathExtent],
293 home=GetEnvironmentValue(
"MAGICK_OPENCL_CACHE_DIR");
294 if (home == (
char *) NULL)
296 home=GetEnvironmentValue(
"XDG_CACHE_HOME");
297 #if defined(MAGICKCORE_WINDOWS_SUPPORT) || defined(__MINGW32__)
298 if (home == (
char *) NULL)
299 home=GetEnvironmentValue(
"LOCALAPPDATA");
300 if (home == (
char *) NULL)
301 home=GetEnvironmentValue(
"APPDATA");
302 if (home == (
char *) NULL)
303 home=GetEnvironmentValue(
"USERPROFILE");
307 if (home != (
char *) NULL)
310 (void) FormatLocaleString(path,MagickPathExtent,
"%s",home);
311 status=GetPathAttributes(path,&attributes);
312 if (status == MagickFalse)
313 status=MagickCreateDirectory(path);
316 if (status != MagickFalse)
318 (void) FormatLocaleString(path,MagickPathExtent,
319 "%s%sImageMagick",home,DirectorySeparator);
321 status=GetPathAttributes(path,&attributes);
322 if (status == MagickFalse)
323 status=MagickCreateDirectory(path);
326 if (status != MagickFalse)
328 temp=(
char*) AcquireCriticalMemory(strlen(path)+1);
329 CopyMagickString(temp,path,strlen(path)+1);
331 home=DestroyString(home);
335 home=GetEnvironmentValue(
"HOME");
336 if (home != (
char *) NULL)
339 (void) FormatLocaleString(path,MagickPathExtent,
"%s%s.cache",
340 home,DirectorySeparator);
341 status=GetPathAttributes(path,&attributes);
342 if (status == MagickFalse)
343 status=MagickCreateDirectory(path);
346 if (status != MagickFalse)
348 (void) FormatLocaleString(path,MagickPathExtent,
349 "%s%s.cache%sImageMagick",home,DirectorySeparator,
351 status=GetPathAttributes(path,&attributes);
352 if (status == MagickFalse)
353 status=MagickCreateDirectory(path);
356 if (status != MagickFalse)
358 temp=(
char*) AcquireCriticalMemory(strlen(path)+1);
359 CopyMagickString(temp,path,strlen(path)+1);
361 home=DestroyString(home);
364 if (temp == (
char *) NULL)
366 temp=AcquireString(
"?");
367 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
368 "Cannot use cache directory: \"%s\"",path);
371 (
void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
372 "Using cache directory: \"%s\"",temp);
373 cache_directory=temp;
375 UnlockSemaphoreInfo(cache_directory_lock);
377 if (*cache_directory ==
'?')
378 return((
const char *) NULL);
379 return(cache_directory);
382 static void SelectOpenCLDevice(MagickCLEnv clEnv,cl_device_type type)
391 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
392 "Selecting device for type: %d",(int) type);
393 for (i = 0; i < clEnv->number_devices; i++)
394 clEnv->devices[i]->enabled=MagickFalse;
396 for (i = 0; i < clEnv->number_devices; i++)
398 device=clEnv->devices[i];
399 if (device->type != type)
402 device->enabled=MagickTrue;
403 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
404 "Selected device: %s",device->name);
405 for (j = i+1; j < clEnv->number_devices; j++)
410 other_device=clEnv->devices[j];
411 if (IsSameOpenCLDevice(device,other_device))
412 other_device->enabled=MagickTrue;
417 static size_t StringSignature(
const char*
string)
432 stringLength=(size_t) strlen(
string);
433 signature=stringLength;
434 n=stringLength/
sizeof(size_t);
436 for (i = 0; i < n; i++)
438 if (n *
sizeof(
size_t) != stringLength)
444 for (i = 0; i < 4; i++, j++)
446 if (j < stringLength)
457 static void DestroyMagickCLCacheInfo(MagickCLCacheInfo info)
462 for (i=0; i < (ssize_t) info->event_count; i++)
463 openCL_library->clReleaseEvent(info->events[i]);
464 info->events=(cl_event *) RelinquishMagickMemory(info->events);
465 if (info->buffer != (cl_mem) NULL)
466 openCL_library->clReleaseMemObject(info->buffer);
467 RelinquishSemaphoreInfo(&info->events_semaphore);
468 ReleaseOpenCLDevice(info->device);
469 RelinquishMagickMemory(info);
476 MagickPrivate cl_mem CreateOpenCLBuffer(MagickCLDevice device,
477 cl_mem_flags flags,
size_t size,
void *host_ptr)
479 return(openCL_library->clCreateBuffer(device->context,flags,size,host_ptr,
483 MagickPrivate
void ReleaseOpenCLKernel(cl_kernel kernel)
485 (void) openCL_library->clReleaseKernel(kernel);
488 MagickPrivate
void ReleaseOpenCLMemObject(cl_mem memobj)
490 (void) openCL_library->clReleaseMemObject(memobj);
493 MagickPrivate
void RetainOpenCLMemObject(cl_mem memobj)
495 (void) openCL_library->clRetainMemObject(memobj);
498 MagickPrivate cl_int SetOpenCLKernelArg(cl_kernel kernel,
size_t arg_index,
499 size_t arg_size,
const void *arg_value)
501 return(openCL_library->clSetKernelArg(kernel,(cl_uint) arg_index,arg_size,
533 MagickPrivate MagickCLCacheInfo AcquireMagickCLCacheInfo(MagickCLDevice device,
534 Quantum *pixels,
const MagickSizeType length)
542 info=(MagickCLCacheInfo) AcquireCriticalMemory(
sizeof(*info));
543 (void) memset(info,0,
sizeof(*info));
544 LockSemaphoreInfo(openCL_lock);
546 UnlockSemaphoreInfo(openCL_lock);
550 info->events_semaphore=AcquireSemaphoreInfo();
551 info->buffer=openCL_library->clCreateBuffer(device->context,
552 CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,(
size_t) length,(
void *) pixels,
554 if (status == CL_SUCCESS)
556 DestroyMagickCLCacheInfo(info);
557 return((MagickCLCacheInfo) NULL);
579 static MagickCLDevice AcquireMagickCLDevice()
584 device=(MagickCLDevice) AcquireMagickMemory(
sizeof(*device));
587 (void) memset(device,0,
sizeof(*device));
588 ActivateSemaphoreInfo(&device->lock);
589 device->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
590 device->command_queues_index=-1;
591 device->enabled=MagickTrue;
611 static MagickCLEnv AcquireMagickCLEnv(
void)
619 clEnv=(MagickCLEnv) AcquireMagickMemory(
sizeof(*clEnv));
620 if (clEnv != (MagickCLEnv) NULL)
622 (void) memset(clEnv,0,
sizeof(*clEnv));
623 ActivateSemaphoreInfo(&clEnv->lock);
624 clEnv->cpu_score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
625 clEnv->enabled=MagickFalse;
626 option=getenv(
"MAGICK_OCL_DEVICE");
627 if (option != (
const char *) NULL)
629 if ((IsStringTrue(option) != MagickFalse) ||
630 (strcmp(option,
"GPU") == 0) ||
631 (strcmp(option,
"CPU") == 0))
632 clEnv->enabled=MagickTrue;
661 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLDevice device)
666 cl_command_queue_properties
669 assert(device != (MagickCLDevice) NULL);
670 LockSemaphoreInfo(device->lock);
671 if ((device->profile_kernels == MagickFalse) &&
672 (device->command_queues_index >= 0))
674 queue=device->command_queues[device->command_queues_index--];
675 UnlockSemaphoreInfo(device->lock);
679 UnlockSemaphoreInfo(device->lock);
681 if (device->profile_kernels != MagickFalse)
682 properties=CL_QUEUE_PROFILING_ENABLE;
683 queue=openCL_library->clCreateCommandQueue(device->context,
684 device->deviceID,properties,(cl_int *) NULL);
717 MagickPrivate cl_kernel AcquireOpenCLKernel(MagickCLDevice device,
718 const char *kernel_name)
723 assert(device != (MagickCLDevice) NULL);
724 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
"Using kernel: %s",
726 kernel=openCL_library->clCreateKernel(device->program,kernel_name,
757 static void LoadOpenCLDeviceBenchmark(MagickCLEnv clEnv,
const char *xml)
760 keyword[MagickPathExtent],
766 MagickCLDeviceBenchmark
773 if (xml == (
char *) NULL)
775 device_benchmark=(MagickCLDeviceBenchmark *) NULL;
776 token=AcquireString(xml);
777 extent=strlen(token)+MagickPathExtent;
778 for (q=(
char *) xml; *q !=
'\0'; )
783 (void) GetNextToken(q,&q,extent,token);
786 (void) CopyMagickString(keyword,token,MagickPathExtent);
787 if (LocaleNCompare(keyword,
"<!DOCTYPE",9) == 0)
792 while ((LocaleNCompare(q,
"]>",2) != 0) && (*q !=
'\0'))
793 (
void) GetNextToken(q,&q,extent,token);
796 if (LocaleNCompare(keyword,
"<!--",4) == 0)
801 while ((LocaleNCompare(q,
"->",2) != 0) && (*q !=
'\0'))
802 (void) GetNextToken(q,&q,extent,token);
805 if (LocaleCompare(keyword,
"<device") == 0)
810 device_benchmark=(MagickCLDeviceBenchmark *) AcquireQuantumMemory(1,
811 sizeof(*device_benchmark));
812 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
814 (void) memset(device_benchmark,0,
sizeof(*device_benchmark));
815 device_benchmark->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
818 if (device_benchmark == (MagickCLDeviceBenchmark *) NULL)
820 if (LocaleCompare(keyword,
"/>") == 0)
822 if (device_benchmark->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
824 if (LocaleCompare(device_benchmark->name,
"CPU") == 0)
825 clEnv->cpu_score=device_benchmark->score;
834 for (i = 0; i < clEnv->number_devices; i++)
836 device=clEnv->devices[i];
837 if (IsBenchmarkedOpenCLDevice(device,device_benchmark))
838 device->score=device_benchmark->score;
843 device_benchmark->platform_name=RelinquishMagickMemory(
844 device_benchmark->platform_name);
845 device_benchmark->vendor_name=RelinquishMagickMemory(
846 device_benchmark->vendor_name);
847 device_benchmark->name=RelinquishMagickMemory(device_benchmark->name);
848 device_benchmark->version=RelinquishMagickMemory(
849 device_benchmark->version);
850 device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
854 (void) GetNextToken(q,(
const char **) NULL,extent,token);
857 (void) GetNextToken(q,&q,extent,token);
858 (void) GetNextToken(q,&q,extent,token);
864 if (LocaleCompare((
char *) keyword,
"maxClockFrequency") == 0)
866 device_benchmark->max_clock_frequency=StringToInteger(token);
869 if (LocaleCompare((
char *) keyword,
"maxComputeUnits") == 0)
871 device_benchmark->max_compute_units=StringToInteger(token);
879 if (LocaleCompare((
char *) keyword,
"name") == 0)
880 device_benchmark->name=ConstantString(token);
886 if (LocaleCompare((
char *) keyword,
"platform") == 0)
887 device_benchmark->platform_name=ConstantString(token);
893 if (LocaleCompare((
char *) keyword,
"score") == 0)
894 device_benchmark->score=StringToDouble(token,(
char **) NULL);
900 if (LocaleCompare((
char *) keyword,
"vendor") == 0)
901 device_benchmark->vendor_name=ConstantString(token);
902 if (LocaleCompare((
char *) keyword,
"version") == 0)
903 device_benchmark->version=ConstantString(token);
910 token=(
char *) RelinquishMagickMemory(token);
911 device_benchmark=(MagickCLDeviceBenchmark *) RelinquishMagickMemory(
915 static MagickBooleanType CanWriteProfileToFile(
const char *filename)
920 profileFile=fopen(filename,
"ab");
922 if (profileFile == (FILE *) NULL)
924 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
925 "Unable to save profile to: \"%s\"",filename);
933 static MagickBooleanType LoadOpenCLBenchmarks(MagickCLEnv clEnv)
935 #if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
937 filename[MagickPathExtent];
942 (void) FormatLocaleString(filename,MagickPathExtent,
"%s%s%s",
943 GetOpenCLCacheDirectory(),DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
949 if (CanWriteProfileToFile(filename) == MagickFalse)
955 for (i = 0; i < clEnv->number_devices; i++)
956 clEnv->devices[i]->score=1.0;
958 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
961 #if !MAGICKCORE_ZERO_CONFIGURATION_SUPPORT
962 option=ConfigureFileToStringInfo(filename);
963 LoadOpenCLDeviceBenchmark(clEnv,(
const char *) GetStringInfoDatum(option));
964 option=DestroyStringInfo(option);
969 static void AutoSelectOpenCLDevices(MagickCLEnv clEnv)
983 option=getenv(
"MAGICK_OCL_DEVICE");
984 if (option != (
const char *) NULL)
986 if (strcmp(option,
"GPU") == 0)
987 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_GPU);
988 else if (strcmp(option,
"CPU") == 0)
989 SelectOpenCLDevice(clEnv,CL_DEVICE_TYPE_CPU);
992 if (LoadOpenCLBenchmarks(clEnv) == MagickFalse)
995 benchmark=MagickFalse;
996 if (clEnv->cpu_score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
997 benchmark=MagickTrue;
1000 for (i = 0; i < clEnv->number_devices; i++)
1002 if (clEnv->devices[i]->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1004 benchmark=MagickTrue;
1010 if (benchmark != MagickFalse)
1011 BenchmarkOpenCLDevices(clEnv);
1013 best_score=clEnv->cpu_score;
1014 for (i = 0; i < clEnv->number_devices; i++)
1015 best_score=MagickMin(clEnv->devices[i]->score,best_score);
1017 for (i = 0; i < clEnv->number_devices; i++)
1019 if (clEnv->devices[i]->score != best_score)
1020 clEnv->devices[i]->enabled=MagickFalse;
1049 static double RunOpenCLBenchmark(MagickBooleanType is_cpu)
1066 exception=AcquireExceptionInfo();
1067 imageInfo=AcquireImageInfo();
1068 CloneString(&imageInfo->size,
"2048x1536");
1069 CopyMagickString(imageInfo->filename,
"xc:none",MagickPathExtent);
1070 inputImage=ReadImage(imageInfo,exception);
1072 InitAccelerateTimer(&timer);
1074 for (i=0; i<=2; i++)
1082 StartAccelerateTimer(&timer);
1084 bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
1085 unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
1087 resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,
1094 if (is_cpu == MagickFalse)
1099 cache_info=(
CacheInfo *) resizedImage->cache;
1100 if (cache_info->opencl != (MagickCLCacheInfo) NULL)
1101 openCL_library->clWaitForEvents(cache_info->opencl->event_count,
1102 cache_info->opencl->events);
1106 StopAccelerateTimer(&timer);
1108 if (bluredImage != (
Image *) NULL)
1109 DestroyImage(bluredImage);
1110 if (unsharpedImage != (
Image *) NULL)
1111 DestroyImage(unsharpedImage);
1112 if (resizedImage != (
Image *) NULL)
1113 DestroyImage(resizedImage);
1115 DestroyImage(inputImage);
1116 return(ReadAccelerateTimer(&timer));
1119 static void RunDeviceBenckmark(MagickCLEnv clEnv,MagickCLEnv testEnv,
1120 MagickCLDevice device)
1122 testEnv->devices[0]=device;
1123 default_CLEnv=testEnv;
1124 device->score=RunOpenCLBenchmark(MagickFalse);
1125 default_CLEnv=clEnv;
1126 testEnv->devices[0]=(MagickCLDevice) NULL;
1129 static void CacheOpenCLBenchmarks(MagickCLEnv clEnv)
1132 filename[MagickPathExtent];
1144 (void) FormatLocaleString(filename,MagickPathExtent,
"%s%s%s",
1145 GetOpenCLCacheDirectory(),DirectorySeparator,
1146 IMAGEMAGICK_PROFILE_FILE);
1148 cache_file=fopen_utf8(filename,
"wb");
1149 if (cache_file == (FILE *) NULL)
1151 fwrite(
"<devices>\n",
sizeof(
char),10,cache_file);
1152 fprintf(cache_file,
" <device name=\"CPU\" score=\"%.4g\"/>\n",
1154 for (i = 0; i < clEnv->number_devices; i++)
1159 device=clEnv->devices[i];
1160 duplicate=MagickFalse;
1161 for (j = 0; j < i; j++)
1163 if (IsSameOpenCLDevice(clEnv->devices[j],device))
1165 duplicate=MagickTrue;
1173 if (device->score != MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1174 fprintf(cache_file,
" <device platform=\"%s\" vendor=\"%s\" name=\"%s\"\
1175 version=\"%s\" maxClockFrequency=\"%d\" maxComputeUnits=\"%d\"\
1176 score=\"%.4g\"/>\n",
1177 device->platform_name,device->vendor_name,device->name,device->version,
1178 (
int)device->max_clock_frequency,(
int)device->max_compute_units,
1181 fwrite(
"</devices>",
sizeof(
char),10,cache_file);
1186 static void BenchmarkOpenCLDevices(MagickCLEnv clEnv)
1198 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
1199 "Starting benchmark");
1200 testEnv=AcquireMagickCLEnv();
1201 testEnv->library=openCL_library;
1202 testEnv->devices=(MagickCLDevice *) AcquireCriticalMemory(
1203 sizeof(MagickCLDevice));
1204 testEnv->number_devices=1;
1205 testEnv->benchmark_thread_id=GetMagickThreadId();
1206 testEnv->initialized=MagickTrue;
1208 for (i = 0; i < clEnv->number_devices; i++)
1209 clEnv->devices[i]->score=MAGICKCORE_OPENCL_UNDEFINED_SCORE;
1211 for (i = 0; i < clEnv->number_devices; i++)
1213 device=clEnv->devices[i];
1214 if (device->score == MAGICKCORE_OPENCL_UNDEFINED_SCORE)
1215 RunDeviceBenckmark(clEnv,testEnv,device);
1218 for (j = i+1; j < clEnv->number_devices; j++)
1223 other_device=clEnv->devices[j];
1224 if (IsSameOpenCLDevice(device,other_device))
1225 other_device->score=device->score;
1229 testEnv->enabled=MagickFalse;
1230 default_CLEnv=testEnv;
1231 clEnv->cpu_score=RunOpenCLBenchmark(MagickTrue);
1232 default_CLEnv=clEnv;
1234 testEnv=RelinquishMagickCLEnv(testEnv);
1235 CacheOpenCLBenchmarks(clEnv);
1272 static void CacheOpenCLKernel(MagickCLDevice device,
char *filename,
1284 status=openCL_library->clGetProgramInfo(device->program,
1285 CL_PROGRAM_BINARY_SIZES,
sizeof(
size_t),&binaryProgramSize,NULL);
1286 if (status != CL_SUCCESS)
1288 binaryProgram=(
unsigned char*) AcquireQuantumMemory(1,binaryProgramSize);
1289 if (binaryProgram == (
unsigned char *) NULL)
1291 (void) ThrowMagickException(exception,GetMagickModule(),
1292 ResourceLimitError,
"MemoryAllocationFailed",
"`%s'",filename);
1295 status=openCL_library->clGetProgramInfo(device->program,
1296 CL_PROGRAM_BINARIES,
sizeof(
unsigned char*),&binaryProgram,NULL);
1297 if (status == CL_SUCCESS)
1299 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
1300 "Creating cache file: \"%s\"",filename);
1301 (void) BlobToFile(filename,binaryProgram,binaryProgramSize,exception);
1303 binaryProgram=(
unsigned char *) RelinquishMagickMemory(binaryProgram);
1306 static MagickBooleanType LoadCachedOpenCLKernels(MagickCLDevice device,
1307 const char *filename)
1322 sans_exception=AcquireExceptionInfo();
1323 binaryProgram=(
unsigned char *) FileToBlob(filename,~0UL,&length,
1325 sans_exception=DestroyExceptionInfo(sans_exception);
1326 if (binaryProgram == (
unsigned char *) NULL)
1327 return(MagickFalse);
1328 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
1329 "Loaded cached kernels: \"%s\"",filename);
1330 device->program=openCL_library->clCreateProgramWithBinary(device->context,1,
1331 &device->deviceID,&length,(
const unsigned char**)&binaryProgram,
1332 &binaryStatus,&status);
1333 binaryProgram=(
unsigned char *) RelinquishMagickMemory(binaryProgram);
1334 return((status != CL_SUCCESS) || (binaryStatus != CL_SUCCESS) ? MagickFalse :
1338 static void LogOpenCLBuildFailure(MagickCLDevice device,
const char *kernel,
1342 filename[MagickPathExtent],
1348 (void) FormatLocaleString(filename,MagickPathExtent,
"%s%s%s",
1349 GetOpenCLCacheDirectory(),DirectorySeparator,
"magick_badcl.cl");
1351 (void) remove_utf8(filename);
1352 (void) BlobToFile(filename,kernel,strlen(kernel),exception);
1354 openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1355 CL_PROGRAM_BUILD_LOG,0,NULL,&log_size);
1356 log=(
char*)AcquireCriticalMemory(log_size);
1357 openCL_library->clGetProgramBuildInfo(device->program,device->deviceID,
1358 CL_PROGRAM_BUILD_LOG,log_size,log,&log_size);
1360 (void) FormatLocaleString(filename,MagickPathExtent,
"%s%s%s",
1361 GetOpenCLCacheDirectory(),DirectorySeparator,
"magick_badcl.log");
1363 (void) remove_utf8(filename);
1364 (void) BlobToFile(filename,log,log_size,exception);
1365 log=(
char*)RelinquishMagickMemory(log);
1368 static MagickBooleanType CompileOpenCLKernel(MagickCLDevice device,
1369 const char *kernel,
const char *options,
size_t signature,
1373 deviceName[MagickPathExtent],
1374 filename[MagickPathExtent],
1386 (void) CopyMagickString(deviceName,device->name,MagickPathExtent);
1389 while (*ptr !=
'\0')
1391 if ((*ptr ==
' ') || (*ptr ==
'\\') || (*ptr ==
'/') || (*ptr ==
':') ||
1392 (*ptr ==
'*') || (*ptr ==
'?') || (*ptr ==
'"') || (*ptr ==
'<') ||
1393 (*ptr ==
'>' || *ptr ==
'|'))
1397 (void) FormatLocaleString(filename,MagickPathExtent,
1398 "%s%s%s_%s_%08x_%.20g.bin",GetOpenCLCacheDirectory(),
1399 DirectorySeparator,
"magick_opencl",deviceName,(
unsigned int) signature,
1400 (
double)
sizeof(
char*)*8);
1401 loaded=LoadCachedOpenCLKernels(device,filename);
1402 if (loaded == MagickFalse)
1405 length=strlen(kernel);
1406 device->program=openCL_library->clCreateProgramWithSource(
1407 device->context,1,&kernel,&length,&status);
1408 if (status != CL_SUCCESS)
1409 return(MagickFalse);
1412 status=openCL_library->clBuildProgram(device->program,1,&device->deviceID,
1414 if (status != CL_SUCCESS)
1416 (void) ThrowMagickException(exception,GetMagickModule(),DelegateWarning,
1417 "clBuildProgram failed.",
"(%d)",(int)status);
1418 LogOpenCLBuildFailure(device,kernel,exception);
1419 return(MagickFalse);
1423 if (loaded == MagickFalse)
1424 CacheOpenCLKernel(device,filename,exception);
1429 static cl_event* CopyOpenCLEvents(MagickCLCacheInfo first,
1430 MagickCLCacheInfo second,cl_uint *event_count)
1441 assert(first != (MagickCLCacheInfo) NULL);
1442 assert(event_count != (cl_uint *) NULL);
1443 events=(cl_event *) NULL;
1444 LockSemaphoreInfo(first->events_semaphore);
1445 if (second != (MagickCLCacheInfo) NULL)
1446 LockSemaphoreInfo(second->events_semaphore);
1447 *event_count=first->event_count;
1448 if (second != (MagickCLCacheInfo) NULL)
1449 *event_count+=second->event_count;
1450 if (*event_count > 0)
1452 events=AcquireQuantumMemory(*event_count,
sizeof(*events));
1453 if (events == (cl_event *) NULL)
1458 for (i=0; i < first->event_count; i++, j++)
1459 events[j]=first->events[i];
1460 if (second != (MagickCLCacheInfo) NULL)
1462 for (i=0; i < second->event_count; i++, j++)
1463 events[j]=second->events[i];
1467 UnlockSemaphoreInfo(first->events_semaphore);
1468 if (second != (MagickCLCacheInfo) NULL)
1469 UnlockSemaphoreInfo(second->events_semaphore);
1495 MagickPrivate MagickCLCacheInfo CopyMagickCLCacheInfo(MagickCLCacheInfo info)
1509 if (info == (MagickCLCacheInfo) NULL)
1510 return((MagickCLCacheInfo) NULL);
1511 events=CopyOpenCLEvents(info,(MagickCLCacheInfo) NULL,&event_count);
1512 if (events != (cl_event *) NULL)
1514 queue=AcquireOpenCLCommandQueue(info->device);
1515 pixels=openCL_library->clEnqueueMapBuffer(queue,info->buffer,CL_TRUE,
1516 CL_MAP_READ | CL_MAP_WRITE,0,info->length,event_count,events,
1517 (cl_event *) NULL,(cl_int *) NULL);
1518 assert(pixels == info->pixels);
1519 ReleaseOpenCLCommandQueue(info->device,queue);
1520 events=(cl_event *) RelinquishMagickMemory(events);
1522 return(RelinquishMagickCLCacheInfo(info,MagickFalse));
1544 MagickPrivate
void DumpOpenCLProfileData()
1546 #define OpenCLLog(message) \
1547 fwrite(message,sizeof(char),strlen(message),log); \
1548 fwrite("\n",sizeof(char),1,log);
1552 filename[MagickPathExtent],
1562 if (default_CLEnv == (MagickCLEnv) NULL)
1565 for (i = 0; i < default_CLEnv->number_devices; i++)
1566 if (default_CLEnv->devices[i]->profile_kernels != MagickFalse)
1568 if (i == default_CLEnv->number_devices)
1571 (void) FormatLocaleString(filename,MagickPathExtent,
"%s%s%s",
1572 GetOpenCLCacheDirectory(),DirectorySeparator,
"ImageMagickOpenCL.log");
1574 log=fopen_utf8(filename,
"wb");
1575 if (log == (FILE *) NULL)
1577 for (i = 0; i < default_CLEnv->number_devices; i++)
1582 device=default_CLEnv->devices[i];
1583 if ((device->profile_kernels == MagickFalse) ||
1587 OpenCLLog(
"====================================================");
1588 fprintf(log,
"Device: %s\n",device->name);
1589 fprintf(log,
"Version: %s\n",device->version);
1590 OpenCLLog(
"====================================================");
1591 OpenCLLog(
" average calls min max");
1592 OpenCLLog(
" ------- ----- --- ---");
1599 profile=device->profile_records[j];
1601 CopyMagickString(indent,profile->kernel_name,MagickMin(strlen(
1602 profile->kernel_name),strlen(indent)));
1603 sprintf(buf,
"%s %7d %7d %7d %7d",indent,(
int) (profile->total/
1604 profile->count),(
int) profile->count,(
int) profile->min,
1605 (
int) profile->max);
1609 OpenCLLog(
"====================================================");
1610 fwrite(
"\n\n",
sizeof(
char),2,log);
1662 static MagickBooleanType RegisterCacheEvent(MagickCLCacheInfo info,
1665 assert(info != (MagickCLCacheInfo) NULL);
1666 assert(event != (cl_event) NULL);
1667 if (openCL_library->clRetainEvent(event) != CL_SUCCESS)
1669 openCL_library->clWaitForEvents(1,&event);
1670 return(MagickFalse);
1672 LockSemaphoreInfo(info->events_semaphore);
1673 if (info->events == (cl_event *) NULL)
1675 info->events=AcquireMagickMemory(
sizeof(*info->events));
1676 info->event_count=1;
1679 info->events=ResizeQuantumMemory(info->events,++info->event_count,
1680 sizeof(*info->events));
1681 if (info->events == (cl_event *) NULL)
1682 ThrowFatalException(ResourceLimitFatalError,
"MemoryAllocationFailed");
1683 info->events[info->event_count-1]=event;
1684 UnlockSemaphoreInfo(info->events_semaphore);
1688 MagickPrivate MagickBooleanType EnqueueOpenCLKernel(cl_command_queue queue,
1689 cl_kernel kernel,cl_uint work_dim,
const size_t *offset,
const size_t *gsize,
1690 const size_t *lsize,
const Image *input_image,
const Image *output_image,
1707 assert(input_image != (
const Image *) NULL);
1708 input_info=(
CacheInfo *) input_image->cache;
1709 assert(input_info != (
CacheInfo *) NULL);
1710 assert(input_info->opencl != (MagickCLCacheInfo) NULL);
1712 if (output_image == (
const Image *) NULL)
1713 events=CopyOpenCLEvents(input_info->opencl,(MagickCLCacheInfo) NULL,
1717 output_info=(
CacheInfo *) output_image->cache;
1718 assert(output_info != (
CacheInfo *) NULL);
1719 assert(output_info->opencl != (MagickCLCacheInfo) NULL);
1720 events=CopyOpenCLEvents(input_info->opencl,output_info->opencl,
1723 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,
1724 gsize,lsize,event_count,events,&event);
1726 if ((status != CL_SUCCESS) && (event_count > 0))
1728 openCL_library->clFinish(queue);
1729 status=openCL_library->clEnqueueNDRangeKernel(queue,kernel,work_dim,
1730 offset,gsize,lsize,event_count,events,&event);
1732 events=(cl_event *) RelinquishMagickMemory(events);
1733 if (status != CL_SUCCESS)
1735 (void) OpenCLThrowMagickException(input_info->opencl->device,exception,
1736 GetMagickModule(),ResourceLimitWarning,
1737 "clEnqueueNDRangeKernel failed.",
"'%s'",
".");
1738 return(MagickFalse);
1740 if (flush != MagickFalse)
1741 openCL_library->clFlush(queue);
1742 if (RecordProfileData(input_info->opencl->device,kernel,event) == MagickFalse)
1744 if (RegisterCacheEvent(input_info->opencl,event) != MagickFalse)
1747 (
void) RegisterCacheEvent(output_info->opencl,event);
1750 openCL_library->clReleaseEvent(event);
1773 MagickPrivate MagickCLEnv GetCurrentOpenCLEnv(
void)
1775 if (default_CLEnv != (MagickCLEnv) NULL)
1777 if ((default_CLEnv->benchmark_thread_id != (MagickThreadType) 0) &&
1778 (default_CLEnv->benchmark_thread_id != GetMagickThreadId()))
1779 return((MagickCLEnv) NULL);
1781 return(default_CLEnv);
1784 if (GetOpenCLCacheDirectory() == (
char *) NULL)
1785 return((MagickCLEnv) NULL);
1788 ActivateSemaphoreInfo(&openCL_lock);
1790 LockSemaphoreInfo(openCL_lock);
1791 if (default_CLEnv == (MagickCLEnv) NULL)
1792 default_CLEnv=AcquireMagickCLEnv();
1793 UnlockSemaphoreInfo(openCL_lock);
1795 return(default_CLEnv);
1822 MagickExport
double GetOpenCLDeviceBenchmarkScore(
1823 const MagickCLDevice device)
1825 if (device == (MagickCLDevice) NULL)
1826 return(MAGICKCORE_OPENCL_UNDEFINED_SCORE);
1827 return(device->score);
1852 MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
1853 const MagickCLDevice device)
1855 if (device == (MagickCLDevice) NULL)
1856 return(MagickFalse);
1857 return(device->enabled);
1882 MagickExport
const char *GetOpenCLDeviceName(
const MagickCLDevice device)
1884 if (device == (MagickCLDevice) NULL)
1885 return((
const char *) NULL);
1886 return(device->name);
1911 MagickExport
const char *GetOpenCLDeviceVendorName(
const MagickCLDevice device)
1913 if (device == (MagickCLDevice) NULL)
1914 return((
const char *) NULL);
1915 return(device->vendor_name);
1945 MagickExport MagickCLDevice *GetOpenCLDevices(
size_t *length,
1951 clEnv=GetCurrentOpenCLEnv();
1952 if (clEnv == (MagickCLEnv) NULL)
1954 if (length != (
size_t *) NULL)
1956 return((MagickCLDevice *) NULL);
1958 InitializeOpenCL(clEnv,exception);
1959 if (length != (
size_t *) NULL)
1960 *length=clEnv->number_devices;
1961 return(clEnv->devices);
1986 MagickExport MagickCLDeviceType GetOpenCLDeviceType(
1987 const MagickCLDevice device)
1989 if (device == (MagickCLDevice) NULL)
1990 return(UndefinedCLDeviceType);
1991 if (device->type == CL_DEVICE_TYPE_GPU)
1992 return(GpuCLDeviceType);
1993 if (device->type == CL_DEVICE_TYPE_CPU)
1994 return(CpuCLDeviceType);
1995 return(UndefinedCLDeviceType);
2020 MagickExport
const char *GetOpenCLDeviceVersion(
const MagickCLDevice device)
2022 if (device == (MagickCLDevice) NULL)
2023 return((
const char *) NULL);
2024 return(device->version);
2046 MagickExport MagickBooleanType GetOpenCLEnabled(
void)
2051 clEnv=GetCurrentOpenCLEnv();
2052 if (clEnv == (MagickCLEnv) NULL)
2053 return(MagickFalse);
2054 return(clEnv->enabled);
2081 const MagickCLDevice device,
size_t *length)
2083 if ((device == (
const MagickCLDevice) NULL) || (device->profile_records ==
2086 if (length != (
size_t *) NULL)
2090 if (length != (
size_t *) NULL)
2093 LockSemaphoreInfo(device->lock);
2096 UnlockSemaphoreInfo(device->lock);
2098 return(device->profile_records);
2129 static MagickBooleanType HasOpenCLDevices(MagickCLEnv clEnv,
2133 *accelerateKernelsBuffer,
2134 options[MagickPathExtent];
2146 for (i = 0; i < clEnv->number_devices; i++)
2148 if ((clEnv->devices[i]->enabled != MagickFalse))
2151 if (i == clEnv->number_devices)
2152 return(MagickFalse);
2156 for (i = 0; i < clEnv->number_devices; i++)
2158 if ((clEnv->devices[i]->enabled != MagickFalse) &&
2159 (clEnv->devices[i]->program == (cl_program) NULL))
2165 if (status != MagickFalse)
2169 (void) FormatLocaleString(options,MagickPathExtent,CLOptions,
2170 (
float)QuantumRange,(float)QuantumScale,(
float)CLCharQuantumScale,
2171 (float)MagickEpsilon,(
float)MagickPI,(
unsigned int)MaxMap,
2172 (
unsigned int)MAGICKCORE_QUANTUM_DEPTH);
2174 signature=StringSignature(options);
2175 accelerateKernelsBuffer=(
char*) AcquireQuantumMemory(1,
2176 strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
2177 if (accelerateKernelsBuffer == (
char*) NULL)
2178 return(MagickFalse);
2179 sprintf(accelerateKernelsBuffer,
"%s%s",accelerateKernels,accelerateKernels2);
2180 signature^=StringSignature(accelerateKernelsBuffer);
2183 for (i = 0; i < clEnv->number_devices; i++)
2191 device=clEnv->devices[i];
2192 if ((device->enabled == MagickFalse) ||
2193 (device->program != (cl_program) NULL))
2196 LockSemaphoreInfo(device->lock);
2197 if (device->program != (cl_program) NULL)
2199 UnlockSemaphoreInfo(device->lock);
2202 device_signature=signature;
2203 device_signature^=StringSignature(device->platform_name);
2204 status=CompileOpenCLKernel(device,accelerateKernelsBuffer,options,
2205 device_signature,exception);
2206 UnlockSemaphoreInfo(device->lock);
2207 if (status == MagickFalse)
2210 accelerateKernelsBuffer=RelinquishMagickMemory(accelerateKernelsBuffer);
2238 static cl_uint GetOpenCLDeviceCount(MagickCLEnv clEnv,cl_platform_id platform)
2241 version[MagickPathExtent];
2246 if (clEnv->library->clGetPlatformInfo(platform,CL_PLATFORM_VERSION,
2247 MagickPathExtent,version,NULL) != CL_SUCCESS)
2249 if (strncmp(version,
"OpenCL 1.0 ",11) == 0)
2251 if (clEnv->library->clGetDeviceIDs(platform,
2252 CL_DEVICE_TYPE_CPU|CL_DEVICE_TYPE_GPU,0,NULL,&num) != CL_SUCCESS)
2257 static inline char *GetOpenCLPlatformString(cl_platform_id platform,
2258 cl_platform_info param_name)
2266 openCL_library->clGetPlatformInfo(platform,param_name,0,NULL,&length);
2267 value=AcquireCriticalMemory(length*
sizeof(*value));
2268 openCL_library->clGetPlatformInfo(platform,param_name,length,value,NULL);
2272 static inline char *GetOpenCLDeviceString(cl_device_id device,
2273 cl_device_info param_name)
2281 openCL_library->clGetDeviceInfo(device,param_name,0,NULL,&length);
2282 value=AcquireCriticalMemory(length*
sizeof(*value));
2283 openCL_library->clGetDeviceInfo(device,param_name,length,value,NULL);
2287 static void LoadOpenCLDevices(MagickCLEnv clEnv)
2289 cl_context_properties
2309 if (openCL_library->clGetPlatformIDs(0,NULL,&number_platforms) != CL_SUCCESS)
2311 if (number_platforms == 0)
2313 platforms=(cl_platform_id *) AcquireQuantumMemory(1,number_platforms*
2314 sizeof(cl_platform_id));
2315 if (platforms == (cl_platform_id *) NULL)
2317 if (openCL_library->clGetPlatformIDs(number_platforms,platforms,NULL) != CL_SUCCESS)
2319 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2322 for (i = 0; i < number_platforms; i++)
2328 platform_name=GetOpenCLPlatformString(platforms[i],CL_PLATFORM_NAME);
2330 if (strncmp(platform_name,
"NVIDIA",6) != 0)
2332 number_devices=GetOpenCLDeviceCount(clEnv,platforms[i]);
2333 clEnv->number_devices+=number_devices;
2335 platform_name=(
char *) RelinquishMagickMemory(platform_name);
2336 if (number_devices == 0)
2337 platforms[i]=(cl_platform_id) NULL;
2339 if (clEnv->number_devices == 0)
2341 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2344 clEnv->devices=(MagickCLDevice *) AcquireQuantumMemory(clEnv->number_devices,
2345 sizeof(MagickCLDevice));
2346 if (clEnv->devices == (MagickCLDevice *) NULL)
2348 RelinquishMagickCLDevices(clEnv);
2349 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2352 (void) memset(clEnv->devices,0,clEnv->number_devices*
sizeof(MagickCLDevice));
2353 devices=(cl_device_id *) AcquireQuantumMemory(clEnv->number_devices,
2354 sizeof(cl_device_id));
2355 if (devices == (cl_device_id *) NULL)
2357 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2358 RelinquishMagickCLDevices(clEnv);
2361 (void) memset(devices,0,clEnv->number_devices*
sizeof(cl_device_id));
2362 clEnv->number_contexts=(size_t) number_platforms;
2363 clEnv->contexts=(cl_context *) AcquireQuantumMemory(clEnv->number_contexts,
2364 sizeof(cl_context));
2365 if (clEnv->contexts == (cl_context *) NULL)
2367 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2368 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2369 RelinquishMagickCLDevices(clEnv);
2372 (void) memset(clEnv->contexts,0,clEnv->number_contexts*
sizeof(cl_context));
2374 for (i = 0; i < number_platforms; i++)
2376 if (platforms[i] == (cl_platform_id) NULL)
2379 status=clEnv->library->clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_CPU |
2380 CL_DEVICE_TYPE_GPU,(cl_uint) clEnv->number_devices,devices,&number_devices);
2381 if (status != CL_SUCCESS)
2384 properties[0]=CL_CONTEXT_PLATFORM;
2385 properties[1]=(cl_context_properties) platforms[i];
2387 clEnv->contexts[i]=openCL_library->clCreateContext(properties,number_devices,
2388 devices,NULL,NULL,&status);
2389 if (status != CL_SUCCESS)
2392 for (j = 0; j < number_devices; j++,next++)
2397 device=AcquireMagickCLDevice();
2398 if (device == (MagickCLDevice) NULL)
2401 device->context=clEnv->contexts[i];
2402 device->deviceID=devices[j];
2404 device->platform_name=GetOpenCLPlatformString(platforms[i],
2407 device->vendor_name=GetOpenCLPlatformString(platforms[i],
2408 CL_PLATFORM_VENDOR);
2410 device->name=GetOpenCLDeviceString(devices[j],CL_DEVICE_NAME);
2412 device->version=GetOpenCLDeviceString(devices[j],CL_DRIVER_VERSION);
2414 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_CLOCK_FREQUENCY,
2415 sizeof(cl_uint),&device->max_clock_frequency,NULL);
2417 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_MAX_COMPUTE_UNITS,
2418 sizeof(cl_uint),&device->max_compute_units,NULL);
2420 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_TYPE,
2421 sizeof(cl_device_type),&device->type,NULL);
2423 openCL_library->clGetDeviceInfo(devices[j],CL_DEVICE_LOCAL_MEM_SIZE,
2424 sizeof(cl_ulong),&device->local_memory_size,NULL);
2426 clEnv->devices[next]=device;
2427 (void) LogMagickEvent(AccelerateEvent,GetMagickModule(),
2428 "Found device: %s (%s)",device->name,device->platform_name);
2431 if (next != clEnv->number_devices)
2432 RelinquishMagickCLDevices(clEnv);
2433 platforms=(cl_platform_id *) RelinquishMagickMemory(platforms);
2434 devices=(cl_device_id *) RelinquishMagickMemory(devices);
2437 MagickPrivate MagickBooleanType InitializeOpenCL(MagickCLEnv clEnv,
2443 LockSemaphoreInfo(clEnv->lock);
2444 if (clEnv->initialized != MagickFalse)
2446 UnlockSemaphoreInfo(clEnv->lock);
2447 return(HasOpenCLDevices(clEnv,exception));
2449 if (LoadOpenCLLibrary() != MagickFalse)
2451 clEnv->library=openCL_library;
2452 LoadOpenCLDevices(clEnv);
2453 if (clEnv->number_devices > 0)
2454 AutoSelectOpenCLDevices(clEnv);
2456 clEnv->initialized=MagickTrue;
2458 for (i=0; i < (ssize_t) clEnv->number_devices; i++)
2460 if (strncmp(clEnv->devices[i]->platform_name,
"NVIDIA",6) == 0)
2461 clEnv->devices[i]->enabled=MagickFalse;
2463 UnlockSemaphoreInfo(clEnv->lock);
2464 return(HasOpenCLDevices(clEnv,exception));
2486 void *OsLibraryGetFunctionAddress(
void *library,
const char *functionName)
2488 if ((library == (
void *) NULL) || (functionName == (
const char *) NULL))
2489 return (
void *) NULL;
2490 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2491 return (
void *) GetProcAddress((HMODULE)library,functionName);
2493 return (
void *) dlsym(library,functionName);
2497 static MagickBooleanType BindOpenCLFunctions()
2499 #ifdef MAGICKCORE_HAVE_OPENCL_CL_H
2500 #define BIND(X) openCL_library->X= &X;
2502 (void) memset(openCL_library,0,
sizeof(MagickLibrary));
2503 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2504 openCL_library->library=(
void *)LoadLibraryA(
"OpenCL.dll");
2506 openCL_library->library=(
void *)dlopen(
"libOpenCL.so",RTLD_NOW);
2509 if ((openCL_library->X=(MAGICKpfn_##X)OsLibraryGetFunctionAddress(openCL_library->library,#X)) == NULL) \
2510 return(MagickFalse);
2513 if (openCL_library->library == (
void*) NULL)
2514 return(MagickFalse);
2516 BIND(clGetPlatformIDs);
2517 BIND(clGetPlatformInfo);
2519 BIND(clGetDeviceIDs);
2520 BIND(clGetDeviceInfo);
2522 BIND(clCreateBuffer);
2523 BIND(clReleaseMemObject);
2524 BIND(clRetainMemObject);
2526 BIND(clCreateContext);
2527 BIND(clReleaseContext);
2529 BIND(clCreateCommandQueue);
2530 BIND(clReleaseCommandQueue);
2534 BIND(clCreateProgramWithSource);
2535 BIND(clCreateProgramWithBinary);
2536 BIND(clReleaseProgram);
2537 BIND(clBuildProgram);
2538 BIND(clGetProgramBuildInfo);
2539 BIND(clGetProgramInfo);
2541 BIND(clCreateKernel);
2542 BIND(clReleaseKernel);
2543 BIND(clSetKernelArg);
2544 BIND(clGetKernelInfo);
2546 BIND(clEnqueueReadBuffer);
2547 BIND(clEnqueueMapBuffer);
2548 BIND(clEnqueueUnmapMemObject);
2549 BIND(clEnqueueNDRangeKernel);
2551 BIND(clGetEventInfo);
2552 BIND(clWaitForEvents);
2553 BIND(clReleaseEvent);
2554 BIND(clRetainEvent);
2555 BIND(clSetEventCallback);
2557 BIND(clGetEventProfilingInfo);
2562 static MagickBooleanType LoadOpenCLLibrary(
void)
2564 openCL_library=(MagickLibrary *) AcquireMagickMemory(
sizeof(MagickLibrary));
2565 if (openCL_library == (MagickLibrary *) NULL)
2566 return(MagickFalse);
2568 if (BindOpenCLFunctions() == MagickFalse)
2570 openCL_library=(MagickLibrary *)RelinquishMagickMemory(openCL_library);
2571 return(MagickFalse);
2596 MagickPrivate
void OpenCLTerminus()
2598 DumpOpenCLProfileData();
2599 if (cache_directory != (
char *) NULL)
2600 cache_directory=DestroyString(cache_directory);
2602 RelinquishSemaphoreInfo(&cache_directory_lock);
2603 if (default_CLEnv != (MagickCLEnv) NULL)
2604 default_CLEnv=RelinquishMagickCLEnv(default_CLEnv);
2606 RelinquishSemaphoreInfo(&openCL_lock);
2607 if (openCL_library != (MagickLibrary *) NULL)
2609 if (openCL_library->library != (
void *) NULL)
2610 (void) lt_dlclose(openCL_library->library);
2611 openCL_library=(MagickLibrary *) RelinquishMagickMemory(openCL_library);
2654 MagickPrivate MagickBooleanType OpenCLThrowMagickException(
2655 MagickCLDevice device,
ExceptionInfo *exception,
const char *module,
2656 const char *
function,
const size_t line,
const ExceptionType severity,
2657 const char *tag,
const char *format,...)
2662 assert(device != (MagickCLDevice) NULL);
2664 assert(exception->signature == MagickCoreSignature);
2669 if (device->type == CL_DEVICE_TYPE_CPU)
2673 if (strncmp(device->platform_name,
"Intel",5) == 0)
2674 default_CLEnv->enabled=MagickFalse;
2678 #ifdef OPENCLLOG_ENABLED
2682 va_start(operands,format);
2683 status=ThrowMagickExceptionList(exception,module,
function,line,severity,tag,
2688 magick_unreferenced(module);
2689 magick_unreferenced(
function);
2690 magick_unreferenced(line);
2691 magick_unreferenced(tag);
2692 magick_unreferenced(format);
2724 MagickPrivate MagickBooleanType RecordProfileData(MagickCLDevice device,
2725 cl_kernel kernel,cl_event event)
2745 if (device->profile_kernels == MagickFalse)
2746 return(MagickFalse);
2747 status=openCL_library->clWaitForEvents(1,&event);
2748 if (status != CL_SUCCESS)
2749 return(MagickFalse);
2750 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,0,NULL,
2752 if (status != CL_SUCCESS)
2754 name=AcquireQuantumMemory(length,
sizeof(*name));
2755 if (name == (
char *) NULL)
2757 start=end=elapsed=0;
2758 status=openCL_library->clGetKernelInfo(kernel,CL_KERNEL_FUNCTION_NAME,length,
2759 name,(
size_t *) NULL);
2760 status|=openCL_library->clGetEventProfilingInfo(event,
2761 CL_PROFILING_COMMAND_START,
sizeof(cl_ulong),&start,NULL);
2762 status|=openCL_library->clGetEventProfilingInfo(event,
2763 CL_PROFILING_COMMAND_END,
sizeof(cl_ulong),&end,NULL);
2764 if (status != CL_SUCCESS)
2766 name=DestroyString(name);
2772 LockSemaphoreInfo(device->lock);
2779 if (LocaleCompare(device->profile_records[i]->kernel_name,name) == 0)
2781 profile_record=device->profile_records[i];
2788 name=DestroyString(name);
2791 profile_record=AcquireCriticalMemory(
sizeof(*profile_record));
2792 (void) memset(profile_record,0,
sizeof(*profile_record));
2793 profile_record->kernel_name=name;
2794 device->profile_records=ResizeQuantumMemory(device->profile_records,(i+2),
2795 sizeof(*device->profile_records));
2797 ThrowFatalException(ResourceLimitFatalError,
"MemoryAllocationFailed");
2798 device->profile_records[i]=profile_record;
2801 if ((elapsed < profile_record->min) || (profile_record->count == 0))
2802 profile_record->min=elapsed;
2803 if (elapsed > profile_record->max)
2804 profile_record->max=elapsed;
2805 profile_record->total+=elapsed;
2806 profile_record->count+=1;
2807 UnlockSemaphoreInfo(device->lock);
2836 MagickPrivate
void ReleaseOpenCLCommandQueue(MagickCLDevice device,
2837 cl_command_queue queue)
2839 if (queue == (cl_command_queue) NULL)
2842 assert(device != (MagickCLDevice) NULL);
2843 LockSemaphoreInfo(device->lock);
2844 if ((device->profile_kernels != MagickFalse) ||
2845 (device->command_queues_index >= MAGICKCORE_OPENCL_COMMAND_QUEUES-1))
2847 UnlockSemaphoreInfo(device->lock);
2848 openCL_library->clFinish(queue);
2849 (void) openCL_library->clReleaseCommandQueue(queue);
2853 openCL_library->clFlush(queue);
2854 device->command_queues[++device->command_queues_index]=queue;
2855 UnlockSemaphoreInfo(device->lock);
2882 MagickPrivate
void ReleaseOpenCLDevice(MagickCLDevice device)
2884 assert(device != (MagickCLDevice) NULL);
2885 LockSemaphoreInfo(openCL_lock);
2886 device->requested--;
2887 UnlockSemaphoreInfo(openCL_lock);
2917 static void CL_API_CALL DestroyMagickCLCacheInfoAndPixels(
2918 cl_event magick_unused(event),
2919 cl_int magick_unused(event_command_exec_status),
void *user_data)
2930 magick_unreferenced(event);
2931 magick_unreferenced(event_command_exec_status);
2932 info=(MagickCLCacheInfo) user_data;
2933 for (i=(ssize_t)info->event_count-1; i >= 0; i--)
2941 status=openCL_library->clGetEventInfo(info->events[i],
2942 CL_EVENT_COMMAND_EXECUTION_STATUS,
sizeof(event_status),&event_status,
2944 if ((status == CL_SUCCESS) && (event_status > CL_COMPLETE))
2946 openCL_library->clSetEventCallback(info->events[i],CL_COMPLETE,
2947 &DestroyMagickCLCacheInfoAndPixels,info);
2951 pixels=info->pixels;
2952 RelinquishMagickResource(MemoryResource,info->length);
2953 DestroyMagickCLCacheInfo(info);
2954 (void) RelinquishAlignedMemory(pixels);
2957 MagickPrivate MagickCLCacheInfo RelinquishMagickCLCacheInfo(
2958 MagickCLCacheInfo info,
const MagickBooleanType relinquish_pixels)
2960 if (info == (MagickCLCacheInfo) NULL)
2961 return((MagickCLCacheInfo) NULL);
2962 if (relinquish_pixels != MagickFalse)
2963 DestroyMagickCLCacheInfoAndPixels((cl_event) NULL,0,info);
2965 DestroyMagickCLCacheInfo(info);
2966 return((MagickCLCacheInfo) NULL);
2992 static MagickCLDevice RelinquishMagickCLDevice(MagickCLDevice device)
2994 if (device == (MagickCLDevice) NULL)
2995 return((MagickCLDevice) NULL);
2997 device->platform_name=RelinquishMagickMemory(device->platform_name);
2998 device->vendor_name=RelinquishMagickMemory(device->vendor_name);
2999 device->name=RelinquishMagickMemory(device->name);
3000 device->version=RelinquishMagickMemory(device->version);
3001 if (device->program != (cl_program) NULL)
3002 (void) openCL_library->clReleaseProgram(device->program);
3003 while (device->command_queues_index >= 0)
3004 (void) openCL_library->clReleaseCommandQueue(
3005 device->command_queues[device->command_queues_index--]);
3006 RelinquishSemaphoreInfo(&device->lock);
3007 return((MagickCLDevice) RelinquishMagickMemory(device));
3033 static MagickCLEnv RelinquishMagickCLEnv(MagickCLEnv clEnv)
3035 if (clEnv == (MagickCLEnv) NULL)
3036 return((MagickCLEnv) NULL);
3038 RelinquishSemaphoreInfo(&clEnv->lock);
3039 RelinquishMagickCLDevices(clEnv);
3040 if (clEnv->contexts != (cl_context *) NULL)
3045 for (i=0; i < clEnv->number_contexts; i++)
3046 if (clEnv->contexts[i] != (cl_context) NULL)
3047 (void) openCL_library->clReleaseContext(clEnv->contexts[i]);
3048 clEnv->contexts=(cl_context *) RelinquishMagickMemory(clEnv->contexts);
3050 return((MagickCLEnv) RelinquishMagickMemory(clEnv));
3075 MagickPrivate MagickCLDevice RequestOpenCLDevice(MagickCLEnv clEnv)
3087 if (clEnv == (MagickCLEnv) NULL)
3088 return((MagickCLDevice) NULL);
3090 if (clEnv->number_devices == 1)
3092 if (clEnv->devices[0]->enabled)
3093 return(clEnv->devices[0]);
3095 return((MagickCLDevice) NULL);
3098 device=(MagickCLDevice) NULL;
3100 LockSemaphoreInfo(openCL_lock);
3101 for (i = 0; i < clEnv->number_devices; i++)
3103 if (clEnv->devices[i]->enabled == MagickFalse)
3106 score=clEnv->devices[i]->score+(clEnv->devices[i]->score*
3107 clEnv->devices[i]->requested);
3108 if ((device == (MagickCLDevice) NULL) || (score < best_score))
3110 device=clEnv->devices[i];
3114 if (device != (MagickCLDevice)NULL)
3115 device->requested++;
3116 UnlockSemaphoreInfo(openCL_lock);
3146 MagickExport
void SetOpenCLDeviceEnabled(MagickCLDevice device,
3147 const MagickBooleanType value)
3149 if (device == (MagickCLDevice) NULL)
3151 device->enabled=value;
3181 MagickExport
void SetOpenCLKernelProfileEnabled(MagickCLDevice device,
3182 const MagickBooleanType value)
3184 if (device == (MagickCLDevice) NULL)
3186 device->profile_kernels=value;
3211 MagickExport MagickBooleanType SetOpenCLEnabled(
const MagickBooleanType value)
3216 clEnv=GetCurrentOpenCLEnv();
3217 if (clEnv == (MagickCLEnv) NULL)
3218 return(MagickFalse);
3219 clEnv->enabled=value;
3220 return(clEnv->enabled);
3225 MagickExport
double GetOpenCLDeviceBenchmarkScore(
3226 const MagickCLDevice magick_unused(device))
3228 magick_unreferenced(device);
3232 MagickExport MagickBooleanType GetOpenCLDeviceEnabled(
3233 const MagickCLDevice magick_unused(device))
3235 magick_unreferenced(device);
3236 return(MagickFalse);
3239 MagickExport
const char *GetOpenCLDeviceName(
3240 const MagickCLDevice magick_unused(device))
3242 magick_unreferenced(device);
3243 return((
const char *) NULL);
3246 MagickExport MagickCLDevice *GetOpenCLDevices(
size_t *length,
3249 magick_unreferenced(exception);
3250 if (length != (
size_t *) NULL)
3252 return((MagickCLDevice *) NULL);
3255 MagickExport MagickCLDeviceType GetOpenCLDeviceType(
3256 const MagickCLDevice magick_unused(device))
3258 magick_unreferenced(device);
3259 return(UndefinedCLDeviceType);
3263 const MagickCLDevice magick_unused(device),
size_t *length)
3265 magick_unreferenced(device);
3266 if (length != (
size_t *) NULL)
3271 MagickExport
const char *GetOpenCLDeviceVersion(
3272 const MagickCLDevice magick_unused(device))
3274 magick_unreferenced(device);
3275 return((
const char *) NULL);
3278 MagickExport MagickBooleanType GetOpenCLEnabled(
void)
3280 return(MagickFalse);
3283 MagickExport
void SetOpenCLDeviceEnabled(
3284 MagickCLDevice magick_unused(device),
3285 const MagickBooleanType magick_unused(value))
3287 magick_unreferenced(device);
3288 magick_unreferenced(value);
3291 MagickExport MagickBooleanType SetOpenCLEnabled(
3292 const MagickBooleanType magick_unused(value))
3294 magick_unreferenced(value);
3295 return(MagickFalse);
3298 MagickExport
void SetOpenCLKernelProfileEnabled(
3299 MagickCLDevice magick_unused(device),
3300 const MagickBooleanType magick_unused(value))
3302 magick_unreferenced(device);
3303 magick_unreferenced(value);