Blender  V3.3
kernel_templates.h
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2021-2022 Intel Corporation */
3 
4 #pragma once
5 
6 /* Some macro magic to generate templates for kernel arguments.
7  * The resulting oneapi_call() template allows to call a SYCL/C++ kernel
8  * with typed arguments by only giving it a void `**args` as given by Cycles.
9  * The template will automatically cast from void* to the expected type. */
10 
11 /* When expanded by the preprocessor, the generated templates will look like this example: */
12 #if 0
13 template<typename T0, typename T1, typename T2>
14 void oneapi_call(
15  KernelGlobalsGPU *kg,
16  sycl::handler &cgh,
17  size_t global_size,
18  size_t local_size,
19  void **args,
20  void (*func)(const KernelGlobalsGPU *, size_t, size_t, sycl::handler &, T0, T1, T2))
21 {
22  func(kg, global_size, local_size, cgh, *(T0 *)(args[0]), *(T1 *)(args[1]), *(T2 *)(args[2]));
23 }
24 #endif
25 
26 /* clang-format off */
27 #define ONEAPI_TYP(x) typename T##x
28 #define ONEAPI_CAST(x) *(T##x *)(args[x])
29 #define ONEAPI_T(x) T##x
30 
31 #define ONEAPI_GET_NTH_ARG(_1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15, _16, _17, _18, _19, _20, _21, _22, N, ...) N
32 #define ONEAPI_0(_call, ...)
33 #define ONEAPI_1(_call, x) _call(x)
34 #define ONEAPI_2(_call, x, ...) _call(x), ONEAPI_1(_call, __VA_ARGS__)
35 #define ONEAPI_3(_call, x, ...) _call(x), ONEAPI_2(_call, __VA_ARGS__)
36 #define ONEAPI_4(_call, x, ...) _call(x), ONEAPI_3(_call, __VA_ARGS__)
37 #define ONEAPI_5(_call, x, ...) _call(x), ONEAPI_4(_call, __VA_ARGS__)
38 #define ONEAPI_6(_call, x, ...) _call(x), ONEAPI_5(_call, __VA_ARGS__)
39 #define ONEAPI_7(_call, x, ...) _call(x), ONEAPI_6(_call, __VA_ARGS__)
40 #define ONEAPI_8(_call, x, ...) _call(x), ONEAPI_7(_call, __VA_ARGS__)
41 #define ONEAPI_9(_call, x, ...) _call(x), ONEAPI_8(_call, __VA_ARGS__)
42 #define ONEAPI_10(_call, x, ...) _call(x), ONEAPI_9(_call, __VA_ARGS__)
43 #define ONEAPI_11(_call, x, ...) _call(x), ONEAPI_10(_call, __VA_ARGS__)
44 #define ONEAPI_12(_call, x, ...) _call(x), ONEAPI_11(_call, __VA_ARGS__)
45 #define ONEAPI_13(_call, x, ...) _call(x), ONEAPI_12(_call, __VA_ARGS__)
46 #define ONEAPI_14(_call, x, ...) _call(x), ONEAPI_13(_call, __VA_ARGS__)
47 #define ONEAPI_15(_call, x, ...) _call(x), ONEAPI_14(_call, __VA_ARGS__)
48 #define ONEAPI_16(_call, x, ...) _call(x), ONEAPI_15(_call, __VA_ARGS__)
49 #define ONEAPI_17(_call, x, ...) _call(x), ONEAPI_16(_call, __VA_ARGS__)
50 #define ONEAPI_18(_call, x, ...) _call(x), ONEAPI_17(_call, __VA_ARGS__)
51 #define ONEAPI_19(_call, x, ...) _call(x), ONEAPI_18(_call, __VA_ARGS__)
52 #define ONEAPI_20(_call, x, ...) _call(x), ONEAPI_19(_call, __VA_ARGS__)
53 #define ONEAPI_21(_call, x, ...) _call(x), ONEAPI_20(_call, __VA_ARGS__)
54 
55 #define ONEAPI_CALL_FOR(x, ...) \
56  ONEAPI_GET_NTH_ARG("ignored", \
57  ##__VA_ARGS__, \
58  ONEAPI_21, \
59  ONEAPI_20, \
60  ONEAPI_19, \
61  ONEAPI_18, \
62  ONEAPI_17, \
63  ONEAPI_16, \
64  ONEAPI_15, \
65  ONEAPI_14, \
66  ONEAPI_13, \
67  ONEAPI_12, \
68  ONEAPI_11, \
69  ONEAPI_10, \
70  ONEAPI_9, \
71  ONEAPI_8, \
72  ONEAPI_7, \
73  ONEAPI_6, \
74  ONEAPI_5, \
75  ONEAPI_4, \
76  ONEAPI_3, \
77  ONEAPI_2, \
78  ONEAPI_1, \
79  ONEAPI_0) \
80  (x, ##__VA_ARGS__)
81 
82 /* This template automatically casts entries in the void **args array to the types requested by the kernel func.
83  Since kernel parameters are passed as void ** to the device, this is the closest that we have to type safety. */
84 #define oneapi_template(...) \
85  template<ONEAPI_CALL_FOR(ONEAPI_TYP, __VA_ARGS__)> \
86  void oneapi_call( \
87  KernelGlobalsGPU *kg, \
88  sycl::handler &cgh, \
89  size_t global_size, \
90  size_t local_size, \
91  void **args, \
92  void (*func)(KernelGlobalsGPU*, size_t, size_t, sycl::handler &, ONEAPI_CALL_FOR(ONEAPI_T, __VA_ARGS__))) \
93  { \
94  func(kg, \
95  global_size, \
96  local_size, \
97  cgh, \
98  ONEAPI_CALL_FOR(ONEAPI_CAST, __VA_ARGS__)); \
99  }
100 
102 oneapi_template(0, 1)
104 oneapi_template(0, 1, 2, 3)
105 oneapi_template(0, 1, 2, 3, 4)
106 oneapi_template(0, 1, 2, 3, 4, 5)
107 oneapi_template(0, 1, 2, 3, 4, 5, 6)
108 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7)
109 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8)
110 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9)
111 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10)
112 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11)
113 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12)
114 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13)
115 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14)
116 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15)
117 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16)
118 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17)
119 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18)
120 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19)
121 oneapi_template(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20)
122 
123  /* clang-format on */
SyclQueue void void size_t num_bytes SyclQueue void const char void *memory_device_pointer KernelContext int size_t global_size
#define oneapi_template(...)
#define T2
Definition: md5.cpp:18
#define T1
Definition: md5.cpp:17