From 0b0388eb3d3ebaf92f15c1fec247f211d779fbce Mon Sep 17 00:00:00 2001
From: Stewart Martin-Haugh <smh@cern.ch>
Date: Tue, 11 Feb 2025 21:42:05 +0100
Subject: [PATCH 1/2] Add HIP logic, move to same naming convention as other
 traccc projects (HOST_DEVICE everywhere)

---
 .../common/covfie/benchmark/lorentz.hpp       |  8 ++---
 .../common/covfie/benchmark/propagate.hpp     |  8 +----
 lib/core/covfie/core/algebra/affine.hpp       | 10 +++---
 lib/core/covfie/core/algebra/matrix.hpp       | 12 +++----
 lib/core/covfie/core/algebra/vector.hpp       | 12 +++----
 lib/core/covfie/core/array.hpp                | 34 +++++++++----------
 .../covfie/core/backend/primitive/array.hpp   |  2 +-
 .../core/backend/primitive/constant.hpp       |  2 +-
 .../core/backend/transformer/affine.hpp       |  2 +-
 .../core/backend/transformer/backup.hpp       |  2 +-
 .../covfie/core/backend/transformer/clamp.hpp |  4 +--
 .../backend/transformer/covariant_cast.hpp    |  4 +--
 .../core/backend/transformer/dereference.hpp  |  2 +-
 .../core/backend/transformer/hilbert.hpp      |  6 ++--
 .../core/backend/transformer/linear.hpp       |  4 +--
 .../core/backend/transformer/morton.hpp       |  4 +--
 .../backend/transformer/nearest_neighbour.hpp |  2 +-
 .../core/backend/transformer/shuffle.hpp      |  4 +--
 .../core/backend/transformer/strided.hpp      |  2 +-
 lib/core/covfie/core/field_view.hpp           |  4 +--
 lib/core/covfie/core/qualifiers.hpp           |  6 ++--
 .../backend/primitive/cuda_device_array.hpp   |  2 +-
 .../cuda/backend/primitive/cuda_texture.hpp   |  2 +-
 23 files changed, 64 insertions(+), 74 deletions(-)

diff --git a/benchmarks/common/covfie/benchmark/lorentz.hpp b/benchmarks/common/covfie/benchmark/lorentz.hpp
index 38a97d3..34b1ab2 100644
--- a/benchmarks/common/covfie/benchmark/lorentz.hpp
+++ b/benchmarks/common/covfie/benchmark/lorentz.hpp
@@ -9,16 +9,12 @@
 #include <covfie/benchmark/propagate.hpp>
 #include <covfie/benchmark/types.hpp>
 #include <covfie/core/definitions.hpp>
+#include <covfie/core/qualifiers.hpp>
 #include <covfie/core/field_view.hpp>
 
-#ifdef __CUDACC__
-#define HOST_DEVICE __host__ __device__
-#else
-#define HOST_DEVICE
-#endif
 
 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
index 574a9bc..502e5b6 100644
--- 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
index 5407fda..7773320 100644
--- a/lib/core/covfie/core/algebra/affine.hpp
+++ b/lib/core/covfie/core/algebra/affine.hpp
@@ -22,12 +22,12 @@ 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 +40,7 @@ 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 +74,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 +99,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
index d1e61e0..c841a2d 100644
--- 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
index c9cc578..f941b81 100644
--- a/lib/core/covfie/core/algebra/vector.hpp
+++ b/lib/core/covfie/core/algebra/vector.hpp
@@ -18,12 +18,12 @@
 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) {
@@ -31,7 +31,7 @@ struct vector : public matrix<N, 1, T, 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
index 9acf7a6..8c304ed 100644
--- a/lib/core/covfie/core/array.hpp
+++ b/lib/core/covfie/core/array.hpp
@@ -19,47 +19,47 @@ 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
+    COVFIE_HOST_DEVICE constexpr const scalar_t & operator[](const std::size_t & n
     ) const
     {
         assert(n < dimensions);
@@ -67,51 +67,51 @@ requires(_size > 0) struct array {
         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
index 404ee7e..7f4a993 100644
--- 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
index 6ccae51..958b3b8 100644
--- 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
index cc5d7e9..9627a01 100644
--- 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
index dbfe4a4..7c65a19 100644
--- 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
index 2e51641..f564b0e 100644
--- 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
index 286d71e..a49c0a4 100644
--- 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
index a84de9f..0b1dd70 100644
--- 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);
diff --git a/lib/core/covfie/core/backend/transformer/hilbert.hpp b/lib/core/covfie/core/backend/transformer/hilbert.hpp
index 36f6981..f26b037 100644
--- a/lib/core/covfie/core/backend/transformer/hilbert.hpp
+++ b/lib/core/covfie/core/backend/transformer/hilbert.hpp
@@ -52,7 +52,7 @@ struct hilbert {
         "Number of dimensions for input must be exactly two."
     );
 
-    COVFIE_DEVICE static void
+    COVFIE_HOST_DEVICE static void
     rot(std::size_t n,
         std::size_t * x,
         std::size_t * y,
@@ -71,7 +71,7 @@ struct hilbert {
         }
     }
 
-    COVFIE_DEVICE static std::size_t calculate_index(
+    COVFIE_HOST_DEVICE static std::size_t calculate_index(
         coordinate_t c,
         utility::nd_size<contravariant_input_t::dimensions> sizes
     )
@@ -261,7 +261,7 @@ struct hilbert {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t at(coordinate_t c
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t at(coordinate_t c
         ) const
         {
 #ifndef NDEBUG
diff --git a/lib/core/covfie/core/backend/transformer/linear.hpp b/lib/core/covfie/core/backend/transformer/linear.hpp
index 9ff31bd..b0eb41b 100644
--- a/lib/core/covfie/core/backend/transformer/linear.hpp
+++ b/lib/core/covfie/core/backend/transformer/linear.hpp
@@ -147,7 +147,7 @@ struct linear {
         }
 
         template <std::size_t... Is>
-        COVFIE_DEVICE typename contravariant_output_t::vector_t
+        COVFIE_HOST_DEVICE typename contravariant_output_t::vector_t
         _backend_index_helper(typename contravariant_output_t::vector_t coord, std::size_t n, std::index_sequence<Is...>)
             const
         {
@@ -157,7 +157,7 @@ struct linear {
             )...};
         }
 
-        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
         {
             if constexpr (covariant_output_t::dimensions == 1) {
diff --git a/lib/core/covfie/core/backend/transformer/morton.hpp b/lib/core/covfie/core/backend/transformer/morton.hpp
index 741d5a3..fe36ef3 100644
--- a/lib/core/covfie/core/backend/transformer/morton.hpp
+++ b/lib/core/covfie/core/backend/transformer/morton.hpp
@@ -97,7 +97,7 @@ struct morton {
 
     static constexpr uint32_t IO_MAGIC_HEADER = 0xAB020006;
 
-    COVFIE_DEVICE static std::size_t
+    COVFIE_HOST_DEVICE static std::size_t
     calculate_index(typename contravariant_input_t::vector_t c)
     {
 #ifdef HAVE_BMI2
@@ -299,7 +299,7 @@ struct morton {
         {
         }
 
-        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
         {
 #ifndef NDEBUG
diff --git a/lib/core/covfie/core/backend/transformer/nearest_neighbour.hpp b/lib/core/covfie/core/backend/transformer/nearest_neighbour.hpp
index 45b42eb..f33f8fb 100644
--- a/lib/core/covfie/core/backend/transformer/nearest_neighbour.hpp
+++ b/lib/core/covfie/core/backend/transformer/nearest_neighbour.hpp
@@ -129,7 +129,7 @@ struct nearest_neighbour {
         {
         }
 
-        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
         {
             typename contravariant_output_t::vector_t nc;
diff --git a/lib/core/covfie/core/backend/transformer/shuffle.hpp b/lib/core/covfie/core/backend/transformer/shuffle.hpp
index 6479af1..9e2d7a8 100644
--- a/lib/core/covfie/core/backend/transformer/shuffle.hpp
+++ b/lib/core/covfie/core/backend/transformer/shuffle.hpp
@@ -99,14 +99,14 @@ struct shuffle {
         }
 
         template <std::size_t... Is>
-        COVFIE_DEVICE typename contravariant_input_t::vector_t
+        COVFIE_HOST_DEVICE typename contravariant_input_t::vector_t
         shuffle(typename contravariant_input_t::vector_t c, std::index_sequence<Is...>)
             const
         {
             return {c.at(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 m_backend.at(shuffle(c, indices{}));
diff --git a/lib/core/covfie/core/backend/transformer/strided.hpp b/lib/core/covfie/core/backend/transformer/strided.hpp
index 5c05f18..8796580 100644
--- a/lib/core/covfie/core/backend/transformer/strided.hpp
+++ b/lib/core/covfie/core/backend/transformer/strided.hpp
@@ -229,7 +229,7 @@ struct strided {
         {
         }
 
-        COVFIE_DEVICE typename covariant_output_t::vector_t at(coordinate_t c
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t at(coordinate_t c
         ) const
         {
             typename contravariant_input_t::scalar_t idx = 0;
diff --git a/lib/core/covfie/core/field_view.hpp b/lib/core/covfie/core/field_view.hpp
index 15b7533..e723517 100644
--- a/lib/core/covfie/core/field_view.hpp
+++ b/lib/core/covfie/core/field_view.hpp
@@ -49,7 +49,7 @@ class field_view
             sizeof...(Args) == backend_t::contravariant_input_t::dimensions,
             bool> = true,
         std::enable_if_t<!std::is_scalar_v<Q>, bool> = true>
-    COVFIE_DEVICE output_t at(Args... c) const
+    COVFIE_HOST_DEVICE output_t at(Args... c) const
     {
         return m_storage.at(coordinate_t{
             static_cast<typename backend_t::contravariant_input_t::scalar_t>(c
@@ -59,7 +59,7 @@ class field_view
     template <
         typename T,
         std::enable_if_t<std::is_same_v<T, coordinate_t>, bool> = true>
-    COVFIE_DEVICE output_t at(T c) const
+    COVFIE_HOST_DEVICE output_t at(T c) const
     {
         return m_storage.at(c);
     }
diff --git a/lib/core/covfie/core/qualifiers.hpp b/lib/core/covfie/core/qualifiers.hpp
index cf7c2d9..7a3e402 100644
--- a/lib/core/covfie/core/qualifiers.hpp
+++ b/lib/core/covfie/core/qualifiers.hpp
@@ -6,8 +6,8 @@
 
 #pragma once
 
-#if defined(__CUDACC__)
-#define COVFIE_DEVICE __host__ __device__
+#if defined(__CUDACC__) || defined(__HIP__)
+#define COVFIE_HOST_DEVICE __host__ __device__
 #else
-#define COVFIE_DEVICE
+#define COVFIE_HOST_DEVICE
 #endif
diff --git a/lib/cuda/covfie/cuda/backend/primitive/cuda_device_array.hpp b/lib/cuda/covfie/cuda/backend/primitive/cuda_device_array.hpp
index 7112a9c..5f32599 100644
--- a/lib/cuda/covfie/cuda/backend/primitive/cuda_device_array.hpp
+++ b/lib/cuda/covfie/cuda/backend/primitive/cuda_device_array.hpp
@@ -187,7 +187,7 @@ struct cuda_device_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
         {
             return m_ptr[i];
diff --git a/lib/cuda/covfie/cuda/backend/primitive/cuda_texture.hpp b/lib/cuda/covfie/cuda/backend/primitive/cuda_texture.hpp
index ea74cbb..001f429 100644
--- a/lib/cuda/covfie/cuda/backend/primitive/cuda_texture.hpp
+++ b/lib/cuda/covfie/cuda/backend/primitive/cuda_texture.hpp
@@ -232,7 +232,7 @@ struct cuda_texture {
         {
         }
 
-        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
         {
             channel_t r;

From 1e8d0eda65023bf974e55a903954e8465552fd55 Mon Sep 17 00:00:00 2001
From: Stewart Martin-Haugh <smh@cern.ch>
Date: Tue, 11 Feb 2025 21:47:20 +0100
Subject: [PATCH 2/2] Formatting/linting

---
 benchmarks/common/covfie/benchmark/lorentz.hpp       | 3 +--
 lib/core/covfie/core/algebra/affine.hpp              | 6 ++++--
 lib/core/covfie/core/array.hpp                       | 7 ++++---
 lib/core/covfie/core/backend/transformer/hilbert.hpp | 4 ++--
 lib/core/covfie/core/backend/transformer/strided.hpp | 4 ++--
 5 files changed, 13 insertions(+), 11 deletions(-)

diff --git a/benchmarks/common/covfie/benchmark/lorentz.hpp b/benchmarks/common/covfie/benchmark/lorentz.hpp
index 34b1ab2..c7e9e8c 100644
--- a/benchmarks/common/covfie/benchmark/lorentz.hpp
+++ b/benchmarks/common/covfie/benchmark/lorentz.hpp
@@ -9,9 +9,8 @@
 #include <covfie/benchmark/propagate.hpp>
 #include <covfie/benchmark/types.hpp>
 #include <covfie/core/definitions.hpp>
-#include <covfie/core/qualifiers.hpp>
 #include <covfie/core/field_view.hpp>
-
+#include <covfie/core/qualifiers.hpp>
 
 template <typename propagator_t, typename backend_t>
 COVFIE_HOST_DEVICE inline __attribute__((always_inline)) void lorentz_step(
diff --git a/lib/core/covfie/core/algebra/affine.hpp b/lib/core/covfie/core/algebra/affine.hpp
index 7773320..5830c9f 100644
--- a/lib/core/covfie/core/algebra/affine.hpp
+++ b/lib/core/covfie/core/algebra/affine.hpp
@@ -27,7 +27,8 @@ struct affine : public matrix<N, N + 1, T, I> {
     {
     }
 
-    COVFIE_HOST_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_HOST_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;
 
diff --git a/lib/core/covfie/core/array.hpp b/lib/core/covfie/core/array.hpp
index 8c304ed..b47d196 100644
--- a/lib/core/covfie/core/array.hpp
+++ b/lib/core/covfie/core/array.hpp
@@ -45,7 +45,8 @@ requires(_size > 0) struct array {
         return m_data[n];
     }
 
-    COVFIE_HOST_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);
 
@@ -59,8 +60,8 @@ requires(_size > 0) struct array {
         return m_data[n];
     }
 
-    COVFIE_HOST_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);
 
diff --git a/lib/core/covfie/core/backend/transformer/hilbert.hpp b/lib/core/covfie/core/backend/transformer/hilbert.hpp
index f26b037..086be2c 100644
--- a/lib/core/covfie/core/backend/transformer/hilbert.hpp
+++ b/lib/core/covfie/core/backend/transformer/hilbert.hpp
@@ -261,8 +261,8 @@ struct hilbert {
         {
         }
 
-        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t at(coordinate_t c
-        ) const
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
+        at(coordinate_t c) const
         {
 #ifndef NDEBUG
             for (std::size_t i = 0; i < contravariant_input_t::dimensions; ++i)
diff --git a/lib/core/covfie/core/backend/transformer/strided.hpp b/lib/core/covfie/core/backend/transformer/strided.hpp
index 8796580..d61df32 100644
--- a/lib/core/covfie/core/backend/transformer/strided.hpp
+++ b/lib/core/covfie/core/backend/transformer/strided.hpp
@@ -229,8 +229,8 @@ struct strided {
         {
         }
 
-        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t at(coordinate_t c
-        ) const
+        COVFIE_HOST_DEVICE typename covariant_output_t::vector_t
+        at(coordinate_t c) const
         {
             typename contravariant_input_t::scalar_t idx = 0;