Refactor types to reduce duplication and cyclic dependencies (#1327)

* Move real2/size2 to geocel types
* Refactor LDG infrastructure to reduce circular dependencies

src/celeritas/grid/GridIdFinder.hh

@@ -10,6 +10,7 @@
 #include "corecel/Assert.hh"
 #include "corecel/Macros.hh"
 #include "corecel/cont/Span.hh"
+#include "corecel/data/LdgIterator.hh"
 #include "corecel/math/Algorithms.hh"
 
 namespace celeritas

src/corecel/OpaqueId.hh

@@ -17,12 +17,6 @@
 
 namespace celeritas
 {
-namespace detail
-{
-template<class, class>
-struct LdgLoader;
-}  // namespace detail
-
 //---------------------------------------------------------------------------//
 /*!
  * Type-safe index for accessing an array or collection of data.
@@ -103,6 +97,9 @@ class OpaqueId
     //! Get the value without checking for validity (atypical)
     CELER_CONSTEXPR_FUNCTION size_type unchecked_get() const { return value_; }
 
+    //! Access the underlying data for more efficient loading from memory
+    CELER_CONSTEXPR_FUNCTION size_type const* data() const { return &value_; }
+
   private:
     size_type value_;
 
@@ -113,7 +110,6 @@ class OpaqueId
     {
         return static_cast<size_type>(-1);
     }
-    friend detail::LdgLoader<OpaqueId<value_type, size_type> const, void>;
 };
 
 //---------------------------------------------------------------------------//
@@ -139,6 +135,7 @@ CELER_DEFINE_OPAQUEID_CMP(>=)
 
 #undef CELER_DEFINE_OPAQUEID_CMP
 
+//---------------------------------------------------------------------------//
 //! Allow less-than comparison with *integer* for container comparison
 template<class V, class S, class U>
 CELER_CONSTEXPR_FUNCTION bool operator<(OpaqueId<V, S> lhs, U rhs)
@@ -147,6 +144,7 @@ CELER_CONSTEXPR_FUNCTION bool operator<(OpaqueId<V, S> lhs, U rhs)
     return lhs && (U(lhs.unchecked_get()) < rhs);
 }
 
+//---------------------------------------------------------------------------//
 //! Allow less-than-equal comparison with *integer* for container comparison
 template<class V, class S, class U>
 CELER_CONSTEXPR_FUNCTION bool operator<=(OpaqueId<V, S> lhs, U rhs)
@@ -155,6 +153,7 @@ CELER_CONSTEXPR_FUNCTION bool operator<=(OpaqueId<V, S> lhs, U rhs)
     return lhs && (U(lhs.unchecked_get()) <= rhs);
 }
 
+//---------------------------------------------------------------------------//
 //! Get the distance between two opaque IDs
 template<class V, class S>
 inline CELER_FUNCTION S operator-(OpaqueId<V, S> self, OpaqueId<V, S> other)
@@ -164,6 +163,7 @@ inline CELER_FUNCTION S operator-(OpaqueId<V, S> self, OpaqueId<V, S> other)
     return self.unchecked_get() - other.unchecked_get();
 }
 
+//---------------------------------------------------------------------------//
 //! Increment an opaque ID by an offset
 template<class V, class S>
 inline CELER_FUNCTION OpaqueId<V, S>
@@ -174,6 +174,7 @@ operator+(OpaqueId<V, S> id, std::make_signed_t<S> offset)
     return OpaqueId<V, S>{id.unchecked_get() + static_cast<S>(offset)};
 }
 
+//---------------------------------------------------------------------------//
 //! Decrement an opaque ID by an offset
 template<class V, class S>
 inline CELER_FUNCTION OpaqueId<V, S>
@@ -184,25 +185,7 @@ operator-(OpaqueId<V, S> id, std::make_signed_t<S> offset)
     return OpaqueId<V, S>{id.unchecked_get() - static_cast<S>(offset)};
 }
 
-namespace detail
-{
-//! Template matching to determine if T is an OpaqueId
-template<class T>
-struct IsOpaqueId : std::false_type
-{
-};
-template<class V, class S>
-struct IsOpaqueId<OpaqueId<V, S>> : std::true_type
-{
-};
-template<class V, class S>
-struct IsOpaqueId<OpaqueId<V, S> const> : std::true_type
-{
-};
-template<class T>
-inline constexpr bool is_opaque_id_v = IsOpaqueId<T>::value;
 //---------------------------------------------------------------------------//
-}  // namespace detail
 }  // namespace celeritas
 
 //---------------------------------------------------------------------------//

src/corecel/cont/Span.hh

@@ -10,9 +10,8 @@
 #include <cstddef>
 #include <type_traits>
 
-#include "corecel/data/LdgIterator.hh"
-
 #include "Array.hh"
+
 #include "detail/SpanImpl.hh"
 
 namespace celeritas
@@ -189,10 +188,6 @@ class Span
 template<class T, std::size_t N>
 constexpr std::size_t Span<T, N>::extent;
 
-//! Alias for a Span iterating over values read using __ldg
-template<class T, std::size_t Extent = dynamic_extent>
-using LdgSpan = Span<LdgValue<T>, Extent>;
-
 //---------------------------------------------------------------------------//
 // FREE FUNCTIONS
 //---------------------------------------------------------------------------//
@@ -248,26 +243,5 @@ CELER_CONSTEXPR_FUNCTION auto make_array(Span<T, N> const& s)
     return result;
 }
 
-//---------------------------------------------------------------------------//
-/*!
- * Construct an array from a fixed-size span, removing LdgValue marker.
- *
- * Note: \code make_array(Span<T,N> const&) \endcode is not reused because:
- * 1. Using this overload reads input data using \c __ldg
- * 2. \code return make_array<T, N>(s) \endcode results in segfault (gcc 11.3).
- *    This might be a compiler bug because temporary lifetime should be
- *    extended until the end of the expression and we return a copy...
- */
-template<class T, std::size_t N>
-CELER_CONSTEXPR_FUNCTION auto make_array(LdgSpan<T, N> const& s)
-{
-    Array<std::remove_cv_t<T>, N> result{};
-    for (std::size_t i = 0; i < N; ++i)
-    {
-        result[i] = s[i];
-    }
-    return result;
-}
-
 //---------------------------------------------------------------------------//
 }  // namespace celeritas

src/corecel/cont/detail/SpanImpl.hh

@@ -13,10 +13,16 @@
 #include "corecel/Assert.hh"
 #include "corecel/Macros.hh"
 #include "corecel/OpaqueId.hh"
-#include "corecel/data/LdgIterator.hh"
 
 namespace celeritas
 {
+//---------------------------------------------------------------------------//
+template<class T>
+struct LdgValue;
+template<class T>
+class LdgIterator;
+
+//---------------------------------------------------------------------------//
 namespace detail
 {
 //---------------------------------------------------------------------------//
@@ -58,7 +64,7 @@ struct SpanTraits<LdgValue<T>>
 
 //---------------------------------------------------------------------------//
 //! Sentinel value for span of dynamic type
-constexpr std::size_t dynamic_extent = std::size_t(-1);
+inline constexpr std::size_t dynamic_extent = std::size_t(-1);
 
 //---------------------------------------------------------------------------//
 //! Calculate the return type for a subspan

src/corecel/data/LdgIterator.hh

@@ -13,6 +13,7 @@
 
 #include "corecel/Macros.hh"
 #include "corecel/Types.hh"
+#include "corecel/cont/Span.hh"
 #include "corecel/math/Algorithms.hh"
 
 #include "detail/LdgIteratorImpl.hh"
@@ -240,19 +241,13 @@ swap(LdgIterator<T>& lhs, LdgIterator<T>& rhs) noexcept
 }
 //!@}
 
-//!@{
-//! Helper
-template<class T>
-inline LdgIterator<T> make_ldg_iterator(T* ptr) noexcept
-{
-    return LdgIterator{ptr};
-}
-
+//---------------------------------------------------------------------------//
 /*!
- * Wrapper struct that containers can use to specialize on types supported by
- * LdgIterator, i.e. Span<LdgValue<T>> specialization can internally use
+ * Wrapper struct for specializing on types supported by LdgIterator.
+ *
+ * For example, Span<LdgValue<T>> specialization can internally use
  * LdgIterator. Specializations should refer to LdgValue<T>::value_type to
- * force the template instantiation of LdgValue and type-check T
+ * force the template instantiation of LdgValue and type-check T .
  */
 template<class T>
 struct LdgValue
@@ -262,7 +257,32 @@ struct LdgValue
                   "const arithmetic, OpaqueId or enum type "
                   "required");
 };
-//!@}
 
 //---------------------------------------------------------------------------//
-}  // namespace celeritas
+//! Alias for a Span iterating over values read using __ldg
+template<class T, std::size_t Extent = dynamic_extent>
+using LdgSpan = Span<LdgValue<T>, Extent>;
+
+//---------------------------------------------------------------------------//
+/*!
+ * Construct an array from a fixed-size span, removing LdgValue marker.
+ *
+ * Note: \code make_array(Span<T,N> const&) \endcode is not reused because:
+ * 1. Using this overload reads input data using \c __ldg
+ * 2. \code return make_array<T, N>(s) \endcode results in segfault (gcc 11.3).
+ *    This might be a compiler bug because temporary lifetime should be
+ *    extended until the end of the expression and we return a copy...
+ */
+template<class T, std::size_t N>
+CELER_CONSTEXPR_FUNCTION auto make_array(LdgSpan<T, N> const& s)
+{
+    Array<std::remove_cv_t<T>, N> result{};
+    for (std::size_t i = 0; i < N; ++i)
+    {
+        result[i] = s[i];
+    }
+    return result;
+}
+
+//---------------------------------------------------------------------------//
+}  // namespace celeritas

src/corecel/data/detail/CollectionImpl.hh

@@ -22,7 +22,7 @@
 #include "corecel/sys/Device.hh"
 
 #include "DisabledStorage.hh"
-#include "LdgIteratorImpl.hh"
+#include "TypeTraits.hh"
 #include "../Copier.hh"
 #include "../LdgIterator.hh"
 #include "../PinnedAllocator.hh"

src/corecel/data/detail/LdgIteratorImpl.hh

@@ -14,30 +14,48 @@
 #include "corecel/Types.hh"
 #include "corecel/math/Quantity.hh"
 
+#include "TypeTraits.hh"
+
 namespace celeritas
 {
 namespace detail
 {
+//---------------------------------------------------------------------------//
+/*!
+ * Wrap the low-level CUDA/HIP "load global memory" function.
+ *
+ * This low-level capability allows improved caching because we're \em
+ * promising that no other thread can modify its value while the kernel is
+ * active.
+ */
+template<class T>
+CELER_CONSTEXPR_FUNCTION T ldg(T const* ptr)
+{
+    static_assert(std::is_arithmetic_v<T>,
+                  "Only const arithmetic types are supported by __ldg");
+#if CELER_DEVICE_COMPILE
+    return __ldg(ptr);
+#else
+    return *ptr;
+#endif
+}
+
 //---------------------------------------------------------------------------//
 /*!
  * Reads a value T using __ldg builtin and return a copy of it
  */
 template<class T, typename = void>
 struct LdgLoader
 {
-    static_assert(std::is_arithmetic_v<T> && std::is_const_v<T>,
-                  "Only const arithmetic types are supported by __ldg");
+    static_assert(std::is_const_v<T>, "Only const data are supported by __ldg");
+
     using value_type = std::remove_const_t<T>;
     using pointer = std::add_pointer_t<value_type const>;
     using reference = value_type;
 
     CELER_CONSTEXPR_FUNCTION static reference read(pointer p)
     {
-#if CELER_DEVICE_COMPILE
-        return __ldg(p);
-#else
-        return *p;
-#endif
+        return ldg(p);
     }
 };
 
@@ -48,20 +66,13 @@ struct LdgLoader
 template<class I, class T>
 struct LdgLoader<OpaqueId<I, T> const, void>
 {
-    static_assert(std::is_arithmetic_v<T>,
-                  "OpaqueId needs to be indexed with a type supported by "
-                  "__ldg");
     using value_type = OpaqueId<I, T>;
     using pointer = std::add_pointer_t<value_type const>;
     using reference = value_type;
 
     CELER_CONSTEXPR_FUNCTION static reference read(pointer p)
     {
-#if CELER_DEVICE_COMPILE
-        return value_type{__ldg(&p->value_)};
-#else
-        return value_type{p->value_};
-#endif
+        return value_type{ldg(p->data())};
     }
 };
 
@@ -72,20 +83,13 @@ struct LdgLoader<OpaqueId<I, T> const, void>
 template<class I, class T>
 struct LdgLoader<Quantity<I, T> const, void>
 {
-    static_assert(std::is_arithmetic_v<T>,
-                  "Quantity needs to be represented by a type supported by "
-                  "__ldg");
     using value_type = Quantity<I, T>;
     using pointer = std::add_pointer_t<value_type const>;
     using reference = value_type;
 
     CELER_CONSTEXPR_FUNCTION static reference read(pointer p)
     {
-#if CELER_DEVICE_COMPILE
-        return value_type{__ldg(&p->value_)};
-#else
-        return value_type{p->value_};
-#endif
+        return value_type{ldg(p->data())};
     }
 };
 
@@ -103,21 +107,13 @@ struct LdgLoader<T const, std::enable_if_t<std::is_enum_v<T>>>
         // Technically breaks aliasing rule but it's not an issue:
         // the compiler doesn't derive any optimization and the pointer doesn't
         // escape the function
-        return value_type{__ldg(reinterpret_cast<underlying_type const*>(p))};
+        return value_type{ldg(reinterpret_cast<underlying_type const*>(p))};
 #else
         return *p;
 #endif
     }
 };
 
-//---------------------------------------------------------------------------//
-//! True if T is supported by a LdgLoader specialization
-template<class T>
-inline constexpr bool is_ldg_supported_v
-    = std::is_const_v<T>
-      && (std::is_arithmetic_v<T> || is_opaque_id_v<T> || is_quantity_v<T>
-          || std::is_enum_v<T>);
-
 //---------------------------------------------------------------------------//
 }  // namespace detail
 }  // namespace celeritas