Replace uses of STL types in device code

This enables us to use the code in CUDA without the use of the experimental constexpr relaxation.
acts-project · Nov 20, 2024 · d4e6098 · d4e6098
1 parent af93396
commit d4e6098
Show file tree

Hide file tree

Showing 27 changed files with 324 additions and 188 deletions.
diff --git a/cmake/covfie-compiler-options-cuda.cmake b/cmake/covfie-compiler-options-cuda.cmake
@@ -31,9 +31,6 @@ if("${CMAKE_CUDA_COMPILER_ID}" MATCHES "NVIDIA")
     # Make CUDA generate debug symbols for the device code as well in a debug
     # build.
     covfie_add_flag( CMAKE_CUDA_FLAGS_DEBUG "-G" )
-    # Allow to use functions in device code that are constexpr, even if they are
-    # not marked with __device__.
-    covfie_add_flag( CMAKE_CUDA_FLAGS "--expt-relaxed-constexpr" )
 endif()
 
 covfie_add_flag( CMAKE_CUDA_FLAGS "-Wfloat-conversion" )

diff --git a/examples/core/generate_test_field.cpp b/examples/core/generate_test_field.cpp
@@ -74,7 +74,7 @@ int main(int argc, char ** argv)
     BOOST_LOG_TRIVIAL(info) << "Allocating space for new vector field...";
 
     core_field_t cf(covfie::make_parameter_pack(core_backend_t::configuration_t{
-        201, 201, 301}));
+        201u, 201u, 301u}));
 
     BOOST_LOG_TRIVIAL(info) << "Filling new vector field...";
 

diff --git a/examples/core/scaleup_bfield.cpp b/examples/core/scaleup_bfield.cpp
@@ -99,7 +99,7 @@ int main(int argc, char ** argv)
     output_field_t new_field(covfie::make_parameter_pack(
         scaling * translation,
         std::monostate{},
-        covfie::utility::nd_size<3>{601, 601, 901}
+        covfie::utility::nd_size<3>{601u, 601u, 901u}
     ));
 
     BOOST_LOG_TRIVIAL(info) << "Copying data from old field to new field...";

diff --git a/examples/core/slice3dto2d.cpp b/examples/core/slice3dto2d.cpp
@@ -122,7 +122,7 @@ int main(int argc, char ** argv)
 
     covfie::utility::nd_size<3> in_size =
         f.backend().get_backend().get_backend().get_configuration();
-    covfie::utility::nd_size<2> out_size{0, 0};
+    covfie::utility::nd_size<2> out_size{0u, 0u};
 
     if (vm["axis"].as<std::string>() == "x") {
         out_size = {in_size[1], in_size[2]};

diff --git a/examples/cuda/render_slice.cu b/examples/cuda/render_slice.cu
@@ -95,12 +95,12 @@ __global__ void render(
 
         typename field_t::output_t p =
             vf.at(fx * 20000.f - 10000.f, fy * 20000.f - 10000.f, z);
-        out[height * x + y] = static_cast<unsigned char>(std::lround(
-            255.f * std::min(
-                        std::sqrt(
-                            std::pow(p[0] / 0.000299792458f, 2.f) +
-                            std::pow(p[1] / 0.000299792458f, 2.f) +
-                            std::pow(p[2] / 0.000299792458f, 2.f)
+        out[height * x + y] = static_cast<unsigned char>(::lroundf(
+            255.f * ::fmin(
+                        ::sqrtf(
+                            ::powf(p[0] / 0.000299792458f, 2.f) +
+                            ::powf(p[1] / 0.000299792458f, 2.f) +
+                            ::powf(p[2] / 0.000299792458f, 2.f)
                         ),
                         1.0f
                     )

diff --git a/examples/cuda/render_slice_texture.cu b/examples/cuda/render_slice_texture.cu
@@ -97,11 +97,11 @@ __global__ void render(
         typename field_t::output_t p =
             vf.at(fx * 20000.f - 10000.f, fy * 20000.f - 10000.f, z);
         out[height * x + y] = static_cast<unsigned char>(std::lround(
-            255.f * std::min(
-                        std::sqrt(
-                            std::pow(p[0] / 0.000299792458f, 2.f) +
-                            std::pow(p[1] / 0.000299792458f, 2.f) +
-                            std::pow(p[2] / 0.000299792458f, 2.f)
+            255.f * fmin(
+                        ::sqrtf(
+                            ::powf(p[0] / 0.000299792458f, 2.f) +
+                            ::powf(p[1] / 0.000299792458f, 2.f) +
+                            ::powf(p[2] / 0.000299792458f, 2.f)
                         ),
                         1.0f
                     )

diff --git a/lib/core/covfie/core/algebra/affine.hpp b/lib/core/covfie/core/algebra/affine.hpp
@@ -14,6 +14,7 @@
 
 #include <covfie/core/algebra/matrix.hpp>
 #include <covfie/core/algebra/vector.hpp>
+#include <covfie/core/array.hpp>
 #include <covfie/core/qualifiers.hpp>
 
 namespace covfie::algebra {
@@ -90,7 +91,7 @@ struct affine : public matrix<N, N + 1, T, I> {
             "of the matrix."
         );
 
-        std::array<T, N> arr{args...};
+        array::array<T, N> arr{args...};
 
         matrix<N, N + 1, T, I> result = matrix<N, N + 1, T, I>::identity();
 
@@ -115,7 +116,7 @@ struct affine : public matrix<N, N + 1, T, I> {
             "the matrix."
         );
 
-        std::array<T, N> arr{args...};
+        array::array<T, N> arr{args...};
 
         matrix<N, N + 1, T, I> result = matrix<N, N + 1, T, I>::identity();
 

diff --git a/lib/core/covfie/core/algebra/matrix.hpp b/lib/core/covfie/core/algebra/matrix.hpp
@@ -10,9 +10,9 @@
 
 #pragma once
 
-#include <array>
 #include <cstddef>
 
+#include <covfie/core/array.hpp>
 #include <covfie/core/qualifiers.hpp>
 
 namespace covfie::algebra {
@@ -26,7 +26,7 @@ struct matrix {
     {
     }
 
-    COVFIE_DEVICE matrix(std::array<std::array<T, M>, N> l)
+    COVFIE_DEVICE matrix(array::array<array::array<T, M>, N> l)
     {
         for (I i = 0; i < l.size(); ++i) {
             for (I j = 0; j < l[i].size(); ++j) {

diff --git a/lib/core/covfie/core/algebra/vector.hpp b/lib/core/covfie/core/algebra/vector.hpp
@@ -16,6 +16,7 @@
 #include <utility>
 
 #include <covfie/core/algebra/matrix.hpp>
+#include <covfie/core/array.hpp>
 #include <covfie/core/qualifiers.hpp>
 
 namespace covfie::algebra {
@@ -26,7 +27,7 @@ struct vector : public matrix<N, 1, T, I> {
     {
     }
 
-    COVFIE_DEVICE vector(std::array<T, N> l)
+    COVFIE_DEVICE vector(array::array<T, N> l)
         : matrix<N, 1, T, I>()
     {
         for (I i = 0; i < N; ++i) {
@@ -47,7 +48,7 @@ struct vector : public matrix<N, 1, T, I> {
                 sizeof...(Args) == N,
             bool> = true>
     COVFIE_DEVICE vector(Args... args)
-        : vector(std::array<T, N>{std::forward<Args>(args)...})
+        : vector(array::array<T, N>{std::forward<Args>(args)...})
     {
     }
 

diff --git a/lib/core/covfie/core/array.hpp b/lib/core/covfie/core/array.hpp
@@ -0,0 +1,125 @@
+/*
+ * This file is part of covfie, a part of the ACTS project
+ *
+ * Copyright (c) 2022 CERN
+ *
+ * This Source Code Form is subject to the terms of the Mozilla Public License,
+ * v. 2.0. If a copy of the MPL was not distributed with this file, You can
+ * obtain one at http://mozilla.org/MPL/2.0/.
+ */
+
+#pragma once
+
+#include <cassert>
+#include <cstddef>
+#include <utility>
+
+#include <covfie/core/qualifiers.hpp>
+
+namespace covfie::array {
+template <typename _scalar_t, std::size_t _size>
+requires(_size > 0) struct array {
+    using scalar_t = _scalar_t;
+    using value_type = _scalar_t;
+    static constexpr std::size_t dimensions = _size;
+
+    COVFIE_DEVICE array() = default;
+
+    COVFIE_DEVICE array(const scalar_t (&arr)[dimensions])
+        requires(dimensions > 1)
+        : array(arr, std::make_index_sequence<dimensions>())
+    {
+    }
+
+    COVFIE_DEVICE array(const scalar_t & val)
+        : array(val, std::make_index_sequence<dimensions>())
+    {
+    }
+
+    template <typename... Ts>
+    requires(sizeof...(Ts) == dimensions) COVFIE_DEVICE array(Ts... args)
+        : m_data{std::forward<Ts>(args)...}
+    {
+    }
+
+    COVFIE_DEVICE constexpr scalar_t & at(const std::size_t & n)
+    {
+        assert(n < dimensions);
+
+        return m_data[n];
+    }
+
+    COVFIE_DEVICE constexpr const scalar_t & at(const std::size_t & n) const
+    {
+        assert(n < dimensions);
+
+        return m_data[n];
+    }
+
+    COVFIE_DEVICE constexpr scalar_t & operator[](const std::size_t & n)
+    {
+        assert(n < dimensions);
+
+        return m_data[n];
+    }
+
+    COVFIE_DEVICE constexpr const scalar_t & operator[](const std::size_t & n
+    ) const
+    {
+        assert(n < dimensions);
+
+        return m_data[n];
+    }
+
+    COVFIE_DEVICE constexpr std::size_t size() const
+    {
+        return dimensions;
+    }
+
+    COVFIE_DEVICE constexpr scalar_t * begin()
+    {
+        return m_data + 0;
+    }
+
+    COVFIE_DEVICE constexpr const scalar_t * begin() const
+    {
+        return m_data + 0;
+    }
+
+    COVFIE_DEVICE constexpr const scalar_t * cbegin() const
+    {
+        return m_data + 0;
+    }
+
+    COVFIE_DEVICE constexpr scalar_t * end()
+    {
+        return m_data + dimensions;
+    }
+
+    COVFIE_DEVICE constexpr const scalar_t * end() const
+    {
+        return m_data + dimensions;
+    }
+
+    COVFIE_DEVICE constexpr const scalar_t * cend() const
+    {
+        return m_data + dimensions;
+    }
+
+private:
+    template <std::size_t... Is>
+    COVFIE_DEVICE
+    array(const scalar_t (&arr)[dimensions], std::index_sequence<Is...>)
+        : m_data{arr[Is]...}
+    {
+    }
+
+    template <std::size_t... Is>
+    COVFIE_DEVICE array(const scalar_t & val, std::index_sequence<Is...>)
+        : m_data{((void)Is, val)...}
+    {
+    }
+
+    scalar_t m_data[dimensions];
+};
+}
diff --git a/lib/core/covfie/core/backend/primitive/array.hpp b/lib/core/covfie/core/backend/primitive/array.hpp
@@ -13,7 +13,6 @@
 #include <cassert>
 #include <cstring>
 #include <memory>
-#include <tuple>
 #include <utility>
 
 #include <covfie/core/concepts.hpp>

diff --git a/lib/core/covfie/core/backend/primitive/constant.hpp b/lib/core/covfie/core/backend/primitive/constant.hpp
@@ -10,7 +10,6 @@
 
 #pragma once
 
-#include <array>
 #include <iostream>
 
 #include <covfie/core/concepts.hpp>

diff --git a/lib/core/covfie/core/backend/transformer/hilbert.hpp b/lib/core/covfie/core/backend/transformer/hilbert.hpp
@@ -17,6 +17,7 @@
 #include <numeric>
 #include <type_traits>
 
+#include <covfie/core/array.hpp>
 #include <covfie/core/backend/primitive/array.hpp>
 #include <covfie/core/concepts.hpp>
 #include <covfie/core/parameter_pack.hpp>
@@ -25,7 +26,6 @@
 #include <covfie/core/utility/nd_map.hpp>
 #include <covfie/core/utility/nd_size.hpp>
 #include <covfie/core/utility/numeric.hpp>
-#include <covfie/core/utility/tuple.hpp>
 #include <covfie/core/vector.hpp>
 
 namespace covfie::backend {
@@ -118,20 +118,13 @@ struct hilbert {
             );
         typename T::parent_t::non_owning_data_t nother(other);
 
-        using tuple_t = decltype(std::tuple_cat(
-            std::declval<std::array<
-                typename contravariant_input_t::scalar_t,
-                contravariant_input_t::dimensions>>()
-        ));
-
-        utility::nd_map<tuple_t>(
-            [&nother, &res](tuple_t t) {
-                auto old_c = utility::to_array(t);
+        utility::nd_map<decltype(sizes)>(
+            [&nother, &res](decltype(sizes) t) {
                 coordinate_t c;
 
                 for (std::size_t i = 0; i < contravariant_input_t::dimensions;
                      ++i) {
-                    c[i] = old_c[i];
+                    c[i] = t[i];
                 }
 
                 std::size_t idx = calculate_index(c);
@@ -141,7 +134,7 @@ struct hilbert {
                     res[idx][i] = nother.at(c)[i];
                 }
             },
-            std::tuple_cat(sizes)
+            sizes
         );
 
         return res;

diff --git a/lib/core/covfie/core/backend/transformer/morton.hpp b/lib/core/covfie/core/backend/transformer/morton.hpp
@@ -32,7 +32,6 @@
 #include <covfie/core/utility/nd_map.hpp>
 #include <covfie/core/utility/nd_size.hpp>
 #include <covfie/core/utility/numeric.hpp>
-#include <covfie/core/utility/tuple.hpp>
 #include <covfie/core/vector.hpp>
 
 namespace covfie::backend {
@@ -165,18 +164,15 @@ struct morton {
             );
         typename T::parent_t::non_owning_data_t nother(other);
 
-        using tuple_t = decltype(std::tuple_cat(sizes));
-
-        utility::nd_map<tuple_t>(
-            [&nother, &res](tuple_t t) {
-                auto old_c = utility::to_array(t);
+        utility::nd_map<decltype(sizes)>(
+            [&nother, &res](decltype(sizes) t) {
                 typename contravariant_input_t::vector_t c;
 
                 for (std::size_t i = 0; i < contravariant_input_t::dimensions;
                      ++i) {
                     c[i] = static_cast<
                         typename contravariant_input_t::vector_t::value_type>(
-                        old_c[i]
+                        t[i]
                     );
                 }
 
@@ -187,7 +183,7 @@ struct morton {
                     res[idx][i] = nother.at(c)[i];
                 }
             },
-            std::tuple_cat(sizes)
+            sizes
         );
 
         return res;

diff --git a/lib/core/covfie/core/backend/transformer/shuffle.hpp b/lib/core/covfie/core/backend/transformer/shuffle.hpp
@@ -107,7 +107,7 @@ struct shuffle {
         shuffle(typename contravariant_input_t::vector_t c, std::index_sequence<Is...>)
             const
         {
-            return {std::get<Is>(c)...};
+            return {c.at(Is)...};
         }
 
         COVFIE_DEVICE typename covariant_output_t::vector_t