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