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 = cl::sycl;
18 namespace d = s::detail;
19 
21 namespace __host_std {
22 
24 __SYCL_EXPORT s::cl_float Dot(s::cl_float2, s::cl_float2);
25 __SYCL_EXPORT s::cl_float Dot(s::cl_float3, s::cl_float3);
26 __SYCL_EXPORT s::cl_float Dot(s::cl_float4, s::cl_float4);
28 __SYCL_EXPORT s::cl_double Dot(s::cl_double2, s::cl_double2);
29 __SYCL_EXPORT s::cl_double Dot(s::cl_double3, s::cl_double3);
30 __SYCL_EXPORT s::cl_double Dot(s::cl_double4, s::cl_double4);
32 __SYCL_EXPORT s::cl_half Dot(s::cl_half2, s::cl_half2);
33 __SYCL_EXPORT s::cl_half Dot(s::cl_half3, s::cl_half3);
34 __SYCL_EXPORT s::cl_half Dot(s::cl_half4, s::cl_half4);
35 
36 __SYCL_EXPORT s::cl_int All(s::vec<int, 1>);
37 __SYCL_EXPORT s::cl_int All(s::cl_int2);
38 __SYCL_EXPORT s::cl_int All(s::cl_int3);
39 __SYCL_EXPORT s::cl_int All(s::cl_int4);
40 
41 namespace {
42 
43 template <typename T> inline T __cross(T p0, T p1) {
44  T result(0);
45  result.x() = p0.y() * p1.z() - p0.z() * p1.y();
46  result.y() = p0.z() * p1.x() - p0.x() * p1.z();
47  result.z() = p0.x() * p1.y() - p0.y() * p1.x();
48  return result;
49 }
50 
51 template <typename T> inline void __FMul_impl(T &r, T p0, T p1) {
52  r += p0 * p1;
53 }
54 
55 template <typename T> inline T __FMul(T p0, T p1) {
56  T result = 0;
57  __FMul_impl(result, p0, p1);
58  return result;
59 }
60 
61 template <typename T>
62 inline typename sycl::detail::enable_if_t<d::is_sgengeo<T>::value, T>
63 __length(T t) {
64  return std::sqrt(__FMul(t, t));
65 }
66 
67 template <typename T>
68 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value,
69  typename T::element_type>
70 __length(T t) {
71  return std::sqrt(Dot(t, t));
72 }
73 
74 template <typename T>
75 inline typename sycl::detail::enable_if_t<d::is_sgengeo<T>::value, T>
76 __normalize(T t) {
77  T r = __length(t);
78  return t / T(r);
79 }
80 
81 template <typename T>
82 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value, T>
83 __normalize(T t) {
84  typename T::element_type r = __length(t);
85  return t / T(r);
86 }
87 
88 template <typename T>
89 inline typename sycl::detail::enable_if_t<d::is_sgengeo<T>::value, T>
90 __fast_length(T t) {
91  return std::sqrt(__FMul(t, t));
92 }
93 
94 template <typename T>
95 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value,
96  typename T::element_type>
97 __fast_length(T t) {
98  return std::sqrt(Dot(t, t));
99 }
100 
101 template <typename T>
102 inline typename sycl::detail::enable_if_t<d::is_vgengeo<T>::value, T>
103 __fast_normalize(T t) {
104  if (All(t == T(0.0f)))
105  return t;
106  typename T::element_type r = std::sqrt(Dot(t, t));
107  return t / T(r);
108 }
109 
110 } // namespace
111 
112 // --------------- 4.13.6 Geometric functions. Host implementations ------------
113 // cross
114 __SYCL_EXPORT s::cl_float3 cross(s::cl_float3 p0, s::cl_float3 p1) __NOEXC {
115  return __cross(p0, p1);
116 }
117 __SYCL_EXPORT s::cl_float4 cross(s::cl_float4 p0, s::cl_float4 p1) __NOEXC {
118  return __cross(p0, p1);
119 }
120 __SYCL_EXPORT s::cl_double3 cross(s::cl_double3 p0, s::cl_double3 p1) __NOEXC {
121  return __cross(p0, p1);
122 }
123 __SYCL_EXPORT s::cl_double4 cross(s::cl_double4 p0, s::cl_double4 p1) __NOEXC {
124  return __cross(p0, p1);
125 }
126 __SYCL_EXPORT s::cl_half3 cross(s::cl_half3 p0, s::cl_half3 p1) __NOEXC {
127  return __cross(p0, p1);
128 }
129 __SYCL_EXPORT s::cl_half4 cross(s::cl_half4 p0, s::cl_half4 p1) __NOEXC {
130  return __cross(p0, p1);
131 }
132 
133 // FMul
134 __SYCL_EXPORT s::cl_float FMul(s::cl_float p0, s::cl_float p1) {
135  return __FMul(p0, p1);
136 }
138  return __FMul(p0, p1);
139 }
140 __SYCL_EXPORT s::cl_float FMul(s::cl_half p0, s::cl_half p1) {
141  return __FMul(p0, p1);
142 }
143 
144 // Dot
148 
149 // length
150 __SYCL_EXPORT s::cl_float length(s::cl_float p) { return __length(p); }
151 __SYCL_EXPORT s::cl_double length(s::cl_double p) { return __length(p); }
152 __SYCL_EXPORT s::cl_half length(s::cl_half p) { return __length(p); }
153 __SYCL_EXPORT s::cl_float length(s::vec<float, 1> p) { return __length(p); }
154 __SYCL_EXPORT s::cl_float length(s::cl_float2 p) { return __length(p); }
155 __SYCL_EXPORT s::cl_float length(s::cl_float3 p) { return __length(p); }
156 __SYCL_EXPORT s::cl_float length(s::cl_float4 p) { return __length(p); }
157 __SYCL_EXPORT s::cl_double length(s::vec<double, 1> p) { return __length(p); }
158 __SYCL_EXPORT s::cl_double length(s::cl_double2 p) { return __length(p); }
159 __SYCL_EXPORT s::cl_double length(s::cl_double3 p) { return __length(p); }
160 __SYCL_EXPORT s::cl_double length(s::cl_double4 p) { return __length(p); }
161 __SYCL_EXPORT s::cl_half length(s::vec<s::half, 1> p) { return __length(p); }
162 __SYCL_EXPORT s::cl_half length(s::cl_half2 p) { return __length(p); }
163 __SYCL_EXPORT s::cl_half length(s::cl_half3 p) { return __length(p); }
164 __SYCL_EXPORT s::cl_half length(s::cl_half4 p) { return __length(p); }
165 
166 // distance
168  return length(p0 - p1);
169 }
171  return length(p0 - p1);
172 }
173 __SYCL_EXPORT s::cl_float distance(s::cl_float2 p0, s::cl_float2 p1) {
174  return length(p0 - p1);
175 }
176 __SYCL_EXPORT s::cl_float distance(s::cl_float3 p0, s::cl_float3 p1) {
177  return length(p0 - p1);
178 }
179 __SYCL_EXPORT s::cl_float distance(s::cl_float4 p0, s::cl_float4 p1) {
180  return length(p0 - p1);
181 }
183  return length(p0 - p1);
184 }
186  return length(p0 - p1);
187 }
188 __SYCL_EXPORT s::cl_double distance(s::cl_double2 p0, s::cl_double2 p1) {
189  return length(p0 - p1);
190 }
191 __SYCL_EXPORT s::cl_double distance(s::cl_double3 p0, s::cl_double3 p1) {
192  return length(p0 - p1);
193 }
194 __SYCL_EXPORT s::cl_double distance(s::cl_double4 p0, s::cl_double4 p1) {
195  return length(p0 - p1);
196 }
197 __SYCL_EXPORT s::cl_half distance(s::cl_half p0, s::cl_half p1) {
198  return length(p0 - p1);
199 }
201  s::vec<s::half, 1> p1) {
202  return length(p0 - p1);
203 }
204 __SYCL_EXPORT s::cl_half distance(s::cl_half2 p0, s::cl_half2 p1) {
205  return length(p0 - p1);
206 }
207 __SYCL_EXPORT s::cl_half distance(s::cl_half3 p0, s::cl_half3 p1) {
208  return length(p0 - p1);
209 }
210 __SYCL_EXPORT s::cl_half distance(s::cl_half4 p0, s::cl_half4 p1) {
211  return length(p0 - p1);
212 }
213 
214 // normalize
215 __SYCL_EXPORT s::cl_float normalize(s::cl_float p) { return __normalize(p); }
217  return __normalize(p);
218 }
219 __SYCL_EXPORT s::cl_float2 normalize(s::cl_float2 p) { return __normalize(p); }
220 __SYCL_EXPORT s::cl_float3 normalize(s::cl_float3 p) { return __normalize(p); }
221 __SYCL_EXPORT s::cl_float4 normalize(s::cl_float4 p) { return __normalize(p); }
222 __SYCL_EXPORT s::cl_double normalize(s::cl_double p) { return __normalize(p); }
224  return __normalize(p);
225 }
226 __SYCL_EXPORT s::cl_double2 normalize(s::cl_double2 p) {
227  return __normalize(p);
228 }
229 __SYCL_EXPORT s::cl_double3 normalize(s::cl_double3 p) {
230  return __normalize(p);
231 }
232 __SYCL_EXPORT s::cl_double4 normalize(s::cl_double4 p) {
233  return __normalize(p);
234 }
235 __SYCL_EXPORT s::cl_half normalize(s::cl_half p) { return __normalize(p); }
236 __SYCL_EXPORT s::cl_half2 normalize(s::cl_half2 p) { return __normalize(p); }
237 __SYCL_EXPORT s::cl_half3 normalize(s::cl_half3 p) { return __normalize(p); }
238 __SYCL_EXPORT s::cl_half4 normalize(s::cl_half4 p) { return __normalize(p); }
239 
240 // fast_length
242  return __fast_length(p);
243 }
245  return __fast_length(p);
246 }
247 __SYCL_EXPORT s::cl_float fast_length(s::cl_float2 p) {
248  return __fast_length(p);
249 }
250 __SYCL_EXPORT s::cl_float fast_length(s::cl_float3 p) {
251  return __fast_length(p);
252 }
253 __SYCL_EXPORT s::cl_float fast_length(s::cl_float4 p) {
254  return __fast_length(p);
255 }
256 
257 // fast_normalize
259  if (p == 0.0f)
260  return p;
261  s::cl_float r = std::sqrt(FMul(p, p));
262  return p / r;
263 }
265  return __fast_normalize(p);
266 }
267 __SYCL_EXPORT s::cl_float2 fast_normalize(s::cl_float2 p) {
268  return __fast_normalize(p);
269 }
270 __SYCL_EXPORT s::cl_float3 fast_normalize(s::cl_float3 p) {
271  return __fast_normalize(p);
272 }
273 __SYCL_EXPORT s::cl_float4 fast_normalize(s::cl_float4 p) {
274  return __fast_normalize(p);
275 }
276 
277 // fast_distance
279  return fast_length(p0 - p1);
280 }
282  s::vec<float, 1> p1) {
283  return fast_length(p0 - p1);
284 }
285 __SYCL_EXPORT s::cl_float fast_distance(s::cl_float2 p0, s::cl_float2 p1) {
286  return fast_length(p0 - p1);
287 }
288 __SYCL_EXPORT s::cl_float fast_distance(s::cl_float3 p0, s::cl_float3 p1) {
289  return fast_length(p0 - p1);
290 }
291 __SYCL_EXPORT s::cl_float fast_distance(s::cl_float4 p0, s::cl_float4 p1) {
292  return fast_length(p0 - p1);
293 }
294 
295 } // namespace __host_std
296 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::cl_double
double cl_double
Definition: aliases.hpp:89
T
cl::sycl::sqrt
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > sqrt(T x) __NOEXC
Definition: builtins.hpp:469
cl::sycl
Definition: access.hpp:14
stl_type_traits.hpp
cl::__host_std::normalize
s::cl_half4 normalize(s::cl_half4 p)
Definition: builtins_geometric.cpp:238
cl::__host_std::length
s::cl_half length(s::cl_half4 p)
Definition: builtins_geometric.cpp:164
cl::sycl::detail::half_impl::half
Definition: half_type.hpp:329
cl::__host_std::FMul
s::cl_float FMul(s::cl_half p0, s::cl_half p1)
Definition: builtins_geometric.cpp:140
__NOEXC
#define __NOEXC
Definition: builtins.hpp:18
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::image_channel_order::r
@ r
builtins_helper.hpp
cl::__host_std::cross
s::cl_half4 cross(s::cl_half4 p0, s::cl_half4 p1) __NOEXC
Definition: builtins_geometric.cpp:129
cl::sycl::cl_float
float cl_float
Definition: aliases.hpp:88
cl::sycl::vec
Provides a cross-patform vector class template that works efficiently on SYCL devices as well as in h...
Definition: aliases.hpp:19
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:83
cl::__host_std::Dot
s::cl_half Dot(s::cl_half4, s::cl_half4)
MAKE_GEO_1V_2V_RS
#define MAKE_GEO_1V_2V_RS(Fun, Call, Ret, Arg1, Arg2)
Definition: builtins_helper.hpp:204
cl::__host_std::fast_distance
s::cl_float fast_distance(s::cl_float4 p0, s::cl_float4 p1)
Definition: builtins_geometric.cpp:291
cl::__host_std::fast_length
s::cl_float fast_length(s::cl_float4 p)
Definition: builtins_geometric.cpp:253
All
@ All
Definition: usm_allocator.hpp:14
cl::__host_std::fast_normalize
s::cl_float4 fast_normalize(s::cl_float4 p)
Definition: builtins_geometric.cpp:273
cl::__host_std::distance
s::cl_half distance(s::cl_half4 p0, s::cl_half4 p1)
Definition: builtins_geometric.cpp:210
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12