opencl-private.h revision be04cd4a903ac006a2d6f9607cad24aa4fe491bf 
1/*
2Copyright 1999-2016 ImageMagick Studio LLC, a non-profit organization
3dedicated to making software imaging solutions freely available.
4
5You may not use this file except in compliance with the License.
6obtain a copy of the License at
7
8http://www.imagemagick.org/script/license.php
9
10Unless required by applicable law or agreed to in writing, software
11distributed under the License is distributed on an "AS IS" BASIS,
12WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13See the License for the specific language governing permissions and
14limitations under the License.
15
16MagickCore OpenCL private methods.
17*/
18#ifndef _MAGICKCORE_OPENCL_PRIVATE_H
19#define _MAGICKCORE_OPENCL_PRIVATE_H
20
21/*
22Include declarations.
23*/
24#include "MagickCore/studio.h"
25#include "MagickCore/opencl.h"
26
27#if defined(__cplusplus) || defined(c_plusplus)
28extern "C" {
29#endif
30
31#if !defined(MAGICKCORE_OPENCL_SUPPORT)
32  typedef void* cl_context;
33  typedef void* cl_command_queue;
34  typedef void* cl_device_id;
35  typedef void* cl_event;
36  typedef void* cl_kernel;
37  typedef void* cl_mem;
38  typedef void* cl_platform_id;
39  typedef struct { unsigned char t[8]; } cl_device_type; /* 64-bit */
40#else
41
42#define MAX_COMMAND_QUEUES 16
43
44/*
45 *
46 * function pointer typedefs
47 *
48 */
49
50/* Platform APIs */
51typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetPlatformIDs)(
52                 cl_uint          num_entries,
53                 cl_platform_id * platforms,
54                 cl_uint *        num_platforms) CL_API_SUFFIX__VERSION_1_0;
55
56typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetPlatformInfo)(
57    cl_platform_id   platform,
58    cl_platform_info param_name,
59    size_t           param_value_size,
60    void *           param_value,
61    size_t *         param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
62
63/* Device APIs */
64typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetDeviceIDs)(
65    cl_platform_id   platform,
66    cl_device_type   device_type,
67    cl_uint          num_entries,
68    cl_device_id *   devices,
69    cl_uint *        num_devices) CL_API_SUFFIX__VERSION_1_0;
70
71typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetDeviceInfo)(
72    cl_device_id    device,
73    cl_device_info  param_name,
74    size_t          param_value_size,
75    void *          param_value,
76    size_t *        param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
77
78/* Context APIs */
79typedef CL_API_ENTRY cl_context (CL_API_CALL *MAGICKpfn_clCreateContext)(
80    const cl_context_properties * properties,
81    cl_uint                 num_devices,
82    const cl_device_id *    devices,
83    void (CL_CALLBACK *pfn_notify)(const char *, const void *, size_t, void *),
84    void *                  user_data,
85    cl_int *                errcode_ret) CL_API_SUFFIX__VERSION_1_0;
86
87typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseContext)(
88    cl_context context) CL_API_SUFFIX__VERSION_1_0;
89
90/* Command Queue APIs */
91typedef CL_API_ENTRY cl_command_queue (CL_API_CALL *MAGICKpfn_clCreateCommandQueue)(
92    cl_context                     context,
93    cl_device_id                   device,
94    cl_command_queue_properties    properties,
95    cl_int *                       errcode_ret) CL_API_SUFFIX__VERSION_1_0;
96
97typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseCommandQueue)(
98    cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
99
100/* Memory Object APIs */
101typedef CL_API_ENTRY cl_mem (CL_API_CALL *MAGICKpfn_clCreateBuffer)(
102    cl_context   context,
103    cl_mem_flags flags,
104    size_t       size,
105    void *       host_ptr,
106    cl_int *     errcode_ret) CL_API_SUFFIX__VERSION_1_0;
107
108typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseMemObject)(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0;
109
110/* Program Object APIs */
111typedef CL_API_ENTRY cl_program (CL_API_CALL *MAGICKpfn_clCreateProgramWithSource)(
112    cl_context        context,
113    cl_uint           count,
114    const char **     strings,
115    const size_t *    lengths,
116    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0;
117
118typedef CL_API_ENTRY cl_program (CL_API_CALL *MAGICKpfn_clCreateProgramWithBinary)(
119    cl_context                     context,
120    cl_uint                        num_devices,
121    const cl_device_id *           device_list,
122    const size_t *                 lengths,
123    const unsigned char **         binaries,
124    cl_int *                       binary_status,
125    cl_int *                       errcode_ret) CL_API_SUFFIX__VERSION_1_0;
126
127typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseProgram)(cl_program program) CL_API_SUFFIX__VERSION_1_0;
128
129typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clBuildProgram)(
130    cl_program           program,
131    cl_uint              num_devices,
132    const cl_device_id * device_list,
133    const char *         options,
134    void (CL_CALLBACK *pfn_notify)(cl_program program, void * user_data),
135    void *               user_data) CL_API_SUFFIX__VERSION_1_0;
136
137typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetProgramInfo)(
138    cl_program         program,
139    cl_program_info    param_name,
140    size_t             param_value_size,
141    void *             param_value,
142    size_t *           param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
143
144typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clGetProgramBuildInfo)(
145    cl_program            program,
146    cl_device_id          device,
147    cl_program_build_info param_name,
148    size_t                param_value_size,
149    void *                param_value,
150    size_t *              param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
151
152/* Kernel Object APIs */
153typedef CL_API_ENTRY cl_kernel (CL_API_CALL *MAGICKpfn_clCreateKernel)(
154    cl_program      program,
155    const char *    kernel_name,
156    cl_int *        errcode_ret) CL_API_SUFFIX__VERSION_1_0;
157
158typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clReleaseKernel)(cl_kernel   kernel) CL_API_SUFFIX__VERSION_1_0;
159
160typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clSetKernelArg)(
161    cl_kernel    kernel,
162    cl_uint      arg_index,
163    size_t       arg_size,
164    const void * arg_value) CL_API_SUFFIX__VERSION_1_0;
165
166/* Flush and Finish APIs */
167typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clFlush)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
168
169typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clFinish)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
170
171/* Enqueued Commands APIs */
172typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueReadBuffer)(
173    cl_command_queue    command_queue,
174    cl_mem              buffer,
175    cl_bool             blocking_read,
176    size_t              offset,
177    size_t              cb,
178    void *              ptr,
179    cl_uint             num_events_in_wait_list,
180    const cl_event *    event_wait_list,
181    cl_event *          event) CL_API_SUFFIX__VERSION_1_0;
182
183typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueWriteBuffer)(
184    cl_command_queue   command_queue,
185    cl_mem             buffer,
186    cl_bool            blocking_write,
187    size_t             offset,
188    size_t             cb,
189    const void *       ptr,
190    cl_uint            num_events_in_wait_list,
191    const cl_event *   event_wait_list,
192    cl_event *         event) CL_API_SUFFIX__VERSION_1_0;
193
194typedef CL_API_ENTRY void * (CL_API_CALL *MAGICKpfn_clEnqueueMapBuffer)(
195    cl_command_queue command_queue,
196    cl_mem           buffer,
197    cl_bool          blocking_map,
198    cl_map_flags     map_flags,
199    size_t           offset,
200    size_t           cb,
201    cl_uint          num_events_in_wait_list,
202    const cl_event * event_wait_list,
203    cl_event *       event,
204    cl_int *         errcode_ret) CL_API_SUFFIX__VERSION_1_0;
205
206typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueUnmapMemObject)(
207    cl_command_queue command_queue,
208    cl_mem           memobj,
209    void *           mapped_ptr,
210    cl_uint          num_events_in_wait_list,
211    const cl_event *  event_wait_list,
212    cl_event *        event) CL_API_SUFFIX__VERSION_1_0;
213
214typedef CL_API_ENTRY cl_int (CL_API_CALL *MAGICKpfn_clEnqueueNDRangeKernel)(
215    cl_command_queue command_queue,
216    cl_kernel        kernel,
217    cl_uint          work_dim,
218    const size_t *   global_work_offset,
219    const size_t *   global_work_size,
220    const size_t *   local_work_size,
221    cl_uint          num_events_in_wait_list,
222    const cl_event * event_wait_list,
223    cl_event *       event) CL_API_SUFFIX__VERSION_1_0;
224
225typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clGetEventProfilingInfo)(
226    cl_event event,
227    cl_profiling_info param_name,
228    size_t param_value_size,
229    void *param_value,
230    size_t *param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
231
232typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clWaitForEvents)(
233    cl_uint num_events,
234    const cl_event *event_list) CL_API_SUFFIX__VERSION_1_0;
235
236typedef CL_API_ENTRY cl_int(CL_API_CALL *MAGICKpfn_clReleaseEvent)(
237    cl_event event) CL_API_SUFFIX__VERSION_1_0;
238
239/*
240 *
241 * vendor dispatch table structure
242 *
243 * note that the types in the structure KHRicdVendorDispatch mirror the function
244 * names listed in the string table khrIcdVendorDispatchFunctionNames
245 *
246 */
247
248typedef struct MagickLibraryRec MagickLibrary;
249
250struct MagickLibraryRec
251{
252  void * base;
253
254  MAGICKpfn_clGetPlatformIDs                         clGetPlatformIDs;
255  MAGICKpfn_clGetPlatformInfo                        clGetPlatformInfo;
256  MAGICKpfn_clGetDeviceIDs                           clGetDeviceIDs;
257  MAGICKpfn_clGetDeviceInfo                          clGetDeviceInfo;
258  MAGICKpfn_clCreateContext                          clCreateContext;
259  MAGICKpfn_clCreateCommandQueue                     clCreateCommandQueue;
260  MAGICKpfn_clReleaseCommandQueue                    clReleaseCommandQueue;
261  MAGICKpfn_clCreateBuffer                           clCreateBuffer;
262  MAGICKpfn_clReleaseMemObject                       clReleaseMemObject;
263  MAGICKpfn_clCreateProgramWithSource                clCreateProgramWithSource;
264  MAGICKpfn_clCreateProgramWithBinary                clCreateProgramWithBinary;
265  MAGICKpfn_clReleaseProgram                         clReleaseProgram;
266  MAGICKpfn_clBuildProgram                           clBuildProgram;
267  MAGICKpfn_clGetProgramInfo                         clGetProgramInfo;
268  MAGICKpfn_clGetProgramBuildInfo                    clGetProgramBuildInfo;
269  MAGICKpfn_clCreateKernel                           clCreateKernel;
270  MAGICKpfn_clReleaseKernel                          clReleaseKernel;
271  MAGICKpfn_clSetKernelArg                           clSetKernelArg;
272  MAGICKpfn_clFlush                                  clFlush;
273  MAGICKpfn_clFinish                                 clFinish;
274  MAGICKpfn_clEnqueueReadBuffer                      clEnqueueReadBuffer;
275  MAGICKpfn_clEnqueueWriteBuffer                     clEnqueueWriteBuffer;
276  MAGICKpfn_clEnqueueMapBuffer                       clEnqueueMapBuffer;
277  MAGICKpfn_clEnqueueUnmapMemObject                  clEnqueueUnmapMemObject;
278  MAGICKpfn_clEnqueueNDRangeKernel                   clEnqueueNDRangeKernel;
279  MAGICKpfn_clGetEventProfilingInfo                  clGetEventProfilingInfo;
280  MAGICKpfn_clWaitForEvents                          clWaitForEvents;
281  MAGICKpfn_clReleaseEvent                           clReleaseEvent;
282};
283
284struct _MagickCLEnv {
285  MagickBooleanType OpenCLInitialized;  /* whether OpenCL environment is initialized. */
286  MagickBooleanType OpenCLDisabled;	/* whether if OpenCL has been explicitely disabled. */
287
288  MagickLibrary * library;
289
290  /*OpenCL objects */
291  cl_platform_id platform;
292  cl_device_type deviceType;
293  cl_device_id device;
294  cl_context context;
295
296  MagickBooleanType disableProgramCache; /* disable the OpenCL program cache */
297  cl_program programs[MAGICK_OPENCL_NUM_PROGRAMS]; /* one program object maps one kernel source file */
298
299  MagickBooleanType regenerateProfile;   /* re-run the microbenchmark in auto device selection mode */
300
301  /* Random number generator seeds */
302  unsigned int numGenerators;
303  float randNormalize;
304  cl_mem seeds;
305  SemaphoreInfo* seedsLock;
306
307  SemaphoreInfo* lock;
308
309  cl_command_queue commandQueues[MAX_COMMAND_QUEUES];
310  ssize_t commandQueuesPos;
311  SemaphoreInfo* commandQueuesLock;
312};
313
314#endif
315
316#if defined(MAGICKCORE_HDRI_SUPPORT)
317#define CLOptions "-cl-single-precision-constant -cl-mad-enable -DMAGICKCORE_HDRI_SUPPORT=1 "\
318  "-DCLQuantum=float -DCLSignedQuantum=float -DCLPixelType=float4 -DQuantumRange=%f " \
319  "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
320  "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
321#define CLQuantum  cl_float
322#define CLPixelPacket  cl_float4
323#define CLCharQuantumScale 1.0f
324#elif (MAGICKCORE_QUANTUM_DEPTH == 8)
325#define CLOptions "-cl-single-precision-constant -cl-mad-enable " \
326  "-DCLQuantum=uchar -DCLSignedQuantum=char -DCLPixelType=uchar4 -DQuantumRange=%ff " \
327  "-DQuantumScale=%ff -DCharQuantumScale=%ff -DMagickEpsilon=%ff -DMagickPI=%ff "\
328  "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
329#define CLQuantum  cl_uchar
330#define CLPixelPacket  cl_uchar4
331#define CLCharQuantumScale 1.0f
332#elif (MAGICKCORE_QUANTUM_DEPTH == 16)
333#define CLOptions "-cl-single-precision-constant -cl-mad-enable " \
334  "-DCLQuantum=ushort -DCLSignedQuantum=short -DCLPixelType=ushort4 -DQuantumRange=%f "\
335  "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
336  "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
337#define CLQuantum  cl_ushort
338#define CLPixelPacket  cl_ushort4
339#define CLCharQuantumScale 257.0f
340#elif (MAGICKCORE_QUANTUM_DEPTH == 32)
341#define CLOptions "-cl-single-precision-constant -cl-mad-enable " \
342  "-DCLQuantum=uint -DCLSignedQuantum=int -DCLPixelType=uint4 -DQuantumRange=%f "\
343  "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
344  "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
345#define CLQuantum  cl_uint
346#define CLPixelPacket  cl_uint4
347#define CLCharQuantumScale 16843009.0f
348#elif (MAGICKCORE_QUANTUM_DEPTH == 64)
349#define CLOptions "-cl-single-precision-constant -cl-mad-enable " \
350  "-DCLQuantum=ulong -DCLSignedQuantum=long -DCLPixelType=ulong4 -DQuantumRange=%f "\
351  "-DQuantumScale=%f -DCharQuantumScale=%f -DMagickEpsilon=%f -DMagickPI=%f "\
352  "-DMaxMap=%u -DMAGICKCORE_QUANTUM_DEPTH=%u"
353#define CLQuantum  cl_ulong
354#define CLPixelPacket  cl_ulong4
355#define CLCharQuantumScale 72340172838076673.0f
356#endif
357
358typedef enum {
359  AddNoiseKernel,
360  BlurRowKernel,
361  BlurColumnKernel,
362  CompositeKernel,
363  ComputeFunctionKernel,
364  ContrastKernel,
365  ContrastStretchKernel,
366  ConvolveKernel,
367  EqualizeKernel,
368  GrayScaleKernel,
369  HistogramKernel,
370  HullPass1Kernel,
371  HullPass2Kernel,
372  LocalContrastBlurRowKernel,
373  LocalContrastBlurApplyColumnKernel,
374  ModulateKernel,
375  MotionBlurKernel,
376  RandomNumberGeneratorKernel,
377  ResizeHorizontalKernel,
378  ResizeVerticalKernel,
379  RotationalBlurKernel,
380  UnsharpMaskBlurColumnKernel,
381  UnsharpMaskKernel,
382  WaveletDenoiseKernel,
383  KERNEL_COUNT
384} ProfiledKernels;
385
386extern MagickPrivate cl_context
387  GetOpenCLContext(MagickCLEnv);
388
389extern MagickPrivate cl_kernel
390  AcquireOpenCLKernel(MagickCLEnv, MagickOpenCLProgram, const char*);
391
392extern MagickPrivate cl_command_queue
393  AcquireOpenCLCommandQueue(MagickCLEnv);
394
395extern MagickPrivate MagickBooleanType
396  OpenCLThrowMagickException(ExceptionInfo *,
397    const char *,const char *,const size_t,
398    const ExceptionType,const char *,const char *,...),
399  RelinquishOpenCLCommandQueue(MagickCLEnv, cl_command_queue),
400  RelinquishOpenCLKernel(MagickCLEnv, cl_kernel);
401
402extern MagickPrivate unsigned long
403  GetOpenCLDeviceLocalMemorySize(MagickCLEnv),
404  GetOpenCLDeviceMaxMemAllocSize(MagickCLEnv);
405
406extern MagickPrivate const char*
407  GetOpenCLCachedFilesDirectory();
408
409extern MagickPrivate void
410  OpenCLLog(const char*),
411  UnlockRandSeedBuffer(MagickCLEnv);
412
413extern MagickPrivate cl_mem
414  GetAndLockRandSeedBuffer(MagickCLEnv);
415
416extern MagickPrivate unsigned int
417  GetNumRandGenerators(MagickCLEnv);
418
419extern MagickPrivate float
420  GetRandNormalize(MagickCLEnv);
421
422extern MagickPrivate void
423  OpenCLTerminus(),
424  RecordProfileData(MagickCLEnv,ProfiledKernels,cl_event);
425
426/* #define OPENCLLOG_ENABLED 1 */
427static inline void OpenCLLogException(const char* function,
428                        const unsigned int line,
429                        ExceptionInfo* exception) {
430#ifdef OPENCLLOG_ENABLED
431  if (exception->severity!=0) {
432    char message[MagickPathExtent];
433    /*  dump the source into a file */
434    (void) FormatLocaleString(message,MagickPathExtent,"%s:%d Exception(%d):%s "
435        ,function,line,exception->severity,exception->reason);
436    OpenCLLog(message);
437  }
438#else
439  magick_unreferenced(function);
440  magick_unreferenced(line);
441  magick_unreferenced(exception);
442#endif
443}
444
445
446#if defined(__cplusplus) || defined(c_plusplus)
447}
448#endif
449
450#endif
451