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