Add cache hints

[j-stephan]

This is a followup to #18. We require portable load/store functionality with cache hints (such as __ldg or __stwb on CUDA). This should integrate nicely with #1249 and can probably be solved in the same or a followup PR.

[bernhardmgruber]

Btw, this can be handled really nice with std::mdspan and LLAMA accessors.

[j-stephan]

Since alpaka is about being explicit: Maybe we should add alpaka "intrinsics" that emulate the behavior of __ldg and friends?

[fwyzard]

A comment based on our recent experience: __ldg has the downside of always returning a copy, not a const reference. So it cannot be used with an interface like:

T& x() {
  return *px;
}

T const& x() const {
  return __ldg(px);
}

It can be used like

T x() const {
  return __ldg(px);
}

but that can be expensive (or wrong) for complex types, and cannot be used to get the pointer to px.

[bernhardmgruber]

I am really not sure whether we can provide/define a cross platform behavior of architecture specific intrinsics. Taking __ldg as an example: It is used for read-only data loads, which may go through a different cache hierarchy and have lower latency. What would be the equivalent for the CPU? x86 has the MOVNTDQA instruction, for streaming loads, but is that actually the same? And what about other backends?

[j-stephan]

Intel has a SYCL extension for FPGAs called "load-store units". See here. Maybe it would be worthwhile to copy the concept into alpaka where it could look like this:

// We are inside an alpaka kernel
using ReadOnlyLSU = alpaka::LoadStoreUnit<alpaka::ReadOnly>;

auto val = ReadOnlyLSU::load(some_ptr); // If using CUDA call __ldg() underneath. Otherwise perform a normal load if there is no equivalent.
ReadOnlyLSU::store(some_ptr, val); // This should cause a compile-time error.

So it cannot be used with an interface like:

T& x() {
    return *px;
}

T const& x() const {
    return __ldg(px);
}

Wouldn't that be kind of illegal in CUDA? __ldg requires that the memory in question is read-only for the lifetime of the kernel. With an interface like that you couldn't guarantee that the user won't overwrite the value.

It can be used like

T x() const {
    return __ldg(px);
}

but that can be expensive (or wrong) for complex types, and cannot be used to get the pointer to px.

__ldg() is restricted to fundamental scalar and vector types IIRC. If it gets more complex the compiler would probably complain anyway.

[fwyzard]

The point was to give a small, self-contained example, that shows the problem with the interface of __ldg. The actual implementation is significantly more complex: .

Wouldn't that be kind of illegal in CUDA? __ldg requires that the memory in question is read-only for the lifetime of the kernel. With an interface like that you couldn't guarantee that the user won't overwrite the value.

Actually we can, because the object is either a mutable View or a ConstView, at the kernel level.

__ldg() is restricted to fundamental scalar and vector types IIRC. If it gets more complex the compiler would probably complain anyway.

Whether __ldg is actually used internally or not can depend on the type:

T x() const {
    if constexpr(can_use_ldg<T>) {
        return __ldg(px);
    } else {
        return *px;
    }
}

But the interface should be the same for all types, and using __ldg() for the fundamental types prevents returning the complex ones by reference.

So far we have managed to make things work using the __restrict__ qualifier instead of explicitly using __ldg.