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