MagickCore  6.9.13-46
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/license/ %
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  memset(binary_program,0,binary_program_size);
914  for (i = 0; i < num_devices; i++)
915  {
916  binary_program[i]=(unsigned char *) AcquireQuantumMemory(
917  MagickMax(*(program_sizes+i),1),sizeof(**binary_program));
918  if (binary_program[i] == (unsigned char *) NULL)
919  {
920  status=CL_OUT_OF_HOST_MEMORY;
921  break;
922  }
923  }
924  if (status == CL_SUCCESS)
925  status=clEnv->library->clGetProgramInfo(clEnv->programs[prog],
926  CL_PROGRAM_BINARIES,binary_program_size,binary_program,NULL);
927  if (status == CL_SUCCESS)
928  {
929  for (i = 0; i < num_devices; i++)
930  {
931  int
932  file;
933 
934  size_t
935  program_size;
936 
937  program_size=*(program_sizes+i);
938  if (program_size < 1)
939  continue;
940  file=open_utf8(filename,O_WRONLY | O_CREAT | O_BINARY,S_MODE);
941  if (file != -1)
942  {
943  write(file,binary_program[i],program_size);
944  file=close_utf8(file);
945  }
946  else
947  (void) ThrowMagickException(exception,GetMagickModule(),
948  DelegateWarning,"Saving kernel failed.","`%s'",filename);
949  break;
950  }
951  }
952  for (i = 0; i < num_devices; i++)
953  binary_program[i]=(unsigned char *) RelinquishMagickMemory(
954  binary_program[i]);
955  binary_program=(unsigned char **) RelinquishMagickMemory(binary_program);
956  }
957  program_sizes=(size_t *) RelinquishMagickMemory(program_sizes);
958 }
959 
960 static MagickBooleanType loadBinaryCLProgram(MagickCLEnv clEnv, MagickOpenCLProgram prog, unsigned int signature)
961 {
962  MagickBooleanType loadSuccessful;
963  unsigned char* binaryProgram;
964  char* binaryFileName;
965  FILE* fileHandle;
966 
967 #ifdef MAGICKCORE_CLPERFMARKER
968  clBeginPerfMarkerAMD(__FUNCTION__,"");
969 #endif
970 
971  binaryProgram = NULL;
972  binaryFileName = NULL;
973  fileHandle = NULL;
974  loadSuccessful = MagickFalse;
975 
976  binaryFileName = getBinaryCLProgramName(clEnv, prog, signature);
977  fileHandle = fopen(binaryFileName, "rb");
978  if (fileHandle != NULL)
979  {
980  int b_error;
981  size_t length;
982  cl_int clStatus;
983  cl_int clBinaryStatus;
984 
985  b_error = 0 ;
986  length = 0;
987  b_error |= fseek( fileHandle, 0, SEEK_END ) < 0;
988  b_error |= ( length = ftell( fileHandle ) ) <= 0;
989  b_error |= fseek( fileHandle, 0, SEEK_SET ) < 0;
990  if( b_error )
991  goto cleanup;
992 
993  binaryProgram = (unsigned char*)AcquireMagickMemory(length);
994  if (binaryProgram == NULL)
995  goto cleanup;
996 
997  memset(binaryProgram, 0, length);
998  b_error |= fread(binaryProgram, 1, length, fileHandle) != length;
999 
1000  clEnv->programs[prog] = clEnv->library->clCreateProgramWithBinary(clEnv->context, 1, &clEnv->device, &length, (const unsigned char**)&binaryProgram, &clBinaryStatus, &clStatus);
1001  if (clStatus != CL_SUCCESS
1002  || clBinaryStatus != CL_SUCCESS)
1003  goto cleanup;
1004 
1005  loadSuccessful = MagickTrue;
1006  }
1007 
1008 cleanup:
1009  if (fileHandle != NULL)
1010  fclose(fileHandle);
1011  if (binaryFileName != NULL)
1012  RelinquishMagickMemory(binaryFileName);
1013  if (binaryProgram != NULL)
1014  RelinquishMagickMemory(binaryProgram);
1015 
1016 #ifdef MAGICKCORE_CLPERFMARKER
1017  clEndPerfMarkerAMD();
1018 #endif
1019 
1020  return loadSuccessful;
1021 }
1022 
1023 static unsigned int stringSignature(const char* string)
1024 {
1025  unsigned int stringLength;
1026  unsigned int n,i,j;
1027  unsigned int signature;
1028  union
1029  {
1030  const char* s;
1031  const unsigned int* u;
1032  }p;
1033 
1034 #ifdef MAGICKCORE_CLPERFMARKER
1035  clBeginPerfMarkerAMD(__FUNCTION__,"");
1036 #endif
1037 
1038  stringLength = (unsigned int) strlen(string);
1039  signature = stringLength;
1040  n = stringLength/sizeof(unsigned int);
1041  p.s = string;
1042  for (i = 0; i < n; i++)
1043  {
1044  signature^=p.u[i];
1045  }
1046  if (n * sizeof(unsigned int) != stringLength)
1047  {
1048  char padded[4];
1049  j = n * sizeof(unsigned int);
1050  for (i = 0; i < 4; i++,j++)
1051  {
1052  if (j < stringLength)
1053  padded[i] = p.s[j];
1054  else
1055  padded[i] = 0;
1056  }
1057  p.s = padded;
1058  signature^=p.u[0];
1059  }
1060 
1061 #ifdef MAGICKCORE_CLPERFMARKER
1062  clEndPerfMarkerAMD();
1063 #endif
1064 
1065  return signature;
1066 }
1067 
1068 /* OpenCL kernels for accelerate.c */
1069 extern const char *accelerateKernels, *accelerateKernels2;
1070 
1071 static MagickBooleanType CompileOpenCLKernels(MagickCLEnv clEnv, ExceptionInfo* exception)
1072 {
1073  MagickBooleanType status = MagickFalse;
1074  cl_int clStatus;
1075  unsigned int i;
1076  char* accelerateKernelsBuffer = NULL;
1077 
1078  /* The index of the program strings in this array has to match the value of the enum MagickOpenCLProgram */
1079  const char* MagickOpenCLProgramStrings[MAGICK_OPENCL_NUM_PROGRAMS];
1080 
1081  char options[MaxTextExtent];
1082  unsigned int optionsSignature;
1083 
1084 #ifdef MAGICKCORE_CLPERFMARKER
1085  clBeginPerfMarkerAMD(__FUNCTION__,"");
1086 #endif
1087 
1088  /* Get additional options */
1089  (void) FormatLocaleString(options, MaxTextExtent, CLOptions, (float)QuantumRange,
1090  (float)QuantumScale, (float)CLCharQuantumScale, (float)MagickEpsilon, (float)MagickPI, (unsigned int)MaxMap, (unsigned int)MAGICKCORE_QUANTUM_DEPTH);
1091 
1092  /*
1093  if (getenv("MAGICK_OCL_DEF"))
1094  {
1095  strcat(options," ");
1096  strcat(options,getenv("MAGICK_OCL_DEF"));
1097  }
1098  */
1099 
1100  /*
1101  if (getenv("MAGICK_OCL_BUILD"))
1102  printf("options: %s\n", options);
1103  */
1104 
1105  optionsSignature = stringSignature(options);
1106 
1107  /* get all the OpenCL program strings here */
1108  accelerateKernelsBuffer = (char*) AcquireQuantumMemory(1,strlen(accelerateKernels)+strlen(accelerateKernels2)+1);
1109  FormatLocaleString(accelerateKernelsBuffer,strlen(accelerateKernels)+
1110  strlen(accelerateKernels2)+1,"%s%s",accelerateKernels,accelerateKernels2);
1111  MagickOpenCLProgramStrings[MAGICK_OPENCL_ACCELERATE] = accelerateKernelsBuffer;
1112 
1113  for (i = 0; i < MAGICK_OPENCL_NUM_PROGRAMS; i++)
1114  {
1115  MagickBooleanType loadSuccessful = MagickFalse;
1116  unsigned int programSignature = stringSignature(MagickOpenCLProgramStrings[i]) ^ optionsSignature;
1117 
1118  /* try to load the binary first */
1119  if (clEnv->disableProgramCache != MagickTrue
1120  && !getenv("MAGICK_OCL_REC"))
1121  loadSuccessful = loadBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature);
1122 
1123  if (loadSuccessful == MagickFalse)
1124  {
1125  /* Binary CL program unavailable, compile the program from source */
1126  size_t programLength = strlen(MagickOpenCLProgramStrings[i]);
1127  clEnv->programs[i] = clEnv->library->clCreateProgramWithSource(clEnv->context, 1, &(MagickOpenCLProgramStrings[i]), &programLength, &clStatus);
1128  if (clStatus!=CL_SUCCESS)
1129  {
1130  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1131  "clCreateProgramWithSource failed.", "(%d)", (int)clStatus);
1132 
1133  goto cleanup;
1134  }
1135  }
1136 
1137  clStatus = clEnv->library->clBuildProgram(clEnv->programs[i], 1, &clEnv->device, options, NULL, NULL);
1138  if (clStatus!=CL_SUCCESS)
1139  {
1140  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1141  "clBuildProgram failed.", "(%d)", (int)clStatus);
1142 
1143  if (loadSuccessful == MagickFalse)
1144  {
1145  char path[MaxTextExtent];
1146  FILE* fileHandle;
1147 
1148  /* dump the source into a file */
1149  (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
1150  ,GetOpenCLCachedFilesDirectory()
1151  ,DirectorySeparator,"magick_badcl.cl");
1152  fileHandle = fopen(path, "wb");
1153  if (fileHandle != NULL)
1154  {
1155  fwrite(MagickOpenCLProgramStrings[i], sizeof(char), strlen(MagickOpenCLProgramStrings[i]), fileHandle);
1156  fclose(fileHandle);
1157  }
1158 
1159  /* dump the build log */
1160  {
1161  char* log;
1162  size_t logSize;
1163  clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
1164  log = (char*)AcquireCriticalMemory(logSize);
1165  clEnv->library->clGetProgramBuildInfo(clEnv->programs[i], clEnv->device, CL_PROGRAM_BUILD_LOG, logSize, log, &logSize);
1166 
1167  (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
1168  ,GetOpenCLCachedFilesDirectory()
1169  ,DirectorySeparator,"magick_badcl_build.log");
1170  fileHandle = fopen(path, "wb");
1171  if (fileHandle != NULL)
1172  {
1173  const char* buildOptionsTitle = "build options: ";
1174  fwrite(buildOptionsTitle, sizeof(char), strlen(buildOptionsTitle), fileHandle);
1175  fwrite(options, sizeof(char), strlen(options), fileHandle);
1176  fwrite("\n",sizeof(char), 1, fileHandle);
1177  fwrite(log, sizeof(char), logSize, fileHandle);
1178  fclose(fileHandle);
1179  }
1180  RelinquishMagickMemory(log);
1181  }
1182  }
1183  goto cleanup;
1184  }
1185 
1186  if (loadSuccessful == MagickFalse)
1187  {
1188  /* Save the binary to a file to avoid re-compilation of the kernels in the future */
1189  saveBinaryCLProgram(clEnv, (MagickOpenCLProgram)i, programSignature, exception);
1190  }
1191 
1192  }
1193  status = MagickTrue;
1194 
1195 cleanup:
1196 
1197  if (accelerateKernelsBuffer!=NULL) RelinquishMagickMemory(accelerateKernelsBuffer);
1198 
1199 #ifdef MAGICKCORE_CLPERFMARKER
1200  clEndPerfMarkerAMD();
1201 #endif
1202 
1203  return status;
1204 }
1205 
1206 static MagickBooleanType InitOpenCLPlatformDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
1207  int i,j;
1208  cl_int status;
1209  cl_uint numPlatforms = 0;
1210  cl_platform_id *platforms = NULL;
1211  char* MAGICK_OCL_DEVICE = NULL;
1212  MagickBooleanType OpenCLAvailable = MagickFalse;
1213 
1214 #ifdef MAGICKCORE_CLPERFMARKER
1215  clBeginPerfMarkerAMD(__FUNCTION__,"");
1216 #endif
1217 
1218  /* check if there's an environment variable overriding the device selection */
1219  MAGICK_OCL_DEVICE = getenv("MAGICK_OCL_DEVICE");
1220  if (MAGICK_OCL_DEVICE == (char *) NULL)
1221  return(MagickFalse);
1222  if (strcmp(MAGICK_OCL_DEVICE, "CPU") == 0)
1223  clEnv->deviceType = CL_DEVICE_TYPE_CPU;
1224  else if (strcmp(MAGICK_OCL_DEVICE, "GPU") == 0)
1225  clEnv->deviceType = CL_DEVICE_TYPE_GPU;
1226  else if (IsStringTrue(MAGICK_OCL_DEVICE) != MagickFalse)
1227  {
1228  if (clEnv->deviceType == 0)
1229  clEnv->deviceType = CL_DEVICE_TYPE_ALL;
1230  }
1231  else
1232  return(MagickFalse);
1233 
1234  if (clEnv->device != NULL)
1235  {
1236  status = clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &clEnv->platform, NULL);
1237  if (status != CL_SUCCESS) {
1238  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1239  "Failed to get OpenCL platform from the selected device.", "(%d)", status);
1240  }
1241  goto cleanup;
1242  }
1243  else if (clEnv->platform != NULL)
1244  {
1245  numPlatforms = 1;
1246  platforms = (cl_platform_id *) AcquireQuantumMemory(1,numPlatforms * sizeof(cl_platform_id));
1247  if (platforms == (cl_platform_id *) NULL)
1248  {
1249  (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1250  "AcquireMagickMemory failed.",".");
1251  goto cleanup;
1252  }
1253  platforms[0] = clEnv->platform;
1254  }
1255  else
1256  {
1257  clEnv->device = NULL;
1258 
1259  /* Get the number of OpenCL platforms available */
1260  status = clEnv->library->clGetPlatformIDs(0, NULL, &numPlatforms);
1261  if (status != CL_SUCCESS)
1262  {
1263  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1264  "clGetplatformIDs failed.", "(%d)", status);
1265  goto cleanup;
1266  }
1267 
1268  /* No OpenCL available, just leave */
1269  if (numPlatforms == 0) {
1270  goto cleanup;
1271  }
1272 
1273  platforms = (cl_platform_id *) AcquireQuantumMemory(1,numPlatforms * sizeof(cl_platform_id));
1274  if (platforms == (cl_platform_id *) NULL)
1275  {
1276  (void) ThrowMagickException(exception, GetMagickModule(), ResourceLimitError,
1277  "AcquireMagickMemory failed.",".");
1278  goto cleanup;
1279  }
1280 
1281  status = clEnv->library->clGetPlatformIDs(numPlatforms, platforms, NULL);
1282  if (status != CL_SUCCESS)
1283  {
1284  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1285  "clGetPlatformIDs failed.", "(%d)", status);
1286  goto cleanup;
1287  }
1288  }
1289 
1290  /* Device selection */
1291  clEnv->device = NULL;
1292  for (j = 0; j < 2; j++)
1293  {
1294 
1295  cl_device_type deviceType;
1296  if (clEnv->deviceType == CL_DEVICE_TYPE_ALL)
1297  {
1298  if (j == 0)
1299  deviceType = CL_DEVICE_TYPE_GPU;
1300  else
1301  deviceType = CL_DEVICE_TYPE_CPU;
1302  }
1303  else if (j == 1)
1304  {
1305  break;
1306  }
1307  else
1308  deviceType = clEnv->deviceType;
1309 
1310  for (i = 0; i < numPlatforms; i++)
1311  {
1312  char version[MaxTextExtent];
1313  cl_uint numDevices;
1314  status = clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_VERSION, MaxTextExtent, version, NULL);
1315  if (status != CL_SUCCESS)
1316  {
1317  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1318  "clGetPlatformInfo failed.", "(%d)", status);
1319  goto cleanup;
1320  }
1321  if (strncmp(version,"OpenCL 1.0 ",11) == 0)
1322  continue;
1323  status = clEnv->library->clGetDeviceIDs(platforms[i], deviceType, 1, &(clEnv->device), &numDevices);
1324  if (status != CL_SUCCESS)
1325  {
1326  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1327  "clGetDeviceIDs failed.", "(%d)", status);
1328  goto cleanup;
1329  }
1330  if (clEnv->device != NULL)
1331  {
1332  clEnv->platform = platforms[i];
1333  goto cleanup;
1334  }
1335  }
1336  }
1337 
1338 cleanup:
1339  if (platforms!=NULL)
1340  RelinquishMagickMemory(platforms);
1341 
1342  OpenCLAvailable = (clEnv->platform!=NULL
1343  && clEnv->device!=NULL)?MagickTrue:MagickFalse;
1344 
1345 #ifdef MAGICKCORE_CLPERFMARKER
1346  clEndPerfMarkerAMD();
1347 #endif
1348 
1349  return OpenCLAvailable;
1350 }
1351 
1352 static MagickBooleanType EnableOpenCLInternal(MagickCLEnv clEnv) {
1353  if (clEnv->OpenCLInitialized != MagickFalse
1354  && clEnv->platform != NULL
1355  && clEnv->device != NULL) {
1356  clEnv->OpenCLDisabled = MagickFalse;
1357  return MagickTrue;
1358  }
1359  clEnv->OpenCLDisabled = MagickTrue;
1360  return MagickFalse;
1361 }
1362 
1363 
1364 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception);
1365 /*
1366 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1367 % %
1368 % %
1369 % %
1370 + I n i t O p e n C L E n v %
1371 % %
1372 % %
1373 % %
1374 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1375 %
1376 % InitOpenCLEnv() initialize the OpenCL environment
1377 %
1378 % The format of the RelinquishMagickOpenCLEnv method is:
1379 %
1380 % MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception)
1381 %
1382 % A description of each parameter follows:
1383 %
1384 % o clEnv: OpenCL environment structure
1385 %
1386 % o exception: return any errors or warnings.
1387 %
1388 */
1389 
1390 static void RelinquishCommandQueues(MagickCLEnv clEnv)
1391 {
1392  if (clEnv == (MagickCLEnv) NULL)
1393  return;
1394 
1395  LockSemaphoreInfo(clEnv->commandQueuesLock);
1396  while (clEnv->commandQueuesPos >= 0)
1397  clEnv->library->clReleaseCommandQueue(
1398  clEnv->commandQueues[clEnv->commandQueuesPos--]);
1399  UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1400 }
1401 
1402 MagickExport
1403 MagickBooleanType InitOpenCLEnvInternal(MagickCLEnv clEnv, ExceptionInfo* exception) {
1404  MagickBooleanType status = MagickTrue;
1405  cl_int clStatus;
1406  cl_context_properties cps[3];
1407 
1408 #ifdef MAGICKCORE_CLPERFMARKER
1409  {
1410  int status = clInitializePerfMarkerAMD();
1411  if (status == AP_SUCCESS) {
1412  /* printf("PerfMarker successfully initialized\n"); */
1413  }
1414  }
1415 #endif
1416  clEnv->OpenCLInitialized = MagickTrue;
1417 
1418  /* check and init the global lib */
1419  OpenCLLib=GetOpenCLLib();
1420  if (OpenCLLib)
1421  {
1422  clEnv->library=OpenCLLib;
1423  }
1424  else
1425  {
1426  /* turn off opencl */
1427  MagickBooleanType flag;
1428  flag = MagickTrue;
1429  SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
1430  , sizeof(MagickBooleanType), &flag, exception);
1431  }
1432 
1433  if (clEnv->OpenCLDisabled != MagickFalse)
1434  goto cleanup;
1435 
1436  clEnv->OpenCLDisabled = MagickTrue;
1437  /* setup the OpenCL platform and device */
1438  status = InitOpenCLPlatformDevice(clEnv, exception);
1439  if (status == MagickFalse) {
1440  /* No OpenCL device available */
1441  goto cleanup;
1442  }
1443 
1444  /* create an OpenCL context */
1445  cps[0] = CL_CONTEXT_PLATFORM;
1446  cps[1] = (cl_context_properties)clEnv->platform;
1447  cps[2] = 0;
1448  clEnv->context = clEnv->library->clCreateContext(cps, 1, &(clEnv->device), NULL, NULL, &clStatus);
1449  if (clStatus != CL_SUCCESS)
1450  {
1451  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1452  "clCreateContext failed.", "(%d)", clStatus);
1453  status = MagickFalse;
1454  goto cleanup;
1455  }
1456 
1457  RelinquishCommandQueues(clEnv);
1458 
1459  status = CompileOpenCLKernels(clEnv, exception);
1460  if (status == MagickFalse) {
1461  (void) ThrowMagickException(exception, GetMagickModule(), DelegateWarning,
1462  "clCreateCommandQueue failed.", "(%d)", status);
1463 
1464  goto cleanup;
1465  }
1466 
1467  status = EnableOpenCLInternal(clEnv);
1468 
1469 cleanup:
1470  return status;
1471 }
1472 
1473 
1474 MagickExport
1475 MagickBooleanType InitOpenCLEnv(MagickCLEnv clEnv, ExceptionInfo* exception) {
1476  MagickBooleanType status = MagickFalse;
1477 
1478  if ((clEnv == NULL) || (getenv("MAGICK_OCL_DEVICE") == (const char *) NULL))
1479  return MagickFalse;
1480 
1481 #ifdef MAGICKCORE_CLPERFMARKER
1482  clBeginPerfMarkerAMD(__FUNCTION__,"");
1483 #endif
1484 
1485  LockSemaphoreInfo(clEnv->lock);
1486  if (clEnv->OpenCLInitialized == MagickFalse) {
1487  if (clEnv->device==NULL && clEnv->OpenCLDisabled == MagickFalse)
1488  status = autoSelectDevice(clEnv, exception);
1489  else
1490  status = InitOpenCLEnvInternal(clEnv, exception);
1491  }
1492  UnlockSemaphoreInfo(clEnv->lock);
1493 
1494 #ifdef MAGICKCORE_CLPERFMARKER
1495  clEndPerfMarkerAMD();
1496 #endif
1497  return status;
1498 }
1499 
1500 
1501 /*
1502 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1503 % %
1504 % %
1505 % %
1506 + 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 %
1507 % %
1508 % %
1509 % %
1510 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1511 %
1512 % AcquireOpenCLCommandQueue() acquires an OpenCL command queue
1513 %
1514 % The format of the AcquireOpenCLCommandQueue method is:
1515 %
1516 % cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1517 %
1518 % A description of each parameter follows:
1519 %
1520 % o clEnv: the OpenCL environment.
1521 %
1522 */
1523 
1524 MagickPrivate cl_command_queue AcquireOpenCLCommandQueue(MagickCLEnv clEnv)
1525 {
1526  cl_command_queue
1527  queue;
1528 
1529  cl_command_queue_properties
1530  properties;
1531 
1532  if (clEnv == (MagickCLEnv) NULL)
1533  return (cl_command_queue) NULL;
1534  LockSemaphoreInfo(clEnv->commandQueuesLock);
1535  if (clEnv->commandQueuesPos >= 0) {
1536  queue=clEnv->commandQueues[clEnv->commandQueuesPos--];
1537  UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1538  }
1539  else {
1540  UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1541  properties=0;
1542 #if PROFILE_OCL_KERNELS
1543  properties=CL_QUEUE_PROFILING_ENABLE;
1544 #endif
1545  queue=clEnv->library->clCreateCommandQueue(clEnv->context,clEnv->device,
1546  properties,NULL);
1547  }
1548  return(queue);
1549 }
1550 
1551 /*
1552 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1553 % %
1554 % %
1555 % %
1556 + 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 %
1557 % %
1558 % %
1559 % %
1560 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1561 %
1562 % RelinquishOpenCLCommandQueue() releases the OpenCL command queue
1563 %
1564 % The format of the RelinquishOpenCLCommandQueue method is:
1565 %
1566 % MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
1567 % cl_command_queue queue)
1568 %
1569 % A description of each parameter follows:
1570 %
1571 % o clEnv: the OpenCL environment.
1572 %
1573 % o queue: the OpenCL queue to be released.
1574 %
1575 %
1576 */
1577 
1578 MagickPrivate MagickBooleanType RelinquishOpenCLCommandQueue(MagickCLEnv clEnv,
1579  cl_command_queue queue)
1580 {
1581  MagickBooleanType
1582  status;
1583 
1584  if (clEnv == NULL)
1585  return(MagickFalse);
1586 
1587  LockSemaphoreInfo(clEnv->commandQueuesLock);
1588 
1589  if (clEnv->commandQueuesPos >= MAX_COMMAND_QUEUES-1)
1590  {
1591  clEnv->library->clFinish(queue);
1592  status=(clEnv->library->clReleaseCommandQueue(queue) == CL_SUCCESS) ?
1593  MagickTrue : MagickFalse;
1594  }
1595  else
1596  {
1597  clEnv->library->clFlush(queue);
1598  clEnv->commandQueues[++clEnv->commandQueuesPos]=queue;
1599  status=MagickTrue;
1600  }
1601 
1602  UnlockSemaphoreInfo(clEnv->commandQueuesLock);
1603 
1604  return(status);
1605 }
1606 
1607 /*
1608 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1609 % %
1610 % %
1611 % %
1612 + A c q u i r e O p e n C L K e r n e l %
1613 % %
1614 % %
1615 % %
1616 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1617 %
1618 % AcquireOpenCLKernel() acquires an OpenCL kernel
1619 %
1620 % The format of the AcquireOpenCLKernel method is:
1621 %
1622 % cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv,
1623 % MagickOpenCLProgram program, const char* kernelName)
1624 %
1625 % A description of each parameter follows:
1626 %
1627 % o clEnv: the OpenCL environment.
1628 %
1629 % o program: the OpenCL program module that the kernel belongs to.
1630 %
1631 % o kernelName: the name of the kernel
1632 %
1633 */
1634 
1635 MagickPrivate
1636  cl_kernel AcquireOpenCLKernel(MagickCLEnv clEnv, MagickOpenCLProgram program, const char* kernelName)
1637 {
1638  cl_int clStatus;
1639  cl_kernel kernel = NULL;
1640  if (clEnv != NULL && kernelName!=NULL)
1641  {
1642  kernel = clEnv->library->clCreateKernel(clEnv->programs[program], kernelName, &clStatus);
1643  }
1644  return kernel;
1645 }
1646 
1647 
1648 /*
1649 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1650 % %
1651 % %
1652 % %
1653 + R e l i n q u i s h O p e n C L K e r n e l %
1654 % %
1655 % %
1656 % %
1657 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1658 %
1659 % RelinquishOpenCLKernel() releases an OpenCL kernel
1660 %
1661 % The format of the RelinquishOpenCLKernel method is:
1662 %
1663 % MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv,
1664 % cl_kernel kernel)
1665 %
1666 % A description of each parameter follows:
1667 %
1668 % o clEnv: the OpenCL environment.
1669 %
1670 % o kernel: the OpenCL kernel object to be released.
1671 %
1672 %
1673 */
1674 
1675 MagickPrivate
1676  MagickBooleanType RelinquishOpenCLKernel(MagickCLEnv clEnv, cl_kernel kernel)
1677 {
1678  MagickBooleanType status = MagickFalse;
1679  if (clEnv != NULL && kernel != NULL)
1680  {
1681  status = ((clEnv->library->clReleaseKernel(kernel) == CL_SUCCESS)?MagickTrue:MagickFalse);
1682  }
1683  return status;
1684 }
1685 
1686 /*
1687 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1688 % %
1689 % %
1690 % %
1691 + 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 %
1692 % %
1693 % %
1694 % %
1695 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
1696 %
1697 % GetOpenCLDeviceLocalMemorySize() returns local memory size of the device
1698 %
1699 % The format of the GetOpenCLDeviceLocalMemorySize method is:
1700 %
1701 % unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1702 %
1703 % A description of each parameter follows:
1704 %
1705 % o clEnv: the OpenCL environment.
1706 %
1707 %
1708 */
1709 
1710 MagickPrivate
1711  unsigned long GetOpenCLDeviceLocalMemorySize(MagickCLEnv clEnv)
1712 {
1713  cl_ulong localMemorySize;
1714  clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemorySize, NULL);
1715  return (unsigned long)localMemorySize;
1716 }
1717 
1718 MagickPrivate
1719  unsigned long GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv clEnv)
1720 {
1721  cl_ulong maxMemAllocSize;
1722  clEnv->library->clGetDeviceInfo(clEnv->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &maxMemAllocSize, NULL);
1723  return (unsigned long)maxMemAllocSize;
1724 }
1725 
1726 
1727 /*
1728  Beginning of the OpenCL device selection infrastructure
1729 */
1730 
1731 
1732 typedef enum {
1733  DS_SUCCESS = 0
1734  ,DS_INVALID_PROFILE = 1000
1735  ,DS_MEMORY_ERROR
1736  ,DS_INVALID_PERF_EVALUATOR_TYPE
1737  ,DS_INVALID_PERF_EVALUATOR
1738  ,DS_PERF_EVALUATOR_ERROR
1739  ,DS_FILE_ERROR
1740  ,DS_UNKNOWN_DEVICE_TYPE
1741  ,DS_PROFILE_FILE_ERROR
1742  ,DS_SCORE_SERIALIZER_ERROR
1743  ,DS_SCORE_DESERIALIZER_ERROR
1744 } ds_status;
1745 
1746 /* device type */
1747 typedef enum {
1748  DS_DEVICE_NATIVE_CPU = 0
1749  ,DS_DEVICE_OPENCL_DEVICE
1750 } ds_device_type;
1751 
1752 
1753 typedef struct {
1754  ds_device_type type;
1755  cl_device_type oclDeviceType;
1756  cl_device_id oclDeviceID;
1757  char* oclDeviceName;
1758  char* oclDriverVersion;
1759  cl_uint oclMaxClockFrequency;
1760  cl_uint oclMaxComputeUnits;
1761  void* score; /* a pointer to the score data, the content/format is application defined */
1762 } ds_device;
1763 
1764 typedef struct {
1765  unsigned int numDevices;
1766  ds_device* devices;
1767  const char* version;
1768 } ds_profile;
1769 
1770 /* deallocate memory used by score */
1771 typedef ds_status (*ds_score_release)(void* score);
1772 
1773 static ds_status releaseDeviceResource(ds_device* device, ds_score_release sr) {
1774  ds_status status = DS_SUCCESS;
1775  if (device) {
1776  if (device->oclDeviceName) RelinquishMagickMemory(device->oclDeviceName);
1777  if (device->oclDriverVersion) RelinquishMagickMemory(device->oclDriverVersion);
1778  if (device->score) status = sr(device->score);
1779  }
1780  return status;
1781 }
1782 
1783 static ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr) {
1784  ds_status status = DS_SUCCESS;
1785  if (profile!=NULL) {
1786  if (profile->devices!=NULL && sr!=NULL) {
1787  unsigned int i;
1788  for (i = 0; i < profile->numDevices; i++) {
1789  status = releaseDeviceResource(profile->devices+i,sr);
1790  if (status != DS_SUCCESS)
1791  break;
1792  }
1793  RelinquishMagickMemory(profile->devices);
1794  }
1795  RelinquishMagickMemory(profile);
1796  }
1797  return status;
1798 }
1799 
1800 
1801 static ds_status initDSProfile(ds_profile** p, const char* version) {
1802  int numDevices = 0;
1803  cl_uint numPlatforms = 0;
1804  cl_platform_id* platforms = NULL;
1805  cl_device_id* devices = NULL;
1806  ds_status status = DS_SUCCESS;
1807  ds_profile* profile = NULL;
1808  unsigned int next = 0;
1809  unsigned int i;
1810 
1811  if (p == NULL)
1812  return DS_INVALID_PROFILE;
1813 
1814  profile = (ds_profile*) AcquireMagickMemory(sizeof(ds_profile));
1815  if (profile == NULL)
1816  return DS_MEMORY_ERROR;
1817 
1818  memset(profile, 0, sizeof(ds_profile));
1819 
1820  OpenCLLib->clGetPlatformIDs(0, NULL, &numPlatforms);
1821  if (numPlatforms > 0) {
1822  platforms = (cl_platform_id*) AcquireQuantumMemory(numPlatforms,sizeof(cl_platform_id));
1823  if (platforms == NULL) {
1824  status = DS_MEMORY_ERROR;
1825  goto cleanup;
1826  }
1827  OpenCLLib->clGetPlatformIDs(numPlatforms, platforms, NULL);
1828  for (i = 0; i < (unsigned int)numPlatforms; i++) {
1829  cl_uint num;
1830  if (OpenCLLib->clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU, 0, NULL, &num) == CL_SUCCESS)
1831  numDevices+=num;
1832  }
1833  }
1834 
1835  profile->numDevices = numDevices+1; /* +1 to numDevices to include the native CPU */
1836 
1837  profile->devices = (ds_device*) AcquireQuantumMemory(profile->numDevices,sizeof(ds_device));
1838  if (profile->devices == NULL) {
1839  profile->numDevices = 0;
1840  status = DS_MEMORY_ERROR;
1841  goto cleanup;
1842  }
1843  memset(profile->devices, 0, profile->numDevices*sizeof(ds_device));
1844 
1845  if (numDevices > 0) {
1846  devices = (cl_device_id*) AcquireQuantumMemory(numDevices,sizeof(cl_device_id));
1847  if (devices == NULL) {
1848  status = DS_MEMORY_ERROR;
1849  goto cleanup;
1850  }
1851  for (i = 0; i < (unsigned int)numPlatforms; i++) {
1852  cl_uint num;
1853 
1854  int d;
1855  for (d = 0; d < 2; d++) {
1856  unsigned int j;
1857  cl_device_type deviceType;
1858  switch(d) {
1859  case 0:
1860  deviceType = CL_DEVICE_TYPE_GPU;
1861  break;
1862  case 1:
1863  deviceType = CL_DEVICE_TYPE_CPU;
1864  break;
1865  default:
1866  continue;
1867  break;
1868  }
1869  if (OpenCLLib->clGetDeviceIDs(platforms[i], deviceType, numDevices, devices, &num) != CL_SUCCESS)
1870  continue;
1871  for (j = 0; j < num; j++, next++) {
1872  size_t length;
1873 
1874  profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
1875  profile->devices[next].oclDeviceID = devices[j];
1876 
1877  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1878  , 0, NULL, &length);
1879  profile->devices[next].oclDeviceName = (char*) AcquireQuantumMemory(length,sizeof(char));
1880  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
1881  , length, profile->devices[next].oclDeviceName, NULL);
1882 
1883  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1884  , 0, NULL, &length);
1885  profile->devices[next].oclDriverVersion = (char*) AcquireQuantumMemory(length,sizeof(char));
1886  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
1887  , length, profile->devices[next].oclDriverVersion, NULL);
1888 
1889  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_CLOCK_FREQUENCY
1890  , sizeof(cl_uint), &profile->devices[next].oclMaxClockFrequency, NULL);
1891 
1892  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_MAX_COMPUTE_UNITS
1893  , sizeof(cl_uint), &profile->devices[next].oclMaxComputeUnits, NULL);
1894 
1895  OpenCLLib->clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_TYPE
1896  , sizeof(cl_device_type), &profile->devices[next].oclDeviceType, NULL);
1897  }
1898  }
1899  }
1900  }
1901 
1902  profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
1903  profile->version = version;
1904 
1905 cleanup:
1906  if (platforms) RelinquishMagickMemory(platforms);
1907  if (devices) RelinquishMagickMemory(devices);
1908  if (status == DS_SUCCESS) {
1909  *p = profile;
1910  }
1911  else {
1912  if (profile) {
1913  if (profile->devices)
1914  RelinquishMagickMemory(profile->devices);
1915  RelinquishMagickMemory(profile);
1916  }
1917  }
1918  return status;
1919 }
1920 
1921 /* Pointer to a function that calculates the score of a device (ex: device->score)
1922  update the data size of score. The encoding and the format of the score data
1923  is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
1924  */
1925 typedef ds_status (*ds_perf_evaluator)(ds_device* device, void* data);
1926 
1927 typedef enum {
1928  DS_EVALUATE_ALL
1929  ,DS_EVALUATE_NEW_ONLY
1930 } ds_evaluation_type;
1931 
1932 static ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type
1933  ,ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates) {
1934  ds_status status = DS_SUCCESS;
1935  unsigned int i;
1936  unsigned int updates = 0;
1937 
1938  if (profile == NULL) {
1939  return DS_INVALID_PROFILE;
1940  }
1941  if (evaluator == NULL) {
1942  return DS_INVALID_PERF_EVALUATOR;
1943  }
1944 
1945  for (i = 0; i < profile->numDevices; i++) {
1946  ds_status evaluatorStatus;
1947 
1948  switch (type) {
1949  case DS_EVALUATE_NEW_ONLY:
1950  if (profile->devices[i].score != NULL)
1951  break;
1952  /* else fall through */
1953  case DS_EVALUATE_ALL:
1954  evaluatorStatus = evaluator(profile->devices+i,evaluatorData);
1955  if (evaluatorStatus != DS_SUCCESS) {
1956  status = evaluatorStatus;
1957  return status;
1958  }
1959  updates++;
1960  break;
1961  default:
1962  return DS_INVALID_PERF_EVALUATOR_TYPE;
1963  break;
1964  };
1965  }
1966  if (numUpdates)
1967  *numUpdates = updates;
1968  return status;
1969 }
1970 
1971 
1972 #define DS_TAG_VERSION "<version>"
1973 #define DS_TAG_VERSION_END "</version>"
1974 #define DS_TAG_DEVICE "<device>"
1975 #define DS_TAG_DEVICE_END "</device>"
1976 #define DS_TAG_SCORE "<score>"
1977 #define DS_TAG_SCORE_END "</score>"
1978 #define DS_TAG_DEVICE_TYPE "<type>"
1979 #define DS_TAG_DEVICE_TYPE_END "</type>"
1980 #define DS_TAG_DEVICE_NAME "<name>"
1981 #define DS_TAG_DEVICE_NAME_END "</name>"
1982 #define DS_TAG_DEVICE_DRIVER_VERSION "<driver>"
1983 #define DS_TAG_DEVICE_DRIVER_VERSION_END "</driver>"
1984 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS "<max cu>"
1985 #define DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END "</max cu>"
1986 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ "<max clock>"
1987 #define DS_TAG_DEVICE_MAX_CLOCK_FREQ_END "</max clock>"
1988 
1989 #define DS_DEVICE_NATIVE_CPU_STRING "native_cpu"
1990 
1991 
1992 
1993 typedef ds_status (*ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
1994 static ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file) {
1995  ds_status status = DS_SUCCESS;
1996  FILE* profileFile = NULL;
1997 
1998 
1999  if (profile == NULL)
2000  return DS_INVALID_PROFILE;
2001 
2002  profileFile = fopen(file, "wb");
2003  if (profileFile==NULL) {
2004  status = DS_FILE_ERROR;
2005  }
2006  else {
2007  unsigned int i;
2008 
2009  /* write version string */
2010  fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
2011  fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
2012  fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
2013  fwrite("\n", sizeof(char), 1, profileFile);
2014 
2015  for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++) {
2016  void* serializedScore;
2017  unsigned int serializedScoreSize;
2018 
2019  fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);
2020 
2021  fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
2022  fwrite(&profile->devices[i].type,sizeof(ds_device_type),1, profileFile);
2023  fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);
2024 
2025  switch(profile->devices[i].type) {
2026  case DS_DEVICE_NATIVE_CPU:
2027  {
2028  /* There's no need to emit a device name for the native CPU device. */
2029  /*
2030  fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
2031  fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
2032  fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
2033  */
2034  }
2035  break;
2036  case DS_DEVICE_OPENCL_DEVICE:
2037  {
2038  char tmp[16];
2039 
2040  fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
2041  fwrite(profile->devices[i].oclDeviceName,sizeof(char),strlen(profile->devices[i].oclDeviceName), profileFile);
2042  fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
2043 
2044  fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
2045  fwrite(profile->devices[i].oclDriverVersion,sizeof(char),strlen(profile->devices[i].oclDriverVersion), profileFile);
2046  fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);
2047 
2048  fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS), profileFile);
2049  (void) FormatLocaleString(tmp,sizeof(tmp),"%d",
2050  profile->devices[i].oclMaxComputeUnits);
2051  fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
2052  fwrite(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END), profileFile);
2053 
2054  fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ), profileFile);
2055  (void) FormatLocaleString(tmp,sizeof(tmp),"%d",
2056  profile->devices[i].oclMaxClockFrequency);
2057  fwrite(tmp,sizeof(char),strlen(tmp), profileFile);
2058  fwrite(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END, sizeof(char), strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ_END), profileFile);
2059  }
2060  break;
2061  default:
2062  status = DS_UNKNOWN_DEVICE_TYPE;
2063  break;
2064  };
2065 
2066  fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
2067  status = serializer(profile->devices+i, &serializedScore, &serializedScoreSize);
2068  if (status == DS_SUCCESS && serializedScore!=NULL && serializedScoreSize > 0) {
2069  fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
2070  RelinquishMagickMemory(serializedScore);
2071  }
2072  fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
2073  fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
2074  fwrite("\n",sizeof(char),1,profileFile);
2075  }
2076  fclose(profileFile);
2077  }
2078  return status;
2079 }
2080 
2081 
2082 static ds_status readProFile(const char* fileName, char** content, size_t* contentSize) {
2083  ds_status status = DS_SUCCESS;
2084  FILE * input = NULL;
2085  size_t size = 0;
2086  size_t rsize = 0;
2087  char* binary = NULL;
2088 
2089  *contentSize = 0;
2090  *content = NULL;
2091 
2092  input = fopen(fileName, "rb");
2093  if(input == NULL) {
2094  return DS_FILE_ERROR;
2095  }
2096 
2097  fseek(input, 0L, SEEK_END);
2098  size = ftell(input);
2099  rewind(input);
2100  binary = (char*) AcquireQuantumMemory(1,size);
2101  if(binary == NULL) {
2102  status = DS_FILE_ERROR;
2103  goto cleanup;
2104  }
2105  rsize = fread(binary, sizeof(char), size, input);
2106  if (rsize!=size
2107  || ferror(input)) {
2108  status = DS_FILE_ERROR;
2109  goto cleanup;
2110  }
2111  *contentSize = size;
2112  *content = binary;
2113 
2114 cleanup:
2115  if (input != NULL) fclose(input);
2116  if (status != DS_SUCCESS
2117  && binary != NULL) {
2118  RelinquishMagickMemory(binary);
2119  *content = NULL;
2120  *contentSize = 0;
2121  }
2122  return status;
2123 }
2124 
2125 
2126 static const char* findString(const char* contentStart, const char* contentEnd, const char* string) {
2127  size_t stringLength;
2128  const char* currentPosition;
2129  const char* found;
2130  found = NULL;
2131  stringLength = strlen(string);
2132  currentPosition = contentStart;
2133  for(currentPosition = contentStart; currentPosition < contentEnd; currentPosition++) {
2134  if (*currentPosition == string[0]) {
2135  if (currentPosition+stringLength < contentEnd) {
2136  if (strncmp(currentPosition, string, stringLength) == 0) {
2137  found = currentPosition;
2138  break;
2139  }
2140  }
2141  }
2142  }
2143  return found;
2144 }
2145 
2146 
2147 typedef ds_status (*ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
2148 static ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file) {
2149 
2150  ds_status status = DS_SUCCESS;
2151  char* contentStart = NULL;
2152  const char* contentEnd = NULL;
2153  size_t contentSize;
2154 
2155  if (profile==NULL)
2156  return DS_INVALID_PROFILE;
2157 
2158  status = readProFile(file, &contentStart, &contentSize);
2159  if (status == DS_SUCCESS) {
2160  const char* currentPosition;
2161  const char* dataStart;
2162  const char* dataEnd;
2163  size_t versionStringLength;
2164 
2165  contentEnd = contentStart + contentSize;
2166  currentPosition = contentStart;
2167 
2168 
2169  /* parse the version string */
2170  dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
2171  if (dataStart == NULL) {
2172  status = DS_PROFILE_FILE_ERROR;
2173  goto cleanup;
2174  }
2175  dataStart += strlen(DS_TAG_VERSION);
2176 
2177  dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
2178  if (dataEnd==NULL) {
2179  status = DS_PROFILE_FILE_ERROR;
2180  goto cleanup;
2181  }
2182 
2183  versionStringLength = strlen(profile->version);
2184  if (versionStringLength!=(size_t)(dataEnd-dataStart)
2185  || strncmp(profile->version, dataStart, versionStringLength)!=(int)0) {
2186  /* version mismatch */
2187  status = DS_PROFILE_FILE_ERROR;
2188  goto cleanup;
2189  }
2190  currentPosition = dataEnd+strlen(DS_TAG_VERSION_END);
2191 
2192  /* parse the device information */
2193 DisableMSCWarning(4127)
2194  while (1) {
2195 RestoreMSCWarning
2196  unsigned int i;
2197 
2198  const char* deviceTypeStart;
2199  const char* deviceTypeEnd;
2200  ds_device_type deviceType;
2201 
2202  const char* deviceNameStart;
2203  const char* deviceNameEnd;
2204 
2205  const char* deviceScoreStart;
2206  const char* deviceScoreEnd;
2207 
2208  const char* deviceDriverStart;
2209  const char* deviceDriverEnd;
2210 
2211  const char* tmpStart;
2212  const char* tmpEnd;
2213  char tmp[16];
2214 
2215  cl_uint maxClockFrequency;
2216  cl_uint maxComputeUnits;
2217 
2218  dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
2219  if (dataStart==NULL) {
2220  /* nothing useful remain, quit...*/
2221  break;
2222  }
2223  dataStart+=strlen(DS_TAG_DEVICE);
2224  dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
2225  if (dataEnd==NULL) {
2226  status = DS_PROFILE_FILE_ERROR;
2227  goto cleanup;
2228  }
2229 
2230  /* parse the device type */
2231  deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
2232  if (deviceTypeStart==NULL) {
2233  status = DS_PROFILE_FILE_ERROR;
2234  goto cleanup;
2235  }
2236  deviceTypeStart+=strlen(DS_TAG_DEVICE_TYPE);
2237  deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
2238  if (deviceTypeEnd==NULL) {
2239  status = DS_PROFILE_FILE_ERROR;
2240  goto cleanup;
2241  }
2242  memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));
2243 
2244 
2245  /* parse the device name */
2246  if (deviceType == DS_DEVICE_OPENCL_DEVICE) {
2247 
2248  deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
2249  if (deviceNameStart==NULL) {
2250  status = DS_PROFILE_FILE_ERROR;
2251  goto cleanup;
2252  }
2253  deviceNameStart+=strlen(DS_TAG_DEVICE_NAME);
2254  deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
2255  if (deviceNameEnd==NULL) {
2256  status = DS_PROFILE_FILE_ERROR;
2257  goto cleanup;
2258  }
2259 
2260 
2261  deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
2262  if (deviceDriverStart==NULL) {
2263  status = DS_PROFILE_FILE_ERROR;
2264  goto cleanup;
2265  }
2266  deviceDriverStart+=strlen(DS_TAG_DEVICE_DRIVER_VERSION);
2267  deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
2268  if (deviceDriverEnd ==NULL) {
2269  status = DS_PROFILE_FILE_ERROR;
2270  goto cleanup;
2271  }
2272 
2273 
2274  tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
2275  if (tmpStart==NULL) {
2276  status = DS_PROFILE_FILE_ERROR;
2277  goto cleanup;
2278  }
2279  tmpStart+=strlen(DS_TAG_DEVICE_MAX_COMPUTE_UNITS);
2280  tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_COMPUTE_UNITS_END);
2281  if (tmpEnd ==NULL) {
2282  status = DS_PROFILE_FILE_ERROR;
2283  goto cleanup;
2284  }
2285  memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2286  tmp[tmpEnd-tmpStart] = '\0';
2287  maxComputeUnits = strtol(tmp,(char **) NULL,10);
2288 
2289 
2290  tmpStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2291  if (tmpStart==NULL) {
2292  status = DS_PROFILE_FILE_ERROR;
2293  goto cleanup;
2294  }
2295  tmpStart+=strlen(DS_TAG_DEVICE_MAX_CLOCK_FREQ);
2296  tmpEnd = findString(tmpStart, contentEnd, DS_TAG_DEVICE_MAX_CLOCK_FREQ_END);
2297  if (tmpEnd ==NULL) {
2298  status = DS_PROFILE_FILE_ERROR;
2299  goto cleanup;
2300  }
2301  memcpy(tmp,tmpStart,tmpEnd-tmpStart);
2302  tmp[tmpEnd-tmpStart] = '\0';
2303  maxClockFrequency = strtol(tmp,(char **) NULL,10);
2304 
2305 
2306  /* check if this device is on the system */
2307  for (i = 0; i < profile->numDevices; i++) {
2308  if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) {
2309  size_t actualDeviceNameLength;
2310  size_t driverVersionLength;
2311 
2312  actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
2313  driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
2314  if (actualDeviceNameLength == (size_t)(deviceNameEnd - deviceNameStart)
2315  && driverVersionLength == (size_t)(deviceDriverEnd - deviceDriverStart)
2316  && maxComputeUnits == profile->devices[i].oclMaxComputeUnits
2317  && maxClockFrequency == profile->devices[i].oclMaxClockFrequency
2318  && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength)==(int)0
2319  && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength)==(int)0) {
2320 
2321  deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2322  if (deviceNameStart==NULL) {
2323  status = DS_PROFILE_FILE_ERROR;
2324  goto cleanup;
2325  }
2326  deviceScoreStart+=strlen(DS_TAG_SCORE);
2327  deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2328  status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2329  if (status != DS_SUCCESS) {
2330  goto cleanup;
2331  }
2332  }
2333  }
2334  }
2335 
2336  }
2337  else if (deviceType == DS_DEVICE_NATIVE_CPU) {
2338  for (i = 0; i < profile->numDevices; i++) {
2339  if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU) {
2340  deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
2341  if (deviceScoreStart==NULL) {
2342  status = DS_PROFILE_FILE_ERROR;
2343  goto cleanup;
2344  }
2345  deviceScoreStart+=strlen(DS_TAG_SCORE);
2346  deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
2347  status = deserializer(profile->devices+i, (const unsigned char*)deviceScoreStart, deviceScoreEnd-deviceScoreStart);
2348  if (status != DS_SUCCESS) {
2349  goto cleanup;
2350  }
2351  }
2352  }
2353  }
2354 
2355  /* skip over the current one to find the next device */
2356  currentPosition = dataEnd+strlen(DS_TAG_DEVICE_END);
2357  }
2358  }
2359 cleanup:
2360  if (contentStart!=NULL) RelinquishMagickMemory(contentStart);
2361  return status;
2362 }
2363 
2364 
2365 #if 0
2366 static ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num) {
2367  unsigned int i;
2368  if (profile == NULL || num==NULL)
2369  return DS_MEMORY_ERROR;
2370  *num=0;
2371  for (i = 0; i < profile->numDevices; i++) {
2372  if (profile->devices[i].score == NULL) {
2373  (*num)++;
2374  }
2375  }
2376  return DS_SUCCESS;
2377 }
2378 #endif
2379 
2380 /*
2381  End of the OpenCL device selection infrastructure
2382 */
2383 
2384 
2385 typedef double AccelerateScoreType;
2386 
2387 static ds_status AcceleratePerfEvaluator(ds_device *device,
2388  void *magick_unused(data))
2389 {
2390 #define ACCELERATE_PERF_DIMEN "2048x1536"
2391 #define NUM_ITER 2
2392 #define ReturnStatus(status) \
2393 { \
2394  if (oldClEnv != (MagickCLEnv) NULL) \
2395  defaultCLEnv=oldClEnv; \
2396  if (clEnv != (MagickCLEnv) NULL) \
2397  (void) RelinquishMagickOpenCLEnv(clEnv); \
2398  return status; \
2399 }
2400 
2401  AccelerateTimer
2402  timer;
2403 
2404  ExceptionInfo
2405  *exception=NULL;
2406 
2407  MagickBooleanType
2408  status;
2409 
2410  MagickCLEnv
2411  clEnv=NULL,
2412  oldClEnv=NULL;
2413 
2414  magick_unreferenced(data);
2415 
2416  if (device == NULL)
2417  ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2418 
2419  clEnv=AcquireMagickOpenCLEnv();
2420  exception=AcquireExceptionInfo();
2421 
2422  if (device->type == DS_DEVICE_NATIVE_CPU)
2423  {
2424  /* CPU device */
2425  MagickBooleanType flag=MagickTrue;
2426  SetMagickOpenCLEnvParamInternal(clEnv,
2427  MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED,sizeof(MagickBooleanType),
2428  &flag,exception);
2429  }
2430  else if (device->type == DS_DEVICE_OPENCL_DEVICE)
2431  {
2432  /* OpenCL device */
2433  SetMagickOpenCLEnvParamInternal(clEnv,MAGICK_OPENCL_ENV_PARAM_DEVICE,
2434  sizeof(cl_device_id),&device->oclDeviceID,exception);
2435  }
2436  else
2437  ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2438 
2439  /* recompile the OpenCL kernels if it needs to */
2440  clEnv->disableProgramCache = defaultCLEnv->disableProgramCache;
2441 
2442  status=InitOpenCLEnvInternal(clEnv,exception);
2443  oldClEnv=defaultCLEnv;
2444  defaultCLEnv=clEnv;
2445 
2446  /* microbenchmark */
2447  if (status != MagickFalse)
2448  {
2449  Image
2450  *inputImage;
2451 
2452  ImageInfo
2453  *imageInfo;
2454 
2455  int
2456  i;
2457 
2458  imageInfo=AcquireImageInfo();
2459  CloneString(&imageInfo->size,ACCELERATE_PERF_DIMEN);
2460  CopyMagickString(imageInfo->filename,"xc:none",MaxTextExtent);
2461  inputImage=ReadImage(imageInfo,exception);
2462  if (inputImage == (Image *) NULL)
2463  ReturnStatus(DS_PERF_EVALUATOR_ERROR);
2464 
2465  initAccelerateTimer(&timer);
2466 
2467  for (i=0; i<=NUM_ITER; i++)
2468  {
2469  cl_uint
2470  event_count;
2471 
2472  cl_event
2473  *events;
2474 
2475  Image
2476  *bluredImage,
2477  *resizedImage,
2478  *unsharpedImage;
2479 
2480  if (i > 0)
2481  startAccelerateTimer(&timer);
2482 
2483 #ifdef MAGICKCORE_CLPERFMARKER
2484  clBeginPerfMarkerAMD("PerfEvaluatorRegion","");
2485 #endif
2486 
2487  bluredImage=BlurImage(inputImage,10.0f,3.5f,exception);
2488  unsharpedImage=UnsharpMaskImage(bluredImage,2.0f,2.0f,50.0f,10.0f,
2489  exception);
2490  resizedImage=ResizeImage(unsharpedImage,640,480,LanczosFilter,1.0,
2491  exception);
2492 
2493  /*
2494  We need this to get a proper performance benchmark, the operations
2495  are executed asynchronous.
2496  */
2497  if (device->type != DS_DEVICE_NATIVE_CPU)
2498  {
2499  events=GetOpenCLEvents(resizedImage,&event_count);
2500  if (event_count > 0)
2501  clEnv->library->clWaitForEvents(event_count,events);
2502  events=(cl_event *) RelinquishMagickMemory(events);
2503  }
2504 
2505 #ifdef MAGICKCORE_CLPERFMARKER
2506  clEndPerfMarkerAMD();
2507 #endif
2508 
2509  if (i > 0)
2510  stopAccelerateTimer(&timer);
2511 
2512  if (bluredImage)
2513  DestroyImage(bluredImage);
2514  if (unsharpedImage)
2515  DestroyImage(unsharpedImage);
2516  if (resizedImage)
2517  DestroyImage(resizedImage);
2518  }
2519  DestroyImage(inputImage);
2520  }
2521  /* end of microbenchmark */
2522 
2523  if (device->score == NULL)
2524  device->score= AcquireMagickMemory(sizeof(AccelerateScoreType));
2525 
2526  if (status != MagickFalse)
2527  *(AccelerateScoreType*) device->score=readAccelerateTimer(&timer);
2528  else
2529  *(AccelerateScoreType*) device->score=42;
2530 
2531  ReturnStatus(DS_SUCCESS);
2532 }
2533 
2534 ds_status AccelerateScoreSerializer(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize) {
2535  if (device
2536  && device->score) {
2537  /* generate a string from the score */
2538  char* s = (char*) AcquireQuantumMemory(256,sizeof(char));
2539  (void) FormatLocaleString(s,256,"%.4f",*((AccelerateScoreType*)
2540  device->score));
2541  *serializedScore = (void*)s;
2542  *serializedScoreSize = (unsigned int) strlen(s);
2543  return DS_SUCCESS;
2544  }
2545  else {
2546  return DS_SCORE_SERIALIZER_ERROR;
2547  }
2548 }
2549 
2550 ds_status AccelerateScoreDeserializer(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize) {
2551  if (device) {
2552  /* convert the string back to an int */
2553  char* s = (char*) AcquireQuantumMemory(1,serializedScoreSize+1);
2554  memcpy(s, serializedScore, serializedScoreSize);
2555  s[serializedScoreSize] = (char)'\0';
2556  device->score = AcquireMagickMemory(sizeof(AccelerateScoreType));
2557  *((AccelerateScoreType*)device->score) = (AccelerateScoreType)
2558  strtod(s, (char **) NULL);
2559  RelinquishMagickMemory(s);
2560  return DS_SUCCESS;
2561  }
2562  else {
2563  return DS_SCORE_DESERIALIZER_ERROR;
2564  }
2565 }
2566 
2567 ds_status AccelerateScoreRelease(void* score) {
2568  if (score!=NULL) {
2569  RelinquishMagickMemory(score);
2570  }
2571  return DS_SUCCESS;
2572 }
2573 
2574 ds_status canWriteProfileToFile(const char *path)
2575 {
2576  FILE* profileFile = fopen(path, "ab");
2577 
2578  if (profileFile==NULL)
2579  return DS_FILE_ERROR;
2580 
2581  fclose(profileFile);
2582  return DS_SUCCESS;
2583 }
2584 
2585 
2586 #define IMAGEMAGICK_PROFILE_VERSION "ImageMagick Device Selection v0.9"
2587 #define IMAGEMAGICK_PROFILE_FILE "ImagemagickOpenCLDeviceProfile"
2588 static MagickBooleanType autoSelectDevice(MagickCLEnv clEnv, ExceptionInfo* exception) {
2589 
2590  MagickBooleanType mStatus = MagickFalse;
2591  ds_status status;
2592  ds_profile* profile;
2593  unsigned int numDeviceProfiled = 0;
2594  unsigned int i;
2595  unsigned int bestDeviceIndex;
2596  AccelerateScoreType bestScore;
2597  char path[MaxTextExtent];
2598  MagickBooleanType flag;
2599  ds_evaluation_type profileType;
2600 
2601  LockDefaultOpenCLEnv();
2602 
2603  /* Initially, just set OpenCL to off */
2604  flag = MagickTrue;
2605  SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2606  , sizeof(MagickBooleanType), &flag, exception);
2607 
2608  /* check and init the global lib */
2609  OpenCLLib=GetOpenCLLib();
2610  if (OpenCLLib==NULL)
2611  {
2612  mStatus=InitOpenCLEnvInternal(clEnv, exception);
2613  goto cleanup;
2614  }
2615 
2616  clEnv->library=OpenCLLib;
2617 
2618  status = initDSProfile(&profile, IMAGEMAGICK_PROFILE_VERSION);
2619  if (status!=DS_SUCCESS) {
2620  (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2621  goto cleanup;
2622  }
2623 
2624  (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
2625  ,GetOpenCLCachedFilesDirectory()
2626  ,DirectorySeparator,IMAGEMAGICK_PROFILE_FILE);
2627 
2628  if (canWriteProfileToFile(path) != DS_SUCCESS) {
2629  /* We can not write out a device profile, so don't run the benchmark */
2630  /* select the first GPU device */
2631 
2632  bestDeviceIndex = 0;
2633  for (i = 1; i < profile->numDevices; i++) {
2634  if ((profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE) && (profile->devices[i].oclDeviceType == CL_DEVICE_TYPE_GPU)) {
2635  bestDeviceIndex = i;
2636  break;
2637  }
2638  }
2639  }
2640  else {
2641  if (clEnv->regenerateProfile != MagickFalse) {
2642  profileType = DS_EVALUATE_ALL;
2643  }
2644  else {
2645  readProfileFromFile(profile, AccelerateScoreDeserializer, path);
2646  profileType = DS_EVALUATE_NEW_ONLY;
2647  }
2648  status = profileDevices(profile, profileType, AcceleratePerfEvaluator, NULL, &numDeviceProfiled);
2649 
2650  if (status!=DS_SUCCESS) {
2651  (void) ThrowMagickException(exception, GetMagickModule(), ModuleFatalError, "Error when initializing the profile", "'%s'", ".");
2652  goto cleanup;
2653  }
2654  if (numDeviceProfiled > 0) {
2655  status = writeProfileToFile(profile, AccelerateScoreSerializer, path);
2656  if (status!=DS_SUCCESS) {
2657  (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when saving the profile into a file", "'%s'", ".");
2658  }
2659  }
2660 
2661  /* pick the best device */
2662  bestDeviceIndex = 0;
2663  bestScore = *(AccelerateScoreType*)profile->devices[bestDeviceIndex].score;
2664  for (i = 1; i < profile->numDevices; i++) {
2665  AccelerateScoreType score = *(AccelerateScoreType*)profile->devices[i].score;
2666  if (score < bestScore) {
2667  bestDeviceIndex = i;
2668  bestScore = score;
2669  }
2670  }
2671  }
2672 
2673  /* set up clEnv with the best device */
2674  if (profile->devices[bestDeviceIndex].type == DS_DEVICE_NATIVE_CPU) {
2675  /* CPU device */
2676  flag = MagickTrue;
2677  SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2678  , sizeof(MagickBooleanType), &flag, exception);
2679  }
2680  else if (profile->devices[bestDeviceIndex].type == DS_DEVICE_OPENCL_DEVICE) {
2681  /* OpenCL device */
2682  flag = MagickFalse;
2683  SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2684  , sizeof(MagickBooleanType), &flag, exception);
2685  SetMagickOpenCLEnvParamInternal(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2686  , sizeof(cl_device_id), &profile->devices[bestDeviceIndex].oclDeviceID,exception);
2687  }
2688  else {
2689  status = DS_PERF_EVALUATOR_ERROR;
2690  goto cleanup;
2691  }
2692  mStatus=InitOpenCLEnvInternal(clEnv, exception);
2693 
2694  status = releaseDSProfile(profile, AccelerateScoreRelease);
2695  if (status!=DS_SUCCESS) {
2696  (void) ThrowMagickException(exception, GetMagickModule(), ModuleWarning, "Error when releasing the profile", "'%s'", ".");
2697  }
2698 
2699 cleanup:
2700 
2701  UnlockDefaultOpenCLEnv();
2702  return mStatus;
2703 }
2704 
2705 
2706 /*
2707 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2708 % %
2709 % %
2710 % %
2711 + I n i t I m a g e M a g i c k O p e n C L %
2712 % %
2713 % %
2714 % %
2715 %%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%
2716 %
2717 % InitImageMagickOpenCL() provides a simplified interface to initialize
2718 % the OpenCL environtment in ImageMagick
2719 %
2720 % The format of the InitImageMagickOpenCL() method is:
2721 %
2722 % MagickBooleanType InitImageMagickOpenCL(ImageMagickOpenCLMode mode,
2723 % void* userSelectedDevice,
2724 % void* selectedDevice)
2725 %
2726 % A description of each parameter follows:
2727 %
2728 % o mode: OpenCL mode in ImageMagick, could be off,auto,user
2729 %
2730 % o userSelectedDevice: when in user mode, a pointer to the selected
2731 % cl_device_id
2732 %
2733 % o selectedDevice: a pointer to cl_device_id where the selected
2734 % cl_device_id by ImageMagick could be returned
2735 %
2736 % o exception: exception
2737 %
2738 */
2739 MagickExport MagickBooleanType InitImageMagickOpenCL(
2740  ImageMagickOpenCLMode mode,void *userSelectedDevice,void *selectedDevice,
2741  ExceptionInfo *exception)
2742 {
2743  MagickBooleanType status = MagickFalse;
2744  MagickCLEnv clEnv = NULL;
2745  MagickBooleanType flag;
2746 
2747  clEnv = GetDefaultOpenCLEnv();
2748  if (clEnv!=NULL) {
2749  switch(mode) {
2750 
2751  case MAGICK_OPENCL_OFF:
2752  flag = MagickTrue;
2753  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2754  , sizeof(MagickBooleanType), &flag, exception);
2755  status = InitOpenCLEnv(clEnv, exception);
2756 
2757  if (selectedDevice)
2758  *(cl_device_id*)selectedDevice = NULL;
2759  break;
2760 
2761  case MAGICK_OPENCL_DEVICE_SELECT_USER:
2762 
2763  if (userSelectedDevice == NULL)
2764  return MagickFalse;
2765 
2766  flag = MagickFalse;
2767  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2768  , sizeof(MagickBooleanType), &flag, exception);
2769 
2770  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2771  , sizeof(cl_device_id), userSelectedDevice,exception);
2772 
2773  status = InitOpenCLEnv(clEnv, exception);
2774  if (selectedDevice) {
2775  GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2776  , sizeof(cl_device_id), selectedDevice, exception);
2777  }
2778  break;
2779 
2780  case MAGICK_OPENCL_DEVICE_SELECT_AUTO_CLEAR_CACHE:
2781  flag = MagickTrue;
2782  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_PROGRAM_CACHE_DISABLED
2783  , sizeof(MagickBooleanType), &flag, exception);
2784  flag = MagickTrue;
2785  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_REGENERATE_PROFILE
2786  , sizeof(MagickBooleanType), &flag, exception);
2787 
2788  /* fall through here!! */
2789  case MAGICK_OPENCL_DEVICE_SELECT_AUTO:
2790  default:
2791  {
2792  cl_device_id d = NULL;
2793  flag = MagickFalse;
2794  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_OPENCL_DISABLED
2795  , sizeof(MagickBooleanType), &flag, exception);
2796  SetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2797  , sizeof(cl_device_id), &d,exception);
2798  status = InitOpenCLEnv(clEnv, exception);
2799  if (selectedDevice) {
2800  GetMagickOpenCLEnvParam(clEnv, MAGICK_OPENCL_ENV_PARAM_DEVICE
2801  , sizeof(cl_device_id), selectedDevice, exception);
2802  }
2803  }
2804  break;
2805  };
2806  }
2807  return status;
2808 }
2809 
2810 
2811 MagickPrivate
2812 MagickBooleanType OpenCLThrowMagickException(ExceptionInfo *exception,
2813  const char *module,const char *function,const size_t line,
2814  const ExceptionType severity,const char *tag,const char *format,...) {
2815  MagickBooleanType
2816  status;
2817 
2818  MagickCLEnv clEnv;
2819 
2820  status = MagickTrue;
2821 
2822  clEnv = GetDefaultOpenCLEnv();
2823 
2824  assert(exception != (ExceptionInfo *) NULL);
2825  assert(exception->signature == MagickCoreSignature);
2826 
2827  if (severity!=0) {
2828  cl_device_type dType;
2829  clEnv->library->clGetDeviceInfo(clEnv->device,CL_DEVICE_TYPE ,sizeof(cl_device_type),&dType,NULL);
2830  if (dType == CL_DEVICE_TYPE_CPU) {
2831  char buffer[MaxTextExtent];
2832  clEnv->library->clGetPlatformInfo(clEnv->platform, CL_PLATFORM_NAME, MaxTextExtent, buffer, NULL);
2833 
2834  /* Workaround for Intel OpenCL CPU runtime bug */
2835  /* Turn off OpenCL when a problem is detected! */
2836  if (strncmp(buffer, "Intel",5) == 0) {
2837 
2838  InitImageMagickOpenCL(MAGICK_OPENCL_OFF, NULL, NULL, exception);
2839  }
2840  }
2841  }
2842 
2843 #ifdef OPENCLLOG_ENABLED
2844  {
2845  va_list
2846  operands;
2847  va_start(operands,format);
2848  status=ThrowMagickExceptionList(exception,module,function,line,severity,tag, format,operands);
2849  va_end(operands);
2850  }
2851 #else
2852  magick_unreferenced(module);
2853  magick_unreferenced(function);
2854  magick_unreferenced(line);
2855  magick_unreferenced(tag);
2856  magick_unreferenced(format);
2857 #endif
2858 
2859  return(status);
2860 }
2861 
2862 char* openclCachedFilesDirectory;
2863 SemaphoreInfo* openclCachedFilesDirectoryLock;
2864 
2865 MagickPrivate
2866 const char* GetOpenCLCachedFilesDirectory() {
2867  if (openclCachedFilesDirectory == NULL) {
2868  if (openclCachedFilesDirectoryLock == NULL)
2869  {
2870  ActivateSemaphoreInfo(&openclCachedFilesDirectoryLock);
2871  }
2872  LockSemaphoreInfo(openclCachedFilesDirectoryLock);
2873  if (openclCachedFilesDirectory == NULL) {
2874  char path[MaxTextExtent];
2875  char *home = NULL;
2876  char *temp = NULL;
2877  struct stat attributes;
2878  MagickBooleanType status;
2879  int mkdirStatus = 0;
2880 
2881 
2882 
2883  home=GetEnvironmentValue("MAGICK_OPENCL_CACHE_DIR");
2884  if (home == (char *) NULL)
2885  {
2886  home=GetEnvironmentValue("XDG_CACHE_HOME");
2887 #if defined(MAGICKCORE_WINDOWS_SUPPORT) || defined(__MINGW32__)
2888  if (home == (char *) NULL)
2889  home=GetEnvironmentValue("LOCALAPPDATA");
2890  if (home == (char *) NULL)
2891  home=GetEnvironmentValue("APPDATA");
2892  if (home == (char *) NULL)
2893  home=GetEnvironmentValue("USERPROFILE");
2894 #endif
2895  }
2896 
2897  if (home != (char *) NULL)
2898  {
2899  /* first check if $HOME exists */
2900  (void) FormatLocaleString(path,MaxTextExtent,"%s",home);
2901  status=GetPathAttributes(path,&attributes);
2902  if (status == MagickFalse)
2903  {
2904 
2905 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2906  mkdirStatus = mkdir(path);
2907 #else
2908  mkdirStatus = mkdir(path, 0777);
2909 #endif
2910  }
2911 
2912  /* first check if $HOME/ImageMagick exists */
2913  if (mkdirStatus==0)
2914  {
2915  (void) FormatLocaleString(path,MaxTextExtent,
2916  "%s%sImageMagick",home,DirectorySeparator);
2917 
2918  status=GetPathAttributes(path,&attributes);
2919  if (status == MagickFalse)
2920  {
2921 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2922  mkdirStatus = mkdir(path);
2923 #else
2924  mkdirStatus = mkdir(path, 0777);
2925 #endif
2926  }
2927  }
2928 
2929  if (mkdirStatus==0)
2930  {
2931  temp = (char*)AcquireCriticalMemory(strlen(path)+1);
2932  CopyMagickString(temp,path,strlen(path)+1);
2933  }
2934  home=DestroyString(home);
2935  } else {
2936  home=GetEnvironmentValue("HOME");
2937  if (home != (char *) NULL)
2938  {
2939  /*
2940  */
2941 
2942  /* first check if $HOME/.cache exists */
2943  (void) FormatLocaleString(path,MaxTextExtent,"%s%s.cache",
2944  home,DirectorySeparator);
2945  status=GetPathAttributes(path,&attributes);
2946  if (status == MagickFalse)
2947  {
2948 
2949 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2950  mkdirStatus = mkdir(path);
2951 #else
2952  mkdirStatus = mkdir(path, 0777);
2953 #endif
2954  }
2955 
2956  /* first check if $HOME/.cache/ImageMagick exists */
2957  if (mkdirStatus==0)
2958  {
2959  (void) FormatLocaleString(path,MaxTextExtent,
2960  "%s%s.cache%sImageMagick",home,DirectorySeparator,
2961  DirectorySeparator);
2962 
2963  status=GetPathAttributes(path,&attributes);
2964  if (status == MagickFalse)
2965  {
2966 #ifdef MAGICKCORE_WINDOWS_SUPPORT
2967  mkdirStatus = mkdir(path);
2968 #else
2969  mkdirStatus = mkdir(path, 0777);
2970 #endif
2971  }
2972  }
2973 
2974  if (mkdirStatus==0)
2975  {
2976  temp = (char*)AcquireCriticalMemory(strlen(path)+1);
2977  CopyMagickString(temp,path,strlen(path)+1);
2978  }
2979  home=DestroyString(home);
2980  }
2981  }
2982  openclCachedFilesDirectory = temp;
2983  }
2984  UnlockSemaphoreInfo(openclCachedFilesDirectoryLock);
2985  }
2986  return openclCachedFilesDirectory;
2987 }
2988 
2989 /* create a function for OpenCL log */
2990 MagickPrivate
2991 void OpenCLLog(const char* message) {
2992 
2993 #ifdef OPENCLLOG_ENABLED
2994 #define OPENCL_LOG_FILE "ImageMagickOpenCL.log"
2995 
2996  FILE* log;
2997  if (getenv("MAGICK_OCL_LOG"))
2998  {
2999  if (message) {
3000  char path[MaxTextExtent];
3001  unsigned long allocSize;
3002 
3003  MagickCLEnv clEnv;
3004 
3005  clEnv = GetDefaultOpenCLEnv();
3006 
3007  /* dump the source into a file */
3008  (void) FormatLocaleString(path,MaxTextExtent,"%s%s%s"
3009  ,GetOpenCLCachedFilesDirectory()
3010  ,DirectorySeparator,OPENCL_LOG_FILE);
3011 
3012 
3013  log = fopen(path, "ab");
3014  if (log == (FILE *) NULL)
3015  return;
3016  fwrite(message, sizeof(char), strlen(message), log);
3017  fwrite("\n", sizeof(char), 1, log);
3018 
3019  if (clEnv->OpenCLInitialized && !clEnv->OpenCLDisabled)
3020  {
3021  allocSize = GetOpenCLDeviceMaxMemAllocSize(clEnv);
3022  fprintf(log, "Devic Max Memory Alloc Size: %lu\n", allocSize);
3023  }
3024 
3025  fclose(log);
3026  }
3027  }
3028 #else
3029  magick_unreferenced(message);
3030 #endif
3031 }
3032 
3033 MagickPrivate void OpenCLTerminus()
3034 {
3035  DumpProfileData();
3036  if (openclCachedFilesDirectory != (char *) NULL)
3037  openclCachedFilesDirectory=DestroyString(openclCachedFilesDirectory);
3038  if (openclCachedFilesDirectoryLock != (SemaphoreInfo*)NULL)
3039  DestroySemaphoreInfo(&openclCachedFilesDirectoryLock);
3040  if (defaultCLEnv != (MagickCLEnv) NULL)
3041  {
3042  (void) RelinquishMagickOpenCLEnv(defaultCLEnv);
3043  defaultCLEnv=(MagickCLEnv)NULL;
3044  }
3045  if (defaultCLEnvLock != (SemaphoreInfo*) NULL)
3046  DestroySemaphoreInfo(&defaultCLEnvLock);
3047  if (OpenCLLib != (MagickLibrary *)NULL)
3048  {
3049  if (OpenCLLib->base != (void *) NULL)
3050  (void) lt_dlclose(OpenCLLib->base);
3051  OpenCLLib=(MagickLibrary *)RelinquishMagickMemory(OpenCLLib);
3052  }
3053  if (OpenCLLibLock != (SemaphoreInfo*)NULL)
3054  DestroySemaphoreInfo(&OpenCLLibLock);
3055 }
3056 
3057 #else
3058 
3059 struct _MagickCLEnv {
3060  MagickBooleanType OpenCLInitialized; /* whether OpenCL environment is initialized. */
3061 };
3062 
3063 /*
3064 * Return the OpenCL environment
3065 */
3066 MagickExport MagickCLEnv GetDefaultOpenCLEnv()
3067 {
3068  return (MagickCLEnv) NULL;
3069 }
3070 
3071 MagickExport MagickBooleanType SetMagickOpenCLEnvParam(
3072  MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
3073  size_t magick_unused(dataSize),void *magick_unused(data),
3074  ExceptionInfo *magick_unused(exception))
3075 {
3076  magick_unreferenced(clEnv);
3077  magick_unreferenced(param);
3078  magick_unreferenced(dataSize);
3079  magick_unreferenced(data);
3080  magick_unreferenced(exception);
3081  return(MagickFalse);
3082 }
3083 
3084 MagickExport MagickBooleanType GetMagickOpenCLEnvParam(
3085  MagickCLEnv magick_unused(clEnv),MagickOpenCLEnvParam magick_unused(param),
3086  size_t magick_unused(dataSize),void *magick_unused(data),
3087  ExceptionInfo *magick_unused(exception))
3088 {
3089  magick_unreferenced(clEnv);
3090  magick_unreferenced(param);
3091  magick_unreferenced(dataSize);
3092  magick_unreferenced(data);
3093  magick_unreferenced(exception);
3094  return(MagickFalse);
3095 }
3096 
3097 MagickExport MagickBooleanType InitOpenCLEnv(MagickCLEnv magick_unused(clEnv),
3098  ExceptionInfo *magick_unused(exception))
3099 {
3100  magick_unreferenced(clEnv);
3101  magick_unreferenced(exception);
3102  return(MagickFalse);
3103 }
3104 
3105 MagickExport MagickBooleanType InitImageMagickOpenCL(
3106  ImageMagickOpenCLMode magick_unused(mode),
3107  void *magick_unused(userSelectedDevice),void *magick_unused(selectedDevice),
3108  ExceptionInfo *magick_unused(exception))
3109 {
3110  magick_unreferenced(mode);
3111  magick_unreferenced(userSelectedDevice);
3112  magick_unreferenced(selectedDevice);
3113  magick_unreferenced(exception);
3114  return(MagickFalse);
3115 }
3116 
3117 #endif /* MAGICKCORE_OPENCL_SUPPORT */
_ExceptionInfo
Definition: exception.h:102
_ImageInfo
Definition: image.h:345
_Image
Definition: image.h:133
_MagickCLEnv
Definition: opencl.c:3059