reconstruct OpaqueId from pointee internal value

celeritas-project · Oct 25, 2023 · f41f64a · f41f64a
1 parent 101bdd9
commit f41f64a
Show file tree

Hide file tree

Showing 4 changed files with 21 additions and 10 deletions.
diff --git a/src/corecel/OpaqueId.hh b/src/corecel/OpaqueId.hh
@@ -17,6 +17,12 @@
 
 namespace celeritas
 {
+namespace detail
+{
+template<class>
+struct LdgLoader;
+}  // namespace detail
+
 //---------------------------------------------------------------------------//
 /*!
  * Type-safe index for accessing an array.
@@ -30,7 +36,8 @@ namespace celeritas
 template<class ValueT, class SizeT = ::celeritas::size_type>
 class OpaqueId
 {
-    static_assert(static_cast<SizeT>(-1) > 0, "SizeT must be unsigned.");
+    static_assert(std::is_unsigned_v<SizeT> && !std::is_same_v<SizeT, bool>,
+                  "SizeT must be unsigned.");
 
   public:
     //!@{
@@ -91,6 +98,7 @@ class OpaqueId
     {
         return static_cast<size_type>(-1);
     }
+    friend detail::LdgLoader<OpaqueId<value_type, size_type>>;
 };
 
 //---------------------------------------------------------------------------//

diff --git a/src/corecel/data/LdgIterator.hh b/src/corecel/data/LdgIterator.hh
@@ -214,7 +214,7 @@ swap(LdgIterator<T>& lhs, LdgIterator<T>& rhs) noexcept
 template<class T>
 inline LdgIterator<T> make_ldg_iterator(T const* ptr) noexcept
 {
-    return LdgIterator<T>{ptr};
+    return LdgIterator{ptr};
 }
 //!@}
 

diff --git a/src/corecel/data/detail/LdgIteratorImpl.hh b/src/corecel/data/detail/LdgIteratorImpl.hh
@@ -46,16 +46,17 @@ struct LdgLoader
 template<class I, class T>
 struct LdgLoader<OpaqueId<I, T>>
 {
+    static_assert(std::is_arithmetic_v<T>);
     using value_type = OpaqueId<I, T>;
-    using pointer = std::add_pointer_t<T const>;
+    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)};
+        return value_type{__ldg(&p->value_)};
 #else
-        return value_type{*p};
+        return value_type{p->value_};
 #endif
     }
 };

diff --git a/test/corecel/data/LdgIterator.test.cc b/test/corecel/data/LdgIterator.test.cc
@@ -73,13 +73,15 @@ TEST_F(LdgIteratorTest, arithmetic_t)
 TEST_F(LdgIteratorTest, opaqueid_t)
 {
     using TestId = OpaqueId<struct LdgIteratorOpaqueIdTest_>;
-    using VecId = std::vector<typename TestId::size_type>;
-    VecId some_data = {1, 2, 3, 4};
+    using VecId = std::vector<TestId>;
+    VecId some_data = {TestId{1}, TestId{2}, TestId{3}, TestId{4}};
     auto n = some_data.size();
-    auto ldg_start = LdgIterator<TestId>(some_data.data());
-    auto ldg_end = LdgIterator<TestId>(some_data.data() + n);
+    auto ldg_start = make_ldg_iterator(some_data.data());
+    auto ldg_end = make_ldg_iterator(some_data.data() + n);
+    LdgIterator ctad_itr{some_data.data()};
+    EXPECT_TRUE((std::is_same_v<decltype(ctad_itr), decltype(ldg_start)>));
     using ptr_type = typename decltype(ldg_start)::pointer;
-    EXPECT_TRUE((std::is_same_v<ptr_type, typename TestId::size_type const*>));
+    EXPECT_TRUE((std::is_same_v<ptr_type, TestId const*>));
     EXPECT_TRUE(ldg_start);
     EXPECT_EQ(static_cast<ptr_type>(ldg_start), some_data.data());
     EXPECT_EQ(*ldg_start++, TestId{1});