Add HIP logic, move to consistent naming convention (#53)

* Add HIP logic, move to same naming convention as other traccc projects (HOST_DEVICE everywhere) * Formatting/linting --------- Co-authored-by: Stewart Martin-Haugh <smh@cern.ch>
acts-project · Feb 13, 2025 · 1507359 · 1507359
1 parent ebd384b
commit 1507359
Show file tree

Hide file tree

Showing 23 changed files with 70 additions and 78 deletions.
diff --git a/benchmarks/common/covfie/benchmark/lorentz.hpp b/benchmarks/common/covfie/benchmark/lorentz.hpp
@@ -10,15 +10,10 @@
 #include <covfie/benchmark/types.hpp>
 #include <covfie/core/definitions.hpp>
 #include <covfie/core/field_view.hpp>
-
-#ifdef __CUDACC__
-#define HOST_DEVICE __host__ __device__
-#else
-#define HOST_DEVICE
-#endif
+#include <covfie/core/qualifiers.hpp>
 
 template <typename propagator_t, typename backend_t>
-HOST_DEVICE inline __attribute__((always_inline)) void lorentz_step(
+COVFIE_HOST_DEVICE inline __attribute__((always_inline)) void lorentz_step(
     const covfie::field_view<backend_t> & f,
     covfie::benchmark::lorentz_agent<3> & o,
     float s

diff --git a/benchmarks/common/covfie/benchmark/propagate.hpp b/benchmarks/common/covfie/benchmark/propagate.hpp
@@ -10,12 +10,6 @@
 #include <covfie/core/definitions.hpp>
 #include <covfie/core/field_view.hpp>
 
-#ifdef __CUDACC__
-#define HOST_DEVICE __host__ __device__
-#else
-#define HOST_DEVICE
-#endif
-
 class Wide
 {
 };
@@ -33,7 +27,7 @@ class RungeKutta4
 };
 
 template <typename propagator_t, typename backend_t>
-HOST_DEVICE inline __attribute__((always_inline)) void propagation_step(
+COVFIE_HOST_DEVICE inline __attribute__((always_inline)) void propagation_step(
     const covfie::field_view<backend_t> & f,
     covfie::benchmark::propagation_agent<3> & o,
     float s

diff --git a/lib/core/covfie/core/algebra/affine.hpp b/lib/core/covfie/core/algebra/affine.hpp
@@ -22,12 +22,13 @@ struct affine : public matrix<N, N + 1, T, I> {
     affine & operator=(const affine &) = default;
     affine & operator=(affine &&) = default;
 
-    COVFIE_DEVICE affine(const matrix<N, N + 1, T, I> & o)
+    COVFIE_HOST_DEVICE affine(const matrix<N, N + 1, T, I> & o)
         : matrix<N, N + 1, T, I>(o)
     {
     }
 
-    COVFIE_DEVICE vector<N, T, I> operator*(const vector<N, T, I> & v) const
+    COVFIE_HOST_DEVICE vector<N, T, I> operator*(const vector<N, T, I> & v
+    ) const
     {
         vector<N + 1, T, I> r;
 
@@ -40,7 +41,8 @@ struct affine : public matrix<N, N + 1, T, I> {
         return matrix<N, N + 1, T, I>::operator*(r);
     }
 
-    COVFIE_DEVICE affine<N, T, I> operator*(const affine<N, T, I> & m) const
+    COVFIE_HOST_DEVICE affine<N, T, I> operator*(const affine<N, T, I> & m
+    ) const
     {
         matrix<N + 1, N + 1, T, I> m1, m2;
 
@@ -74,7 +76,7 @@ struct affine : public matrix<N, N + 1, T, I> {
     }
 
     template <typename... Args>
-    COVFIE_DEVICE static affine<N, T, I> translation(const Args &... args)
+    COVFIE_HOST_DEVICE static affine<N, T, I> translation(const Args &... args)
     {
         static_assert(
             (std::is_convertible_v<Args, T> && ...),
@@ -99,7 +101,7 @@ struct affine : public matrix<N, N + 1, T, I> {
     }
 
     template <typename... Args>
-    COVFIE_DEVICE static affine<N, T, I> scaling(const Args &... args)
+    COVFIE_HOST_DEVICE static affine<N, T, I> scaling(const Args &... args)
     {
         static_assert(
             (std::is_convertible_v<Args, T> && ...),

diff --git a/lib/core/covfie/core/algebra/matrix.hpp b/lib/core/covfie/core/algebra/matrix.hpp
@@ -18,11 +18,11 @@ template <
     typename T = float,
     typename I = std::size_t>
 struct matrix {
-    COVFIE_DEVICE matrix()
+    COVFIE_HOST_DEVICE matrix()
     {
     }
 
-    COVFIE_DEVICE matrix(array::array<array::array<T, M>, N> l)
+    COVFIE_HOST_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) {
@@ -33,18 +33,18 @@ struct matrix {
 
     matrix(const matrix<N, M, T, I> &) = default;
 
-    COVFIE_DEVICE T operator()(const I i, const I j) const
+    COVFIE_HOST_DEVICE T operator()(const I i, const I j) const
     {
         return m_elems[i][j];
     }
 
-    COVFIE_DEVICE T & operator()(const I i, const I j)
+    COVFIE_HOST_DEVICE T & operator()(const I i, const I j)
     {
         return m_elems[i][j];
     }
 
     template <std::size_t P>
-    COVFIE_DEVICE matrix<N, P, T, I> operator*(const matrix<M, P, T, I> & o
+    COVFIE_HOST_DEVICE matrix<N, P, T, I> operator*(const matrix<M, P, T, I> & o
     ) const
     {
         matrix<N, P, T, I> r;
@@ -64,7 +64,7 @@ struct matrix {
         return r;
     }
 
-    COVFIE_DEVICE static matrix<N, M, T, I> identity()
+    COVFIE_HOST_DEVICE static matrix<N, M, T, I> identity()
     {
         matrix<N, M, T, I> result;
 

diff --git a/lib/core/covfie/core/algebra/vector.hpp b/lib/core/covfie/core/algebra/vector.hpp
@@ -18,20 +18,20 @@
 namespace covfie::algebra {
 template <std::size_t N, typename T = float, typename I = std::size_t>
 struct vector : public matrix<N, 1, T, I> {
-    COVFIE_DEVICE vector()
+    COVFIE_HOST_DEVICE vector()
         : matrix<N, 1, T, I>()
     {
     }
 
-    COVFIE_DEVICE vector(array::array<T, N> l)
+    COVFIE_HOST_DEVICE vector(array::array<T, N> l)
         : matrix<N, 1, T, I>()
     {
         for (I i = 0; i < N; ++i) {
             matrix<N, 1, T, I>::operator()(i, 0) = l[i];
         }
     }
 
-    COVFIE_DEVICE vector(const matrix<N, 1, T, I> & o)
+    COVFIE_HOST_DEVICE vector(const matrix<N, 1, T, I> & o)
         : matrix<N, 1, T, I>(o)
     {
     }
@@ -43,17 +43,17 @@ struct vector : public matrix<N, 1, T, I> {
                 std::conjunction_v<std::is_convertible<Args, T>...> &&
                 sizeof...(Args) == N,
             bool> = true>
-    COVFIE_DEVICE vector(Args... args)
+    COVFIE_HOST_DEVICE vector(Args... args)
         : vector(array::array<T, N>{std::forward<Args>(args)...})
     {
     }
 
-    COVFIE_DEVICE T operator()(const I & i) const
+    COVFIE_HOST_DEVICE T operator()(const I & i) const
     {
         return matrix<N, 1, T, I>::operator()(i, 0);
     }
 
-    COVFIE_DEVICE T & operator()(const I & i)
+    COVFIE_HOST_DEVICE T & operator()(const I & i)
     {
         return matrix<N, 1, T, I>::operator()(i, 0);
     }

diff --git a/lib/core/covfie/core/array.hpp b/lib/core/covfie/core/array.hpp
@@ -19,99 +19,100 @@ requires(_size > 0) struct array {
     using value_type = _scalar_t;
     static constexpr std::size_t dimensions = _size;
 
-    COVFIE_DEVICE array() = default;
+    COVFIE_HOST_DEVICE array() = default;
 
-    COVFIE_DEVICE array(const scalar_t (&arr)[dimensions])
+    COVFIE_HOST_DEVICE array(const scalar_t (&arr)[dimensions])
         requires(dimensions > 1)
         : array(arr, std::make_index_sequence<dimensions>())
     {
     }
 
-    COVFIE_DEVICE array(const scalar_t & val)
+    COVFIE_HOST_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)
+    requires(sizeof...(Ts) == dimensions) COVFIE_HOST_DEVICE array(Ts... args)
         : m_data{std::forward<Ts>(args)...}
     {
     }
 
-    COVFIE_DEVICE constexpr scalar_t & at(const std::size_t & n)
+    COVFIE_HOST_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
+    COVFIE_HOST_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)
+    COVFIE_HOST_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
+    COVFIE_HOST_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
+    COVFIE_HOST_DEVICE constexpr std::size_t size() const
     {
         return dimensions;
     }
 
-    COVFIE_DEVICE constexpr scalar_t * begin()
+    COVFIE_HOST_DEVICE constexpr scalar_t * begin()
     {
         return m_data + 0;
     }
 
-    COVFIE_DEVICE constexpr const scalar_t * begin() const
+    COVFIE_HOST_DEVICE constexpr const scalar_t * begin() const
     {
         return m_data + 0;
     }
 
-    COVFIE_DEVICE constexpr const scalar_t * cbegin() const
+    COVFIE_HOST_DEVICE constexpr const scalar_t * cbegin() const
     {
         return m_data + 0;
     }
 
-    COVFIE_DEVICE constexpr scalar_t * end()
+    COVFIE_HOST_DEVICE constexpr scalar_t * end()
     {
         return m_data + dimensions;
     }
 
-    COVFIE_DEVICE constexpr const scalar_t * end() const
+    COVFIE_HOST_DEVICE constexpr const scalar_t * end() const
     {
         return m_data + dimensions;
     }
 
-    COVFIE_DEVICE constexpr const scalar_t * cend() const
+    COVFIE_HOST_DEVICE constexpr const scalar_t * cend() const
     {
         return m_data + dimensions;
     }
 
 private:
     template <std::size_t... Is>
-    COVFIE_DEVICE
+    COVFIE_HOST_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...>)
+    COVFIE_HOST_DEVICE array(const scalar_t & val, std::index_sequence<Is...>)
         : m_data{((void)Is, val)...}
     {
     }

diff --git a/lib/core/covfie/core/backend/primitive/array.hpp b/lib/core/covfie/core/backend/primitive/array.hpp
@@ -207,7 +207,7 @@ struct array {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at(typename contravariant_input_t::vector_t i) const
         {
             assert(i < m_size);

diff --git a/lib/core/covfie/core/backend/primitive/constant.hpp b/lib/core/covfie/core/backend/primitive/constant.hpp
@@ -96,7 +96,7 @@ struct constant {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
             at(typename contravariant_input_t::vector_t) const
         {
             return m_value;

diff --git a/lib/core/covfie/core/backend/transformer/affine.hpp b/lib/core/covfie/core/backend/transformer/affine.hpp
@@ -136,7 +136,7 @@ struct affine {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at(typename contravariant_input_t::vector_t c) const
         {
             covfie::algebra::vector<

diff --git a/lib/core/covfie/core/backend/transformer/backup.hpp b/lib/core/covfie/core/backend/transformer/backup.hpp
@@ -171,7 +171,7 @@ struct backup {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at(typename contravariant_input_t::vector_t coord) const
         {
             for (std::size_t i = 0; i < contravariant_input_t::dimensions; ++i)

diff --git a/lib/core/covfie/core/backend/transformer/clamp.hpp b/lib/core/covfie/core/backend/transformer/clamp.hpp
@@ -144,14 +144,14 @@ struct clamp {
         }
 
         template <std::size_t... Is>
-        COVFIE_DEVICE typename contravariant_input_t::vector_t
+        COVFIE_HOST_DEVICE typename contravariant_input_t::vector_t
         adjust(typename contravariant_input_t::vector_t coord, std::index_sequence<Is...>)
             const
         {
             return {std::clamp(coord[Is], m_min[Is], m_max[Is])...};
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at(typename contravariant_input_t::vector_t coord) const
         {
             return m_backend.at(adjust(

diff --git a/lib/core/covfie/core/backend/transformer/covariant_cast.hpp b/lib/core/covfie/core/backend/transformer/covariant_cast.hpp
@@ -101,14 +101,14 @@ struct covariant_cast {
         }
 
         template <std::size_t... Is>
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at_helper(typename contravariant_input_t::vector_t c, std::index_sequence<Is...>)
             const
         {
             return {static_cast<target_type>(m_backend.at(c)[Is])...};
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at(typename contravariant_input_t::vector_t c) const
         {
             return at_helper(

diff --git a/lib/core/covfie/core/backend/transformer/dereference.hpp b/lib/core/covfie/core/backend/transformer/dereference.hpp
@@ -99,7 +99,7 @@ struct dereference {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
         at(typename contravariant_input_t::vector_t c) const
         {
             return m_backend.at(c);