Implement LdgIterator abstraction

[esseivaju]

Implement iterator for read-only device data using __ldg primitive

[sethrj]

Do we want to do a single PR that includes both the new LDG implementation and its use in Collection<T, const_reference>?

[esseivaju]

I'd do both in the same PR. The iterator implementation should be ready for review, I'll update the collection.

[esseivaju]

since __ldg should only ever return constant primitive type it should be cheaper to return a copy instead of const ref in most cases?

[sethrj]

I would think the compiler would optimize to the same thing if it returns an lvalue or not... but isn't the signature T __ldg(T*) i.e. it always returns an rvalue?

[esseivaju]

yeah, a function call that returns a non-reference is rvalue. In that case:

T const& f(T const* p)
{
    return __ldg(p);
}

I'd assume we bind a reference to temporary which results in UB?

[sethrj]

Since this isn't trying to conform to any existing interface, perhaps LdgLoader with operator()? Accessing with an instance could make sense if we extend LDG loading to "read member data from a struct" (e.g., for struct T, the loader could hold U T::* as class data and read member data from it using ldg?) I think this and the traits class should be in data/detail/LdgIteratorImpl.hh since they're more than a couple of lines.

[esseivaju]

I think it's easier to make this class stateless otherwise the LdgIterator has to make sure they both stay consistent.

[sethrj]

Maybe use make_ldg_iterator in one place and CTAD in the other? And then add

    EXPECT_TRUE((std::is_same_v<decltype(ldg_start), decltype(ldg_end)>));

[esseivaju]

can we use CTAD here? We have to somehow tell if the pointer points to regular builtin type or if it should reconstruct an OpaqueId.

[sethrj]

🤔 but that's just about how it's loaded under the hood right? Both should be LdgIterator<T> and the under-the-hood loader handles the reconstruction implementation (and value type).

[sethrj]

Oh wait, I see: your underlying data is size_type not OpaqueId. We have some collections of OpaqueId itself, and I was thinking we __ldg the underlying data. (Do we need to add a data() accessor to OpaqueId, or is it permissible to use a cast?)

[esseivaju]

Oh I see, if we want to do that we need the address of OpaqueId::value_ which we probably don't want to expose using an accessor, maybe declare the LdgLoader as friend.

[sethrj]

Good idea. As I said, I'm fine postponing that stuff to another PR if it increases the scope too much.

[esseivaju]

Looks like it's actually not that much change, done in f41f64a so it's probably fine to do it here. Maybe separate the collection update in another PR.

[sethrj]

We could also postpone the OpaqueId stuff until later: being able to do ldg on native simple types would be a great start I think.

[sethrj]

LGTM now! Do you want to merge with the test as is (backing out the Collection header include) or implement an LdgSpan for Collection<T, device, const_reference> ?

Also, do we want to require LdgIterator work with LDG-compatible data (const and fundamental type) or fall back to not using LDG? I'm thinking the former.

[esseivaju]

Let's do the LdgSpan and Collection in a separate PR. I agree, we should require LGD-compatible data. I've updated how constness is specified. Falling back to not using LDG can lead to misusing the iterator.

[esseivaju]

I'm done pushing to this branch unless you have more comments.

[esseivaju]

Realised that I was missing comparison operators while working on the LdgSpan

[sethrj]

Great! Even though the const requirement makes testing a little more awkward, it's the right thing to do from a device perspective: we never want to accidentally use LDG on non-const data.