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