This will make mdspan and its tests work with HIP #233
Conversation
Note we need the CTAD stuff because clang wants to see device markup on deduction guides.
Compiler errors out when calling std::swap. Note since it was constexpr function it was apparently automatically host device, but the tuple one isn't.
This test doesn't fully work due to using std::vector
| @@ -115,11 +115,15 @@ static_assert(_MDSPAN_CPLUSPLUS >= MDSPAN_CXX_STD_14, "mdspan requires C++14 or | |||
| # define _MDSPAN_NO_UNIQUE_ADDRESS | |||
| #endif | |||
|
|
|||
| // AMDs HIP compiler seems to have issues with concepts | |||
| // it pretends concepts exist, but doesn't ship <concept> | |||
| #ifndef __HIPCC__ | |||
There was a problem hiding this comment.
But that would be defined regardless of whether it is amdclang++ or some other HIP-capable compiler.
Should we guard depending on whether HIP was enabled at configuration time?
(not asking you to change, wondering what the right thing to do might be)
I see that we don't have such a "HIP enabled" macro define anyway.
There was a problem hiding this comment.
No for now whether HIPCC is defined is good enough. We can revisit later and refine based on compiler version when we have evidence which ones work.
| #endif | ||
|
|
||
| #ifdef _MDSPAN_HAS_HIP | ||
| # define MDSPAN_CTAD _MDSPAN_HOST_DEVICE |
There was a problem hiding this comment.
I'd prefer CTAD -> DEDUCTION_GUIDE
| @@ -321,9 +321,16 @@ class mdspan | |||
|
|
|||
| MDSPAN_INLINE_FUNCTION | |||
| friend constexpr void swap(mdspan& x, mdspan& y) noexcept { | |||
| // can't call the std::swap inside on HIP | |||
| #ifndef _MDSPAN_HAS_HIP | |||
There was a problem hiding this comment.
is that a FIXME? do you want to annotate with something greppable?
There was a problem hiding this comment.
not really. There is nothing to fix, its just what we have to do.
There was a problem hiding this comment.
do we consider it as a defect in the toolchain?
There was a problem hiding this comment.
Not really. I mean no one says we can call random std:: function on device ...
There was a problem hiding this comment.
That said it would be great if we could :-)
tests/test_mdspan_swap.cpp
Outdated
| @@ -102,7 +102,9 @@ TEST(TestMDSpanSwap, std_swap_dynamic_extents) { | |||
| __MDSPAN_TESTS_RUN_TEST(test_mdspan_std_swap_dynamic_extents()) | |||
| } | |||
|
|
|||
|
|
|||
| // On HIP we actually don't call through to swap via ADL | |||
| // so the foo swap test which has sideeffects will fail | |||
There was a problem hiding this comment.
sideeffects -> side effects
Consider annotating here too
There was a problem hiding this comment.
as I said I don't think there is anything to fix?
| template<class T> | ||
| void mallocManaged(T** ptr, size_t size) { (void) cudaMallocManaged(ptr, size); } | ||
| template<class T> | ||
| void freeManaged(T* ptr) { (void) cudaFree(ptr); } |
There was a problem hiding this comment.
I don't like the name but I suppose you were afraid of ambiguities?
There was a problem hiding this comment.
partly, partly making clear that this is managed memory equivalent and that it needs to be that, because we use the allocations from both host and device in the tests.
Note I didn't address benchmarks and examples, just tests. Also the biggest change is a patch gtest so it compiles with that clang version with C++20.