WIP

celerity · Dec 29, 2023 · 9ec2ff5 · 9ec2ff5
1 parent a8d5f69
commit 9ec2ff5
Show file tree

Hide file tree

Showing 8 changed files with 117 additions and 59 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -67,6 +67,7 @@ add_library(simsycl
     include/simsycl/detail/config.hh
     include/simsycl/detail/coordinate.hh
     include/simsycl/detail/hash.hh
+    include/simsycl/detail/math_utils.hh
     include/simsycl/detail/subscript.hh
     include/simsycl/detail/utils.hh
     include/simsycl/detail/vec_swizzles.inc
@@ -76,11 +77,12 @@ add_library(simsycl
     include/simsycl/sycl/atomic_fence.hh
     include/simsycl/sycl/atomic_ref.hh
     include/simsycl/sycl/backend.hh
+    include/simsycl/sycl/binary_ops.hh
     include/simsycl/sycl/buffer.hh
     include/simsycl/sycl/concepts.hh
     include/simsycl/sycl/context.hh
-    include/simsycl/sycl/device.hh
     include/simsycl/sycl/device_selector.hh
+    include/simsycl/sycl/device.hh
     include/simsycl/sycl/enums.hh
     include/simsycl/sycl/event.hh
     include/simsycl/sycl/exception.hh
@@ -91,18 +93,18 @@ add_library(simsycl
     include/simsycl/sycl/h_item.hh
     include/simsycl/sycl/handler.hh
     include/simsycl/sycl/id.hh
+    include/simsycl/sycl/image.hh
     include/simsycl/sycl/info.hh
     include/simsycl/sycl/interop_handle.hh
-    include/simsycl/sycl/image.hh
     include/simsycl/sycl/item.hh
     include/simsycl/sycl/kernel.hh
-    include/simsycl/sycl/math.hh
-    include/simsycl/sycl/math_geometric.hh
     include/simsycl/sycl/marray.hh
+    include/simsycl/sycl/math_common.hh
+    include/simsycl/sycl/math_geometric.hh
+    include/simsycl/sycl/math.hh
     include/simsycl/sycl/multi_ptr.hh
     include/simsycl/sycl/nd_item.hh
     include/simsycl/sycl/nd_range.hh
-    include/simsycl/sycl/binary_ops.hh
     include/simsycl/sycl/platform.hh
     include/simsycl/sycl/private_memory.hh
     include/simsycl/sycl/property.hh

diff --git a/include/simsycl/detail/math_utils.hh b/include/simsycl/detail/math_utils.hh
@@ -0,0 +1,74 @@
+#pragma once
+
+#include "../sycl/marray.hh"
+#include "../sycl/vec.hh"
+
+namespace simsycl::detail {
+
+template<typename T>
+concept SyclFloat = std::is_same_v<T, float> || std::is_same_v<T, double>
+#if SIMSYCL_FEATURE_HALF_TYPE
+    || std::is_same_v<T, sycl::half>
+#endif
+    ;
+
+template<SyclFloat T>
+struct num_elements<T> : std::integral_constant<int, 1> {};
+
+template<typename T>
+concept GenFloat
+    = SyclFloat<T> || ((is_swizzle_v<T> || is_vec_v<T> || is_marray_v<T>)&&SyclFloat<typename T::element_type>);
+
+template<typename T>
+concept GeoFloat = SyclFloat<T>
+    || ((is_swizzle_v<T> || is_vec_v<T> || is_marray_v<T>)&&(num_elements_v<T> > 0 && num_elements_v<T> <= 4)
+        && SyclFloat<typename T::value_type>);
+
+template<typename T>
+    requires(is_vec_v<T> || is_swizzle_v<T> || is_marray_v<T>)
+auto sum(const T &f) {
+    auto ret = f[0];
+    for(int i = 1; i < num_elements_v<T>; ++i) { ret += f[i]; }
+    return ret;
+}
+template<SyclFloat T>
+auto sum(const T &f) {
+    return f;
+}
+
+template<typename T>
+struct element_type {
+    using type = T;
+};
+template<typename T>
+    requires(is_vec_v<T> || is_swizzle_v<T>)
+struct element_type<T> {
+    using type = typename T::element_type;
+};
+template<typename T>
+    requires(is_marray_v<T>)
+struct element_type<T> {
+    using type = typename T::value_type;
+};
+template<typename T>
+using element_type_t = typename element_type<T>::type;
+
+template<typename DataT, int NumElements>
+sycl::vec<DataT, NumElements> marray_to_vec(const sycl::marray<DataT, NumElements> &v) {
+    sycl::vec<DataT, NumElements> ret;
+    for(int i = 0; i < NumElements; ++i) { ret[i] = v[i]; }
+    return ret;
+}
+
+template<typename VT, typename T>
+    requires(!is_marray_v<T>)
+sycl::vec<element_type_t<VT>, num_elements_v<VT>> to_matching_vec(const T &v) {
+    return to_vec<element_type_t<VT>, num_elements_v<VT>>(v);
+}
+template<typename VT, typename T>
+    requires(is_marray_v<T>)
+sycl::vec<element_type_t<VT>, num_elements_v<VT>> to_matching_vec(const T &v) {
+    return marray_to_vec<element_type_t<VT>, num_elements_v<VT>>(v);
+}
+
+} // namespace simsycl::detail
diff --git a/include/simsycl/sycl.hh b/include/simsycl/sycl.hh
@@ -27,6 +27,7 @@
 #include "sycl/kernel.hh"
 #include "sycl/marray.hh"
 #include "sycl/math.hh"
+#include "sycl/math_common.hh"
 #include "sycl/math_geometric.hh"
 #include "sycl/multi_ptr.hh"
 #include "sycl/nd_item.hh"

diff --git a/include/simsycl/sycl/marray.hh b/include/simsycl/sycl/marray.hh
@@ -28,11 +28,11 @@ struct marray_init_arg_traits<DataT, sycl::marray<DataT, N>, ArgTN...> {
     static constexpr size_t num_elements = N + marray_init_arg_traits<DataT, ArgTN...>::num_elements;
 };
 
-template<typename MArray, typename DataT>
+template<typename T>
 constexpr bool is_marray_v = false;
 
 template<typename DataT, size_t NumElements>
-constexpr bool is_marray_v<sycl::marray<DataT, NumElements>, DataT> = true;
+constexpr bool is_marray_v<sycl::marray<DataT, NumElements>> = true;
 
 } // namespace simsycl::detail
 

diff --git a/include/simsycl/sycl/math_common.hh b/include/simsycl/sycl/math_common.hh
@@ -0,0 +1,11 @@
+#pragma once
+
+#include "../detail/math_utils.hh"
+
+namespace simsycl::sycl {
+
+// Note: arguments are passed by const ref rather than value to avoid gcc warnings
+//       the standard requires pass-by-value, but I'm not sure if this is visible to the user
+
+
+} // namespace simsycl::sycl
diff --git a/include/simsycl/sycl/math_geometric.hh b/include/simsycl/sycl/math_geometric.hh
@@ -1,53 +1,6 @@
 #pragma once
 
-#include "math.hh"
-#include "vec.hh"
-
-namespace simsycl::detail {
-
-template<typename T>
-concept SyclFloat = std::is_same_v<T, float> || std::is_same_v<T, double>
-#if SIMSYCL_FEATURE_HALF_TYPE
-    || std::is_same_v<T, sycl::half>
-#endif
-    ;
-
-template<typename T>
-concept GeoFloat = //
-    SyclFloat<T>
-    || ((is_swizzle_v<T> || is_vec_v<T>)&&(num_elements_v<T> > 0 && num_elements_v<T> <= 4)
-        && SyclFloat<typename T::element_type>); // TODO: marray
-
-template<typename T>
-    requires(is_vec_v<T> || is_swizzle_v<T>)
-auto sum(const T &f) {
-    auto ret = f[0];
-    for(int i = 1; i < num_elements_v<T>; ++i) { ret += f[i]; }
-    return ret;
-}
-template<SyclFloat T>
-auto sum(const T &f) {
-    return f;
-}
-
-template<GeoFloat T>
-struct element_type {
-    using type = T;
-};
-template<GeoFloat T>
-    requires(is_vec_v<T> || is_swizzle_v<T>)
-struct element_type<T> {
-    using type = typename T::element_type;
-};
-template<GeoFloat T>
-using element_type_t = typename element_type<T>::type;
-
-template<typename VT, typename T>
-auto to_matching_vec(const T &v) {
-    return detail::to_vec<detail::element_type_t<VT>, detail::num_elements_v<VT>>(v);
-}
-
-} // namespace simsycl::detail
+#include "../detail/math_utils.hh"
 
 namespace simsycl::sycl {
 

diff --git a/include/simsycl/sycl/vec.hh b/include/simsycl/sycl/vec.hh
@@ -75,17 +75,22 @@ template<typename DataT, int... Indices>
 class swizzled_vec;
 
 
-template<typename VecOrSwizzle>
-struct num_elements : std::integral_constant<int, 1> {};
+template<typename T>
+struct num_elements {
+    static_assert(sizeof(T) < 0, "num_elements instantiated with unsupported type");
+};
 
 template<typename T, int NumElements>
 struct num_elements<sycl::vec<T, NumElements>> : std::integral_constant<int, NumElements> {};
 
 template<typename T, int... Indices>
 struct num_elements<detail::swizzled_vec<T, Indices...>> : std::integral_constant<int, sizeof...(Indices)> {};
 
-template<typename VecOrSwizzle>
-static constexpr int num_elements_v = num_elements<VecOrSwizzle>::value;
+template<typename T, int NumElements>
+struct num_elements<sycl::marray<T, NumElements>> : std::integral_constant<int, NumElements> {};
+
+template<typename T>
+static constexpr int num_elements_v = num_elements<T>::value;
 
 
 template<int... Is>

diff --git a/test/math_tests.cc b/test/math_tests.cc
@@ -27,6 +27,9 @@ TEST_CASE("Length function works as expected", "[math][geometric]") {
     CHECK(length(v1.xx()) == Catch::Approx(1.4142135624));
     CHECK(length(v6.argb()) == Catch::Approx(33.07567f));
 
+    marray<float, 4> m1 = {1.0f, 2.0f, 3.0f, 4.0f};
+    CHECK(length(m1) == Catch::Approx(5.4772255751f));
+
 #if SIMSYCL_FEATURE_HALF_TYPE
     using sycl::half;
     half h = 7.0f;
@@ -59,3 +62,12 @@ TEST_CASE("Distance function works as expected", "[math][geometric]") {
     CHECK(distance(v6.argb(), vec<float, 4>{0.0f, 0.0f, 0.0f, 0.0f}) == Catch::Approx(33.07567f));
     CHECK(distance(v6.argb(), v2.xyxy()) == Catch::Approx(26.15339f));
 }
+
+TEST_CASE("Clamp function works as expected", "[math]") {
+    int x = 8;
+    CHECK(clamp(x, 0, 10) == 8);
+    CHECK(clamp(x, 0, 5) == 5);
+    CHECK(clamp(x, 10, 20) == 10);
+    // vec<int, 4> v1 = {1, 2, 3, 4};
+    // CHECK(clamp(v1, 0, 10) == v1);
+}