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