DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins_geometric.cpp
Go to the documentation of this file.
1 //==--------- builtins_geometric.cpp - SYCL built-in geometric functions ---==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 // This file defines the host versions of functions defined
10 // in SYCL SPEC section - 4.13.6 Geometric functions.
11 
12 #include "builtins_helper.hpp"
14 
15 #include <cmath>
16 
17 namespace s = sycl;
18 namespace d = s::detail;
19 
20 namespace __host_std {
21 
22 __SYCL_EXPORT s::cl_float sycl_host_Dot(s::vec<float, 1>, s::vec<float, 1>);
23 __SYCL_EXPORT s::cl_float sycl_host_Dot(s::cl_float2, s::cl_float2);
24 __SYCL_EXPORT s::cl_float sycl_host_Dot(s::cl_float3, s::cl_float3);
25 __SYCL_EXPORT s::cl_float sycl_host_Dot(s::cl_float4, s::cl_float4);
26 __SYCL_EXPORT s::cl_double sycl_host_Dot(s::vec<double, 1>, s::vec<double, 1>);
27 __SYCL_EXPORT s::cl_double sycl_host_Dot(s::cl_double2, s::cl_double2);
28 __SYCL_EXPORT s::cl_double sycl_host_Dot(s::cl_double3, s::cl_double3);
29 __SYCL_EXPORT s::cl_double sycl_host_Dot(s::cl_double4, s::cl_double4);
30 __SYCL_EXPORT s::cl_half sycl_host_Dot(s::vec<s::half, 1>, s::vec<s::half, 1>);
31 __SYCL_EXPORT s::cl_half sycl_host_Dot(s::cl_half2, s::cl_half2);
32 __SYCL_EXPORT s::cl_half sycl_host_Dot(s::cl_half3, s::cl_half3);
33 __SYCL_EXPORT s::cl_half sycl_host_Dot(s::cl_half4, s::cl_half4);
34 
35 __SYCL_EXPORT s::cl_int sycl_host_All(s::vec<int, 1>);
36 __SYCL_EXPORT s::cl_int sycl_host_All(s::cl_int2);
37 __SYCL_EXPORT s::cl_int sycl_host_All(s::cl_int3);
38 __SYCL_EXPORT s::cl_int sycl_host_All(s::cl_int4);
39 
40 namespace {
41 
42 template <typename T> inline T __cross(T p0, T p1) {
43  T result(0);
44  result.x() = p0.y() * p1.z() - p0.z() * p1.y();
45  result.y() = p0.z() * p1.x() - p0.x() * p1.z();
46  result.z() = p0.x() * p1.y() - p0.y() * p1.x();
47  return result;
48 }
49 
50 template <typename T> inline void __FMul_impl(T &r, T p0, T p1) {
51  r += p0 * p1;
52 }
53 
54 template <typename T> inline T __FMul(T p0, T p1) {
55  T result = 0;
56  __FMul_impl(result, p0, p1);
57  return result;
58 }
59 
60 template <typename T>
61 inline typename sycl::detail::enable_if_t<d::is_sgengeo<T>::value, T>
62 __length(T t) {
63  return std::sqrt(__FMul(t, t));
64 }
65 
66 template <typename T>
67 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value,
68  typename T::element_type>
69 __length(T t) {
70  return std::sqrt(sycl_host_Dot(t, t));
71 }
72 
73 template <typename T>
74 inline typename sycl::detail::enable_if_t<d::is_sgengeo<T>::value, T>
75 __normalize(T t) {
76  T r = __length(t);
77  return t / T(r);
78 }
79 
80 template <typename T>
81 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value, T>
82 __normalize(T t) {
83  typename T::element_type r = __length(t);
84  return t / T(r);
85 }
86 
87 template <typename T>
88 inline typename sycl::detail::enable_if_t<d::is_sgengeo<T>::value, T>
89 __fast_length(T t) {
90  return std::sqrt(__FMul(t, t));
91 }
92 
93 template <typename T>
94 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value,
95  typename T::element_type>
96 __fast_length(T t) {
97  return std::sqrt(sycl_host_Dot(t, t));
98 }
99 
100 template <typename T>
101 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value, T>
102 __fast_normalize(T t) {
103  if (sycl_host_All(t == T(0.0f)))
104  return t;
105  typename T::element_type r = std::sqrt(sycl_host_Dot(t, t));
106  return t / T(r);
107 }
108 
109 } // namespace
110 
111 // --------------- 4.13.6 Geometric functions. Host implementations ------------
112 // cross
113 __SYCL_EXPORT s::cl_float3 sycl_host_cross(s::cl_float3 p0,
114  s::cl_float3 p1) __NOEXC {
115  return __cross(p0, p1);
116 }
117 __SYCL_EXPORT s::cl_float4 sycl_host_cross(s::cl_float4 p0,
118  s::cl_float4 p1) __NOEXC {
119  return __cross(p0, p1);
120 }
121 __SYCL_EXPORT s::cl_double3 sycl_host_cross(s::cl_double3 p0,
122  s::cl_double3 p1) __NOEXC {
123  return __cross(p0, p1);
124 }
125 __SYCL_EXPORT s::cl_double4 sycl_host_cross(s::cl_double4 p0,
126  s::cl_double4 p1) __NOEXC {
127  return __cross(p0, p1);
128 }
129 __SYCL_EXPORT s::cl_half3 sycl_host_cross(s::cl_half3 p0,
130  s::cl_half3 p1) __NOEXC {
131  return __cross(p0, p1);
132 }
133 __SYCL_EXPORT s::cl_half4 sycl_host_cross(s::cl_half4 p0,
134  s::cl_half4 p1) __NOEXC {
135  return __cross(p0, p1);
136 }
137 
138 // FMul
140  return __FMul(p0, p1);
141 }
143  return __FMul(p0, p1);
144 }
146  return __FMul(p0, p1);
147 }
148 
149 // Dot
151  s::cl_float)
153  s::cl_double)
155  s::cl_half)
156 
157 // length
158 __SYCL_EXPORT s::cl_float sycl_host_length(s::cl_float p) {
159  return __length(p);
160 }
162  return __length(p);
163 }
164 __SYCL_EXPORT s::cl_half sycl_host_length(s::cl_half p) { return __length(p); }
165 __SYCL_EXPORT s::cl_float sycl_host_length(s::vec<float, 1> p) {
166  return __length(p);
167 }
168 __SYCL_EXPORT s::cl_float sycl_host_length(s::cl_float2 p) {
169  return __length(p);
170 }
171 __SYCL_EXPORT s::cl_float sycl_host_length(s::cl_float3 p) {
172  return __length(p);
173 }
174 __SYCL_EXPORT s::cl_float sycl_host_length(s::cl_float4 p) {
175  return __length(p);
176 }
177 __SYCL_EXPORT s::cl_double sycl_host_length(s::vec<double, 1> p) {
178  return __length(p);
179 }
180 __SYCL_EXPORT s::cl_double sycl_host_length(s::cl_double2 p) {
181  return __length(p);
182 }
183 __SYCL_EXPORT s::cl_double sycl_host_length(s::cl_double3 p) {
184  return __length(p);
185 }
186 __SYCL_EXPORT s::cl_double sycl_host_length(s::cl_double4 p) {
187  return __length(p);
188 }
189 __SYCL_EXPORT s::cl_half sycl_host_length(s::vec<s::half, 1> p) {
190  return __length(p);
191 }
192 __SYCL_EXPORT s::cl_half sycl_host_length(s::cl_half2 p) { return __length(p); }
193 __SYCL_EXPORT s::cl_half sycl_host_length(s::cl_half3 p) { return __length(p); }
194 __SYCL_EXPORT s::cl_half sycl_host_length(s::cl_half4 p) { return __length(p); }
195 
196 // distance
198  return sycl_host_length(p0 - p1);
199 }
200 __SYCL_EXPORT s::cl_float sycl_host_distance(s::vec<float, 1> p0,
201  s::vec<float, 1> p1) {
202  return sycl_host_length(p0 - p1);
203 }
204 __SYCL_EXPORT s::cl_float sycl_host_distance(s::cl_float2 p0, s::cl_float2 p1) {
205  return sycl_host_length(p0 - p1);
206 }
207 __SYCL_EXPORT s::cl_float sycl_host_distance(s::cl_float3 p0, s::cl_float3 p1) {
208  return sycl_host_length(p0 - p1);
209 }
210 __SYCL_EXPORT s::cl_float sycl_host_distance(s::cl_float4 p0, s::cl_float4 p1) {
211  return sycl_host_length(p0 - p1);
212 }
214  s::cl_double p1) {
215  return sycl_host_length(p0 - p1);
216 }
217 __SYCL_EXPORT s::cl_float sycl_host_distance(s::vec<double, 1> p0,
218  s::vec<double, 1> p1) {
219  return sycl_host_length(p0 - p1);
220 }
221 __SYCL_EXPORT s::cl_double sycl_host_distance(s::cl_double2 p0,
222  s::cl_double2 p1) {
223  return sycl_host_length(p0 - p1);
224 }
225 __SYCL_EXPORT s::cl_double sycl_host_distance(s::cl_double3 p0,
226  s::cl_double3 p1) {
227  return sycl_host_length(p0 - p1);
228 }
229 __SYCL_EXPORT s::cl_double sycl_host_distance(s::cl_double4 p0,
230  s::cl_double4 p1) {
231  return sycl_host_length(p0 - p1);
232 }
234  return sycl_host_length(p0 - p1);
235 }
236 __SYCL_EXPORT s::cl_float sycl_host_distance(s::vec<s::half, 1> p0,
237  s::vec<s::half, 1> p1) {
238  return sycl_host_length(p0 - p1);
239 }
240 __SYCL_EXPORT s::cl_half sycl_host_distance(s::cl_half2 p0, s::cl_half2 p1) {
241  return sycl_host_length(p0 - p1);
242 }
243 __SYCL_EXPORT s::cl_half sycl_host_distance(s::cl_half3 p0, s::cl_half3 p1) {
244  return sycl_host_length(p0 - p1);
245 }
246 __SYCL_EXPORT s::cl_half sycl_host_distance(s::cl_half4 p0, s::cl_half4 p1) {
247  return sycl_host_length(p0 - p1);
248 }
249 
250 // normalize
252  return __normalize(p);
253 }
254 __SYCL_EXPORT s::cl_float sycl_host_normalize(s::vec<float, 1> p) {
255  return __normalize(p);
256 }
257 __SYCL_EXPORT s::cl_float2 sycl_host_normalize(s::cl_float2 p) {
258  return __normalize(p);
259 }
260 __SYCL_EXPORT s::cl_float3 sycl_host_normalize(s::cl_float3 p) {
261  return __normalize(p);
262 }
263 __SYCL_EXPORT s::cl_float4 sycl_host_normalize(s::cl_float4 p) {
264  return __normalize(p);
265 }
267  return __normalize(p);
268 }
269 __SYCL_EXPORT s::cl_double sycl_host_normalize(s::vec<double, 1> p) {
270  return __normalize(p);
271 }
272 __SYCL_EXPORT s::cl_double2 sycl_host_normalize(s::cl_double2 p) {
273  return __normalize(p);
274 }
275 __SYCL_EXPORT s::cl_double3 sycl_host_normalize(s::cl_double3 p) {
276  return __normalize(p);
277 }
278 __SYCL_EXPORT s::cl_double4 sycl_host_normalize(s::cl_double4 p) {
279  return __normalize(p);
280 }
282  return __normalize(p);
283 }
284 __SYCL_EXPORT s::cl_half2 sycl_host_normalize(s::cl_half2 p) {
285  return __normalize(p);
286 }
287 __SYCL_EXPORT s::cl_half3 sycl_host_normalize(s::cl_half3 p) {
288  return __normalize(p);
289 }
290 __SYCL_EXPORT s::cl_half4 sycl_host_normalize(s::cl_half4 p) {
291  return __normalize(p);
292 }
293 
294 // fast_length
296  return __fast_length(p);
297 }
298 __SYCL_EXPORT s::cl_float sycl_host_fast_length(s::vec<float, 1> p) {
299  return __fast_length(p);
300 }
301 __SYCL_EXPORT s::cl_float sycl_host_fast_length(s::cl_float2 p) {
302  return __fast_length(p);
303 }
304 __SYCL_EXPORT s::cl_float sycl_host_fast_length(s::cl_float3 p) {
305  return __fast_length(p);
306 }
307 __SYCL_EXPORT s::cl_float sycl_host_fast_length(s::cl_float4 p) {
308  return __fast_length(p);
309 }
310 
311 // fast_normalize
313  if (p == 0.0f)
314  return p;
316  return p / r;
317 }
318 __SYCL_EXPORT s::cl_float sycl_host_fast_normalize(s::vec<float, 1> p) {
319  return __fast_normalize(p);
320 }
321 __SYCL_EXPORT s::cl_float2 sycl_host_fast_normalize(s::cl_float2 p) {
322  return __fast_normalize(p);
323 }
324 __SYCL_EXPORT s::cl_float3 sycl_host_fast_normalize(s::cl_float3 p) {
325  return __fast_normalize(p);
326 }
327 __SYCL_EXPORT s::cl_float4 sycl_host_fast_normalize(s::cl_float4 p) {
328  return __fast_normalize(p);
329 }
330 
331 // fast_distance
333  s::cl_float p1) {
334  return sycl_host_fast_length(p0 - p1);
335 }
336 __SYCL_EXPORT s::cl_float sycl_host_fast_distance(s::vec<float, 1> p0,
337  s::vec<float, 1> p1) {
338  return sycl_host_fast_length(p0 - p1);
339 }
340 __SYCL_EXPORT s::cl_float sycl_host_fast_distance(s::cl_float2 p0,
341  s::cl_float2 p1) {
342  return sycl_host_fast_length(p0 - p1);
343 }
344 __SYCL_EXPORT s::cl_float sycl_host_fast_distance(s::cl_float3 p0,
345  s::cl_float3 p1) {
346  return sycl_host_fast_length(p0 - p1);
347 }
348 __SYCL_EXPORT s::cl_float sycl_host_fast_distance(s::cl_float4 p0,
349  s::cl_float4 p1) {
350  return sycl_host_fast_length(p0 - p1);
351 }
352 
353 } // namespace __host_std
#define __NOEXC
Definition: builtins.hpp:18
s::cl_double sycl_host_length(s::cl_double p)
s::cl_float sycl_host_distance(s::cl_float p0, s::cl_float p1)
s::cl_float sycl_host_fast_length(s::cl_float p)
s::cl_float sycl_host_fast_normalize(s::cl_float p)
s::cl_int sycl_host_All(s::vec< int, 1 >)
s::cl_float sycl_host_Dot(s::vec< float, 1 >, s::vec< float, 1 >)
s::cl_float sycl_host_fast_distance(s::cl_float p0, s::cl_float p1)
MAKE_GEO_1V_2V_RS(sycl_host_Dot, __FMul_impl, s::cl_float, s::cl_float, s::cl_float) MAKE_GEO_1V_2V_RS(sycl_host_Dot
s::cl_float sycl_host_FMul(s::cl_float p0, s::cl_float p1)
s::cl_float sycl_host_normalize(s::cl_float p)
s::cl_float3 sycl_host_cross(s::cl_float3 p0, s::cl_float3 p1) __NOEXC
half cl_half
Definition: aliases.hpp:87
double cl_double
Definition: aliases.hpp:89
std::int32_t cl_int
Definition: aliases.hpp:83
float cl_float
Definition: aliases.hpp:88
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > sqrt(T x) __NOEXC
Definition: builtins.hpp:469
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14