Skip to content

Commit 073a36b

Browse files
author
Alexander Batashev
authored
[SYCL] Add missing builtins for vec<X, 1> (#1984)
On host using builtins with `vec<X, 1>` types causes linkage errors. Add missing overloads. Signed-off-by: Alexander Batashev <[email protected]>
1 parent 7c15695 commit 073a36b

File tree

4 files changed

+618
-22
lines changed

4 files changed

+618
-22
lines changed

‎sycl/source/detail/builtins_geometric.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,16 +19,20 @@ namespace d = s::detail;
1919
__SYCL_INLINE_NAMESPACE(cl) {
2020
namespace __host_std {
2121

22+
__SYCL_EXPORT s::cl_float Dot(s::vec<float, 1>, s::vec<float, 1>);
2223
__SYCL_EXPORT s::cl_float Dot(s::cl_float2, s::cl_float2);
2324
__SYCL_EXPORT s::cl_float Dot(s::cl_float3, s::cl_float3);
2425
__SYCL_EXPORT s::cl_float Dot(s::cl_float4, s::cl_float4);
26+
__SYCL_EXPORT s::cl_double Dot(s::vec<double, 1>, s::vec<double, 1>);
2527
__SYCL_EXPORT s::cl_double Dot(s::cl_double2, s::cl_double2);
2628
__SYCL_EXPORT s::cl_double Dot(s::cl_double3, s::cl_double3);
2729
__SYCL_EXPORT s::cl_double Dot(s::cl_double4, s::cl_double4);
30+
__SYCL_EXPORT s::cl_half Dot(s::vec<half, 1>, s::vec<half, 1>);
2831
__SYCL_EXPORT s::cl_half Dot(s::cl_half2, s::cl_half2);
2932
__SYCL_EXPORT s::cl_half Dot(s::cl_half3, s::cl_half3);
3033
__SYCL_EXPORT s::cl_half Dot(s::cl_half4, s::cl_half4);
3134

35+
__SYCL_EXPORT s::cl_int All(s::vec<int, 1>);
3236
__SYCL_EXPORT s::cl_int All(s::cl_int2);
3337
__SYCL_EXPORT s::cl_int All(s::cl_int3);
3438
__SYCL_EXPORT s::cl_int All(s::cl_int4);
@@ -144,12 +148,15 @@ MAKE_GEO_1V_2V_RS(Dot, __FMul_impl, s::cl_half, s::cl_half, s::cl_half)
144148
__SYCL_EXPORT s::cl_float length(s::cl_float p) { return __length(p); }
145149
__SYCL_EXPORT s::cl_double length(s::cl_double p) { return __length(p); }
146150
__SYCL_EXPORT s::cl_half length(s::cl_half p) { return __length(p); }
151+
__SYCL_EXPORT s::cl_float length(s::vec<float, 1> p) { return __length(p); }
147152
__SYCL_EXPORT s::cl_float length(s::cl_float2 p) { return __length(p); }
148153
__SYCL_EXPORT s::cl_float length(s::cl_float3 p) { return __length(p); }
149154
__SYCL_EXPORT s::cl_float length(s::cl_float4 p) { return __length(p); }
155+
__SYCL_EXPORT s::cl_double length(s::vec<double, 1> p) { return __length(p); }
150156
__SYCL_EXPORT s::cl_double length(s::cl_double2 p) { return __length(p); }
151157
__SYCL_EXPORT s::cl_double length(s::cl_double3 p) { return __length(p); }
152158
__SYCL_EXPORT s::cl_double length(s::cl_double4 p) { return __length(p); }
159+
__SYCL_EXPORT s::cl_half length(s::vec<half, 1> p) { return __length(p); }
153160
__SYCL_EXPORT s::cl_half length(s::cl_half2 p) { return __length(p); }
154161
__SYCL_EXPORT s::cl_half length(s::cl_half3 p) { return __length(p); }
155162
__SYCL_EXPORT s::cl_half length(s::cl_half4 p) { return __length(p); }
@@ -158,6 +165,9 @@ __SYCL_EXPORT s::cl_half length(s::cl_half4 p) { return __length(p); }
158165
__SYCL_EXPORT s::cl_float distance(s::cl_float p0, s::cl_float p1) {
159166
return length(p0 - p1);
160167
}
168+
__SYCL_EXPORT s::cl_float distance(s::vec<float, 1> p0, s::vec<float, 1> p1) {
169+
return length(p0 - p1);
170+
}
161171
__SYCL_EXPORT s::cl_float distance(s::cl_float2 p0, s::cl_float2 p1) {
162172
return length(p0 - p1);
163173
}
@@ -170,6 +180,9 @@ __SYCL_EXPORT s::cl_float distance(s::cl_float4 p0, s::cl_float4 p1) {
170180
__SYCL_EXPORT s::cl_double distance(s::cl_double p0, s::cl_double p1) {
171181
return length(p0 - p1);
172182
}
183+
__SYCL_EXPORT s::cl_float distance(s::vec<double, 1> p0, s::vec<double, 1> p1) {
184+
return length(p0 - p1);
185+
}
173186
__SYCL_EXPORT s::cl_double distance(s::cl_double2 p0, s::cl_double2 p1) {
174187
return length(p0 - p1);
175188
}
@@ -182,6 +195,9 @@ __SYCL_EXPORT s::cl_double distance(s::cl_double4 p0, s::cl_double4 p1) {
182195
__SYCL_EXPORT s::cl_half distance(s::cl_half p0, s::cl_half p1) {
183196
return length(p0 - p1);
184197
}
198+
__SYCL_EXPORT s::cl_float distance(s::vec<half, 1> p0, s::vec<half, 1> p1) {
199+
return length(p0 - p1);
200+
}
185201
__SYCL_EXPORT s::cl_half distance(s::cl_half2 p0, s::cl_half2 p1) {
186202
return length(p0 - p1);
187203
}
@@ -194,10 +210,16 @@ __SYCL_EXPORT s::cl_half distance(s::cl_half4 p0, s::cl_half4 p1) {
194210

195211
// normalize
196212
__SYCL_EXPORT s::cl_float normalize(s::cl_float p) { return __normalize(p); }
213+
__SYCL_EXPORT s::cl_float normalize(s::vec<float, 1> p) {
214+
return __normalize(p);
215+
}
197216
__SYCL_EXPORT s::cl_float2 normalize(s::cl_float2 p) { return __normalize(p); }
198217
__SYCL_EXPORT s::cl_float3 normalize(s::cl_float3 p) { return __normalize(p); }
199218
__SYCL_EXPORT s::cl_float4 normalize(s::cl_float4 p) { return __normalize(p); }
200219
__SYCL_EXPORT s::cl_double normalize(s::cl_double p) { return __normalize(p); }
220+
__SYCL_EXPORT s::cl_double normalize(s::vec<double, 1> p) {
221+
return __normalize(p);
222+
}
201223
__SYCL_EXPORT s::cl_double2 normalize(s::cl_double2 p) {
202224
return __normalize(p);
203225
}
@@ -216,6 +238,9 @@ __SYCL_EXPORT s::cl_half4 normalize(s::cl_half4 p) { return __normalize(p); }
216238
__SYCL_EXPORT s::cl_float fast_length(s::cl_float p) {
217239
return __fast_length(p);
218240
}
241+
__SYCL_EXPORT s::cl_float fast_length(s::vec<float, 1> p) {
242+
return __fast_length(p);
243+
}
219244
__SYCL_EXPORT s::cl_float fast_length(s::cl_float2 p) {
220245
return __fast_length(p);
221246
}
@@ -233,6 +258,9 @@ __SYCL_EXPORT s::cl_float fast_normalize(s::cl_float p) {
233258
s::cl_float r = std::sqrt(FMul(p, p));
234259
return p / r;
235260
}
261+
__SYCL_EXPORT s::cl_float fast_normalize(s::vec<float, 1> p) {
262+
return __fast_normalize(p);
263+
}
236264
__SYCL_EXPORT s::cl_float2 fast_normalize(s::cl_float2 p) {
237265
return __fast_normalize(p);
238266
}
@@ -247,6 +275,10 @@ __SYCL_EXPORT s::cl_float4 fast_normalize(s::cl_float4 p) {
247275
__SYCL_EXPORT s::cl_float fast_distance(s::cl_float p0, s::cl_float p1) {
248276
return fast_length(p0 - p1);
249277
}
278+
__SYCL_EXPORT s::cl_float fast_distance(s::vec<float, 1> p0,
279+
s::vec<float, 1> p1) {
280+
return fast_length(p0 - p1);
281+
}
250282
__SYCL_EXPORT s::cl_float fast_distance(s::cl_float2 p0, s::cl_float2 p1) {
251283
return fast_length(p0 - p1);
252284
}

‎sycl/source/detail/builtins_helper.hpp

Lines changed: 42 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -16,23 +16,24 @@
1616
#define __NOEXC /*noexcept*/
1717

1818
#define __MAKE_1V(Fun, Call, N, Ret, Arg1) \
19-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x) { \
20-
Ret##N r; \
19+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x) { \
20+
sycl::vec<Ret, N> r; \
2121
detail::helper<N - 1>().run_1v( \
2222
r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \
2323
return r; \
2424
}
2525

2626
#define __MAKE_1V_2V(Fun, Call, N, Ret, Arg1, Arg2) \
27-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y) { \
28-
Ret##N r; \
27+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, \
28+
sycl::vec<Arg2, N> y) { \
29+
sycl::vec<Ret, N> r; \
2930
detail::helper<N - 1>().run_1v_2v( \
3031
r, [](Arg1 x, Arg2 y) { return cl::__host_std::Call(x, y); }, x, y); \
3132
return r; \
3233
}
3334

3435
#define __MAKE_1V_2V_RS(Fun, Call, N, Ret, Arg1, Arg2) \
35-
__SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x, Arg2##N y) { \
36+
__SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x, sycl::vec<Arg2, N> y) { \
3637
Ret r = Ret(); \
3738
detail::helper<N - 1>().run_1v_2v_rs( \
3839
r, \
@@ -42,16 +43,17 @@
4243
}
4344

4445
#define __MAKE_1V_RS(Fun, Call, N, Ret, Arg1) \
45-
__SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x) { \
46+
__SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x) { \
4647
Ret r = Ret(); \
4748
detail::helper<N - 1>().run_1v_rs( \
4849
r, [](Ret &r, Arg1 x) { return cl::__host_std::Call(r, x); }, x); \
4950
return r; \
5051
}
5152

5253
#define __MAKE_1V_2V_3V(Fun, Call, N, Ret, Arg1, Arg2, Arg3) \
53-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y, Arg3##N z) { \
54-
Ret##N r; \
54+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC( \
55+
sycl::vec<Arg1, N> x, sycl::vec<Arg2, N> y, sycl::vec<Arg3, N> z) { \
56+
sycl::vec<Ret, N> r; \
5557
detail::helper<N - 1>().run_1v_2v_3v( \
5658
r, \
5759
[](Arg1 x, Arg2 y, Arg3 z) { return cl::__host_std::Call(x, y, z); }, \
@@ -60,8 +62,9 @@
6062
}
6163

6264
#define __MAKE_1V_2S_3S(Fun, N, Ret, Arg1, Arg2, Arg3) \
63-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2 y, Arg3 z) { \
64-
Ret##N r; \
65+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, Arg2 y, \
66+
Arg3 z) { \
67+
sycl::vec<Ret, N> r; \
6568
detail::helper<N - 1>().run_1v_2s_3s( \
6669
r, \
6770
[](Arg1 x, Arg2 y, Arg3 z) { return cl::__host_std::Fun(x, y, z); }, \
@@ -70,40 +73,42 @@
7073
}
7174

7275
#define __MAKE_1V_2S(Fun, N, Ret, Arg1, Arg2) \
73-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2 y) { \
74-
Ret##N r; \
76+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, Arg2 y) { \
77+
sycl::vec<Ret, N> r; \
7578
detail::helper<N - 1>().run_1v_2s( \
7679
r, [](Arg1 x, Arg2 y) { return cl::__host_std::Fun(x, y); }, x, y); \
7780
return r; \
7881
}
7982

8083
#define __MAKE_SR_1V_AND(Fun, Call, N, Ret, Arg1) \
81-
__SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x) { \
84+
__SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x) { \
8285
Ret r; \
8386
detail::helper<N - 1>().run_1v_sr_and( \
8487
r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \
8588
return r; \
8689
}
8790

8891
#define __MAKE_SR_1V_OR(Fun, Call, N, Ret, Arg1) \
89-
__SYCL_EXPORT Ret Fun __NOEXC(Arg1##N x) { \
92+
__SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x) { \
9093
Ret r; \
9194
detail::helper<N - 1>().run_1v_sr_or( \
9295
r, [](Arg1 x) { return cl::__host_std::Call(x); }, x); \
9396
return r; \
9497
}
9598

9699
#define __MAKE_1V_2P(Fun, N, Ret, Arg1, Arg2) \
97-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N *y) { \
98-
Ret##N r; \
100+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, \
101+
sycl::vec<Arg2, N> *y) { \
102+
sycl::vec<Ret, N> r; \
99103
detail::helper<N - 1>().run_1v_2p( \
100104
r, [](Arg1 x, Arg2 *y) { return cl::__host_std::Fun(x, y); }, x, y); \
101105
return r; \
102106
}
103107

104108
#define __MAKE_1V_2V_3P(Fun, N, Ret, Arg1, Arg2, Arg3) \
105-
__SYCL_EXPORT Ret##N Fun __NOEXC(Arg1##N x, Arg2##N y, Arg3##N *z) { \
106-
Ret##N r; \
109+
__SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC( \
110+
sycl::vec<Arg1, N> x, sycl::vec<Arg2, N> y, sycl::vec<Arg3, N> *z) { \
111+
sycl::vec<Ret, N> r; \
107112
detail::helper<N - 1>().run_1v_2v_3p( \
108113
r, \
109114
[](Arg1 x, Arg2 y, Arg3 *z) { return cl::__host_std::Fun(x, y, z); }, \
@@ -114,15 +119,18 @@
114119
#define MAKE_1V(Fun, Ret, Arg1) MAKE_1V_FUNC(Fun, Fun, Ret, Arg1)
115120

116121
#define MAKE_1V_FUNC(Fun, Call, Ret, Arg1) \
122+
__MAKE_1V(Fun, Call, 1, Ret, Arg1) \
117123
__MAKE_1V(Fun, Call, 2, Ret, Arg1) \
118124
__MAKE_1V(Fun, Call, 3, Ret, Arg1) \
119125
__MAKE_1V(Fun, Call, 4, Ret, Arg1) \
120-
__MAKE_1V(Fun, Call, 8, Ret, Arg1) __MAKE_1V(Fun, Call, 16, Ret, Arg1)
126+
__MAKE_1V(Fun, Call, 8, Ret, Arg1) \
127+
__MAKE_1V(Fun, Call, 16, Ret, Arg1)
121128

122129
#define MAKE_1V_2V(Fun, Ret, Arg1, Arg2) \
123130
MAKE_1V_2V_FUNC(Fun, Fun, Ret, Arg1, Arg2)
124131

125132
#define MAKE_1V_2V_FUNC(Fun, Call, Ret, Arg1, Arg2) \
133+
__MAKE_1V_2V(Fun, Call, 1, Ret, Arg1, Arg2) \
126134
__MAKE_1V_2V(Fun, Call, 2, Ret, Arg1, Arg2) \
127135
__MAKE_1V_2V(Fun, Call, 3, Ret, Arg1, Arg2) \
128136
__MAKE_1V_2V(Fun, Call, 4, Ret, Arg1, Arg2) \
@@ -133,6 +141,7 @@
133141
MAKE_1V_2V_3V_FUNC(Fun, Fun, Ret, Arg1, Arg2, Arg3)
134142

135143
#define MAKE_1V_2V_3V_FUNC(Fun, Call, Ret, Arg1, Arg2, Arg3) \
144+
__MAKE_1V_2V_3V(Fun, Call, 1, Ret, Arg1, Arg2, Arg3) \
136145
__MAKE_1V_2V_3V(Fun, Call, 2, Ret, Arg1, Arg2, Arg3) \
137146
__MAKE_1V_2V_3V(Fun, Call, 3, Ret, Arg1, Arg2, Arg3) \
138147
__MAKE_1V_2V_3V(Fun, Call, 4, Ret, Arg1, Arg2, Arg3) \
@@ -153,44 +162,55 @@
153162
}
154163

155164
#define MAKE_1V_2S(Fun, Ret, Arg1, Arg2) \
165+
__MAKE_1V_2S(Fun, 1, Ret, Arg1, Arg2) \
156166
__MAKE_1V_2S(Fun, 2, Ret, Arg1, Arg2) \
157167
__MAKE_1V_2S(Fun, 3, Ret, Arg1, Arg2) \
158168
__MAKE_1V_2S(Fun, 4, Ret, Arg1, Arg2) \
159-
__MAKE_1V_2S(Fun, 8, Ret, Arg1, Arg2) __MAKE_1V_2S(Fun, 16, Ret, Arg1, Arg2)
169+
__MAKE_1V_2S(Fun, 8, Ret, Arg1, Arg2) \
170+
__MAKE_1V_2S(Fun, 16, Ret, Arg1, Arg2)
160171

161172
#define MAKE_1V_2S_3S(Fun, Ret, Arg1, Arg2, Arg3) \
173+
__MAKE_1V_2S_3S(Fun, 1, Ret, Arg1, Arg2, Arg3) \
162174
__MAKE_1V_2S_3S(Fun, 2, Ret, Arg1, Arg2, Arg3) \
163175
__MAKE_1V_2S_3S(Fun, 3, Ret, Arg1, Arg2, Arg3) \
164176
__MAKE_1V_2S_3S(Fun, 4, Ret, Arg1, Arg2, Arg3) \
165177
__MAKE_1V_2S_3S(Fun, 8, Ret, Arg1, Arg2, Arg3) \
166178
__MAKE_1V_2S_3S(Fun, 16, Ret, Arg1, Arg2, Arg3)
167179

168180
#define MAKE_SR_1V_AND(Fun, Call, Ret, Arg1) \
181+
__MAKE_SR_1V_AND(Fun, Call, 1, Ret, Arg1) \
169182
__MAKE_SR_1V_AND(Fun, Call, 2, Ret, Arg1) \
170183
__MAKE_SR_1V_AND(Fun, Call, 3, Ret, Arg1) \
171184
__MAKE_SR_1V_AND(Fun, Call, 4, Ret, Arg1) \
172185
__MAKE_SR_1V_AND(Fun, Call, 8, Ret, Arg1) \
173186
__MAKE_SR_1V_AND(Fun, Call, 16, Ret, Arg1)
174187

175188
#define MAKE_SR_1V_OR(Fun, Call, Ret, Arg1) \
189+
__MAKE_SR_1V_OR(Fun, Call, 1, Ret, Arg1) \
176190
__MAKE_SR_1V_OR(Fun, Call, 2, Ret, Arg1) \
177191
__MAKE_SR_1V_OR(Fun, Call, 3, Ret, Arg1) \
178192
__MAKE_SR_1V_OR(Fun, Call, 4, Ret, Arg1) \
179193
__MAKE_SR_1V_OR(Fun, Call, 8, Ret, Arg1) \
180194
__MAKE_SR_1V_OR(Fun, Call, 16, Ret, Arg1)
181195

182196
#define MAKE_1V_2P(Fun, Ret, Arg1, Arg2) \
197+
__MAKE_1V_2P(Fun, 1, Ret, Arg1, Arg2) \
183198
__MAKE_1V_2P(Fun, 2, Ret, Arg1, Arg2) \
184199
__MAKE_1V_2P(Fun, 3, Ret, Arg1, Arg2) \
185200
__MAKE_1V_2P(Fun, 4, Ret, Arg1, Arg2) \
186-
__MAKE_1V_2P(Fun, 8, Ret, Arg1, Arg2) __MAKE_1V_2P(Fun, 16, Ret, Arg1, Arg2)
201+
__MAKE_1V_2P(Fun, 8, Ret, Arg1, Arg2) \
202+
__MAKE_1V_2P(Fun, 16, Ret, Arg1, Arg2)
187203

188204
#define MAKE_GEO_1V_2V_RS(Fun, Call, Ret, Arg1, Arg2) \
205+
__MAKE_1V_2V_RS(Fun, Call, 1, Ret, Arg1, Arg2) \
189206
__MAKE_1V_2V_RS(Fun, Call, 2, Ret, Arg1, Arg2) \
190207
__MAKE_1V_2V_RS(Fun, Call, 3, Ret, Arg1, Arg2) \
191-
__MAKE_1V_2V_RS(Fun, Call, 4, Ret, Arg1, Arg2)
208+
__MAKE_1V_2V_RS(Fun, Call, 4, Ret, Arg1, Arg2) \
209+
__MAKE_1V_2V_RS(Fun, Call, 8, Ret, Arg1, Arg2) \
210+
__MAKE_1V_2V_RS(Fun, Call, 16, Ret, Arg1, Arg2)
192211

193212
#define MAKE_1V_2V_3P(Fun, Ret, Arg1, Arg2, Arg3) \
213+
__MAKE_1V_2V_3P(Fun, 1, Ret, Arg1, Arg2, Arg3) \
194214
__MAKE_1V_2V_3P(Fun, 2, Ret, Arg1, Arg2, Arg3) \
195215
__MAKE_1V_2V_3P(Fun, 3, Ret, Arg1, Arg2, Arg3) \
196216
__MAKE_1V_2V_3P(Fun, 4, Ret, Arg1, Arg2, Arg3) \

0 commit comments

Comments
 (0)