diff --git a/.upstream-tests/test/cuda/exception/exception.pass.cpp b/.upstream-tests/test/cuda/exception/exception.pass.cpp new file mode 100644 index 0000000000..adfb8ff41c --- /dev/null +++ b/.upstream-tests/test/cuda/exception/exception.pass.cpp @@ -0,0 +1,58 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include +#include + +#include + +#ifndef _LIBCUDACXX_NO_EXCEPTIONS +void test_throw() { throw cuda::cuda_error{cudaErrorInvalidKernelImage}; } +#endif + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ + static_assert(std::is_base_of<::std::runtime_error, cuda::cuda_error>::value, + ""); + + cuda::cuda_error e{cudaErrorMemoryAllocation}; + assert(e.what() == std::string{"cudaErrorMemoryAllocation: out of memory"}); + assert(e.code() == cudaErrorMemoryAllocation); + + std::string msg = "test message"; + cuda::cuda_error e2{cudaErrorMemoryAllocation, msg}; + assert(e2.what() == msg + ": cudaErrorMemoryAllocation: out of memory"); + assert(e2.code() == cudaErrorMemoryAllocation); + +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + try { + test_throw(); + } catch (cuda::cuda_error const &e) { + assert(e.what() == + std::string{ + "cudaErrorInvalidKernelImage: device kernel image is invalid"}); + assert(e.code() == cudaErrorInvalidKernelImage); + } + + try { + test_throw(); + } catch (::std::runtime_error const &e) { + assert(e.what() == + std::string{ + "cudaErrorInvalidKernelImage: device kernel image is invalid"}); + } +#endif // _LIBCUDACXX_NO_EXCEPTIONS +#endif // __CUDA_ARCH__ + + return 0; +} diff --git a/.upstream-tests/test/cuda/exception_helper.h b/.upstream-tests/test/cuda/exception_helper.h new file mode 100644 index 0000000000..6672fc663f --- /dev/null +++ b/.upstream-tests/test/cuda/exception_helper.h @@ -0,0 +1,9 @@ +#pragma once + +#ifndef _LIBCUDACXX_NO_EXCEPTIONS +#define _LIBCUDACXX_TEST_TRY try +#define _LIBCUDACXX_TEST_CATCH(...) catch(__VA_ARGS__) +#else +#define _LIBCUDACXX_TEST_TRY if (1) +#define _LIBCUDACXX_TEST_CATCH(...) else if (0) +#endif // _LIBCUDACXX_NO_EXCEPTIONS diff --git a/.upstream-tests/test/cuda/memory_resource/memres_derived.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memres_derived.pass.cpp new file mode 100644 index 0000000000..88ed9f7b43 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memres_derived.pass.cpp @@ -0,0 +1,100 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include + +struct event { + enum action { ALLOCATE, DEALLOCATE }; + action act; + std::uintptr_t pointer; + cuda::std::size_t bytes; + cuda::std::size_t alignment; +}; + +bool operator==(event const& lhs, event const& rhs){ + return std::tie(lhs.act, lhs.pointer, lhs.bytes, lhs.alignment) == + std::tie(rhs.act, rhs.pointer, rhs.bytes, rhs.alignment); +} + +template +class derived_resource : public cuda::pmr::memory_resource { +public: + std::vector &events() { return events_; } +private: + void *do_allocate(cuda::std::size_t bytes, + cuda::std::size_t alignment) override { + auto p = 0xDEADBEEF; + events().push_back(event{event::ALLOCATE, p, bytes, alignment}); + return reinterpret_cast(p); + } + + void do_deallocate(void *p, cuda::std::size_t bytes, + cuda::std::size_t alignment) override { + events().push_back(event{event::DEALLOCATE, + reinterpret_cast(p), bytes, + alignment}); + } + + std::vector events_; +}; + +template +void test_derived_resource(){ + using derived = derived_resource; + using base = cuda::pmr::memory_resource; + + derived d; + base * b = &d; + + assert(b->is_equal(*b)); + assert(b->is_equal(d)); + + auto p0 = b->allocate(100); + assert(d.events().size() == 1); + assert((d.events().back() == event{event::ALLOCATE, + reinterpret_cast(p0), + 100, derived::default_alignment})); + + auto p1 = b->allocate(42, 32); + assert(d.events().size() == 2); + assert( + (d.events().back() == + event{event::ALLOCATE, reinterpret_cast(p1), 42, 32})); + + b->deallocate(p0, 100); + assert(d.events().size() == 3); + assert((d.events().back() == event{event::DEALLOCATE, + reinterpret_cast(p0), + 100, derived::default_alignment})); + + b->deallocate(p1, 42, 32); + assert(d.events().size() == 4); + assert((d.events().back() == event{event::DEALLOCATE, + reinterpret_cast(p1), 42, + 32})); +} + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ + test_derived_resource(); + test_derived_resource(); + test_derived_resource(); + test_derived_resource(); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memres_members.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memres_members.pass.cpp new file mode 100644 index 0000000000..1ae55812e7 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memres_members.pass.cpp @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include + +template +constexpr bool test_memory_kind() { + using mr = cuda::pmr::memory_resource; + return std::is_same::value; +} + +template +constexpr bool test_alignment() { + using mr = cuda::pmr::memory_resource; + return mr::default_alignment == Alignment; +} + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ + namespace memory_kind = cuda::pmr::memory_kind; + static_assert(test_memory_kind(), ""); + static_assert(test_memory_kind(), ""); + static_assert(test_memory_kind(), ""); + static_assert(test_memory_kind(), ""); + + static_assert(test_alignment(), ""); + static_assert(test_alignment(), ""); + static_assert(test_alignment(), ""); + static_assert(test_alignment(), ""); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/pmr_adaptor/pmr_adaptor.pass.cpp b/.upstream-tests/test/cuda/memory_resource/pmr_adaptor/pmr_adaptor.pass.cpp new file mode 100644 index 0000000000..f2e2cbb83f --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/pmr_adaptor/pmr_adaptor.pass.cpp @@ -0,0 +1,114 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++98, c++03, c++11, c++14 + +#include +#include +#include +#include +#include +#include +#include +#include + + +#if __has_include() +#include +namespace pmr = ::std::pmr; +#elif __has_include() +#include +namespace pmr = ::std::experimental::pmr; +#endif + + +struct event { + enum action { ALLOCATE, DEALLOCATE }; + action act; + std::uintptr_t pointer; + cuda::std::size_t bytes; + cuda::std::size_t alignment; +}; + +bool operator==(event const& lhs, event const& rhs){ + return std::tie(lhs.act, lhs.pointer, lhs.bytes, lhs.alignment) == + std::tie(rhs.act, rhs.pointer, rhs.bytes, rhs.alignment); +} + +class derived_resource : public cuda::pmr::memory_resource { +public: + std::vector &events() { return events_; } +private: + void *do_allocate(cuda::std::size_t bytes, + cuda::std::size_t alignment) override { + auto p = 0xDEADBEEF; + events().push_back(event{event::ALLOCATE, p, bytes, alignment}); + return reinterpret_cast(p); + } + + void do_deallocate(void *p, cuda::std::size_t bytes, + cuda::std::size_t alignment) override { + events().push_back(event{event::DEALLOCATE, + reinterpret_cast(p), bytes, + alignment}); + } + + std::vector events_; +}; + + +template +void test_adaptor(Pointer mr){ + auto p = &*mr; + cuda::pmr_adaptor adapted{std::move(mr)}; + assert(p == adapted.resource()); + assert(adapted.is_equal(adapted)); + + pmr::memory_resource * std_mr = &adapted; + assert(std_mr->is_equal(adapted)); + assert(adapted.is_equal(*std_mr)); + + auto p0 = std_mr->allocate(42); + assert(p->events().size() == 1); + assert((p->events().back() == event{event::ALLOCATE, + reinterpret_cast(p0), 42, + alignof(std::max_align_t)})); + + std_mr->deallocate(p0, 42); + assert(p->events().size() == 2); + assert((p->events().back() == event{event::DEALLOCATE, + reinterpret_cast(p0), 42, + alignof(std::max_align_t)})); + + auto p1 = std_mr->allocate(42, 16); + assert(p->events().size() == 3); + assert( + (p->events().back() == + event{event::ALLOCATE, reinterpret_cast(p1), 42, 16})); + + std_mr->deallocate(p1, 42, 16); + assert(p->events().size() == 4); + assert( + (p->events().back() == + event{event::DEALLOCATE, reinterpret_cast(p1), 42, 16})); +} + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ +#if defined(_LIBCUDACXX_STD_PMR_NS) + derived_resource mr_raw; + test_adaptor(&mr_raw); + test_adaptor(std::make_unique()); + test_adaptor(std::make_shared()); +#endif +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/pmr_adaptor/pmr_adaptor_equal.pass.cpp b/.upstream-tests/test/cuda/memory_resource/pmr_adaptor/pmr_adaptor_equal.pass.cpp new file mode 100644 index 0000000000..84d3f80ce9 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/pmr_adaptor/pmr_adaptor_equal.pass.cpp @@ -0,0 +1,120 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++98, c++03, c++11, c++14 + +#include +#include +#include +#include +#include +#include +#include +#include + +#if __has_include() +#include +namespace pmr = ::std::pmr; +#elif __has_include() +#include +namespace pmr = ::std::experimental::pmr; +#endif + +template +class derived_resource : public cuda::pmr::memory_resource { +public: +private: + void *do_allocate(cuda::std::size_t, cuda::std::size_t) override { + return nullptr; + } + + void do_deallocate(void *, cuda::std::size_t, cuda::std::size_t) override {} + + bool + do_is_equal(cuda::pmr::memory_resource const &other) const + noexcept override { + return dynamic_cast(&other) != nullptr; + } +}; + +template +class more_derived : public derived_resource { +public: +private: + void *do_allocate(cuda::std::size_t, cuda::std::size_t) override { + return nullptr; + } + void do_deallocate(void *, cuda::std::size_t, cuda::std::size_t) override {} +}; + +template +void assert_equal(T1 const &lhs, T2 const &rhs) { + assert(lhs.is_equal(lhs)); + assert(rhs.is_equal(rhs)); + assert(lhs.is_equal(rhs)); + assert(rhs.is_equal(lhs)); +} + +template +void test_equal(cuda::pmr_adaptor const &lhs, + cuda::pmr_adaptor const &rhs) { + assert_equal(lhs, rhs); + + pmr::memory_resource const *pmr_lhs{&lhs}; + assert_equal(lhs, *pmr_lhs); + assert_equal(rhs, *pmr_lhs); + + pmr::memory_resource const *pmr_rhs{&rhs}; + assert_equal(lhs, *pmr_rhs); + assert_equal(rhs, *pmr_rhs); + + assert_equal(*pmr_rhs, *pmr_lhs); +} + +template +void test_pmr_adaptor_equality(){ + derived_resource d; + cuda::pmr_adaptor a_raw{&d}; + cuda::pmr_adaptor a_unique{std::make_unique>()}; + cuda::pmr_adaptor a_shared{std::make_shared>()}; + + test_equal(a_raw, a_unique); + test_equal(a_raw, a_shared); + test_equal(a_unique, a_shared); + + more_derived m; + assert(d.is_equal(m)); + assert(m.is_equal(d)); + + cuda::pmr_adaptor m_raw{&m}; + test_equal(a_raw, m_raw); + test_equal(a_unique, m_raw); + test_equal(a_shared, m_raw); + + cuda::pmr_adaptor m_unique{std::make_unique>()}; + test_equal(a_raw, m_unique); + test_equal(a_unique, m_unique); + test_equal(a_shared, m_unique); + + cuda::pmr_adaptor m_shared{std::make_shared>()}; + test_equal(a_raw, m_shared); + test_equal(a_unique, m_shared); + test_equal(a_shared, m_shared); +} + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ +#if defined(_LIBCUDACXX_STD_PMR_NS) + test_pmr_adaptor_equality(); +#endif +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/adapt_view_by_properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/adapt_view_by_properties.pass.cpp new file mode 100644 index 0000000000..25abbe920f --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/adapt_view_by_properties.pass.cpp @@ -0,0 +1,99 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include + +namespace memory_kind = cuda::pmr::memory_kind; +namespace memory_access = cuda::pmr::memory_access; +using cuda::pmr::memory_resource; +using cuda::pmr::resource_ptr; +using cuda::pmr::property_list; + +struct event { + enum action { ALLOCATE, DEALLOCATE }; + action act; + std::uintptr_t pointer; + cuda::std::size_t bytes; + cuda::std::size_t alignment; +}; + +bool operator==(event const& lhs, event const& rhs){ + return std::tie(lhs.act, lhs.pointer, lhs.bytes, lhs.alignment) == + std::tie(rhs.act, rhs.pointer, rhs.bytes, rhs.alignment); +} + +template +class dummy_resource : public cuda::pmr::memory_resource { +public: + std::vector &events() { return events_; } +private: + void *do_allocate(cuda::std::size_t bytes, + cuda::std::size_t alignment) override { + auto p = 0xDEADBEEF; + events().push_back(event{event::ALLOCATE, p, bytes, alignment}); + return reinterpret_cast(p); + } + + void do_deallocate(void *p, cuda::std::size_t bytes, + cuda::std::size_t alignment) override { + events().push_back(event{event::DEALLOCATE, + reinterpret_cast(p), bytes, + alignment}); + } + + std::vector events_; +}; + +template +struct resource_adaptor : public memory_resource> { + public: + resource_adaptor(resource_ptr upstream) : upstream(upstream) { + } + + private: + void *do_allocate(size_t size, size_t alignment) override { + return upstream->allocate(size, alignment); + } + + void do_deallocate(void *ptr, size_t size, size_t alignment) override { + upstream->deallocate(ptr, size, alignment); + } + + resource_ptr upstream; +}; + +template +resource_adaptor adapt(cuda::pmr::basic_resource_ptr v) { + return resource_adaptor(v); +} + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + + dummy_resource res; + resource_ptr> v1 = &res; + resource_ptr> v2 = &res; + resource_ptr v3 = &res; + cuda::pmr::resource_ptr> v4 = &res; + + (void)adapt(v1); + (void)adapt(v2); + (void)adapt(v3); + (void)adapt(v4); + (void)adapt(view_resource(&res)); + +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_convert_compatible.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_convert_compatible.pass.cpp new file mode 100644 index 0000000000..1ddbf937dc --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_convert_compatible.pass.cpp @@ -0,0 +1,78 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include +#include "resource_hierarchy.h" + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + { + view_D2 derived2_view; + derived2_view = derived2_view; + + view_D1 derived1_view = derived2_view; + derived1_view = derived2_view; + derived1_view = derived1_view; + + + view_M memres_view = derived2_view; + memres_view = derived2_view; + memres_view = derived1_view; + memres_view = memres_view; + + view_B base_view = derived2_view; + + base_view = derived2_view; + base_view = derived1_view; + base_view = memres_view; + base_view = base_view; + } + + { + + view_DA2 derived_async2_view; + derived_async2_view = derived_async2_view; + + view_DA1 derived_async1_view = derived_async2_view; + derived_async1_view = derived_async2_view; + derived_async1_view = derived_async1_view; + + view_MA memres_async_view = derived_async2_view; + memres_async_view = derived_async2_view; + memres_async_view = derived_async1_view; + memres_async_view = memres_async_view; + + view_M memres_view = derived_async2_view; + memres_view = derived_async2_view; + memres_view = derived_async1_view; + memres_view = memres_async_view; + memres_view = memres_view; + + view_BA base_async_view = derived_async2_view; + base_async_view = derived_async2_view; + base_async_view = derived_async1_view; + base_async_view = memres_async_view; + + view_B base_view = derived_async2_view; + base_view = derived_async2_view; + base_view = derived_async1_view; + base_view = memres_async_view; + base_view = memres_view; + base_view = base_async_view; + } +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_interface_propagation.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_interface_propagation.pass.cpp new file mode 100644 index 0000000000..b63a8c4a07 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_interface_propagation.pass.cpp @@ -0,0 +1,197 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + + +#include +#include +#include +#include +#include +#include +#include +#include + +class sync_resource : public cuda::pmr::memory_resource { +public: + int extra_sync() { + return 42; + } + + size_t allocated_size = 0, allocated_alignment = 0; + void *allocated_pointer = nullptr; + mutable const cuda::pmr::memory_resource *compared_resource = nullptr; +private: + void *do_allocate(size_t size, size_t alignment) override { + allocated_size = size; + allocated_alignment = alignment; + return allocated_pointer = reinterpret_cast(0x123400); + } + + void do_deallocate(void *mem, size_t size, size_t alignment) { + assert(mem == allocated_pointer); + assert(size == allocated_size); + assert(alignment == allocated_alignment); + allocated_pointer = 0; + allocated_size = 0; + allocated_alignment = 0; + } + + bool do_is_equal(const cuda::pmr::memory_resource &other) const noexcept override { + compared_resource = &other; + return this == &other; + } +}; + +class async_resource : public cuda::pmr::stream_ordered_resource { +public: + int extra_async() { + return 42; + } + + size_t allocated_size = 0, allocated_alignment = 0; + void *allocated_pointer = nullptr; + cuda::stream_view allocation_stream = {}; + mutable const cuda::pmr::memory_resource *compared_resource = nullptr; +private: + void *do_allocate(size_t size, size_t alignment) override { + allocated_size = size; + allocated_alignment = alignment; + return allocated_pointer = reinterpret_cast(0x123400); + } + + void *do_allocate_async(size_t size, size_t alignment, cuda::stream_view stream) override { + allocated_size = size; + allocated_alignment = alignment; + allocation_stream = stream; + return allocated_pointer = reinterpret_cast(0x123400); + } + + void do_deallocate(void *mem, size_t size, size_t alignment) { + assert(mem == allocated_pointer); + assert(size == allocated_size); + assert(alignment == allocated_alignment); + allocated_pointer = 0; + allocated_size = 0; + allocated_alignment = 0; + } + + void do_deallocate_async(void *mem, size_t size, size_t alignment, cuda::stream_view stream) override { + assert(mem == allocated_pointer); + assert(size == allocated_size); + assert(stream == allocation_stream); + assert(alignment == allocated_alignment); + assert(stream == allocation_stream); + allocated_pointer = 0; + allocated_size = 0; + allocated_alignment = 0; + allocation_stream = {}; + } + + bool do_is_equal(const cuda::pmr::memory_resource &other) const noexcept override { + compared_resource = &other; + return this == &other; + } +}; + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + // syncrhonous resource + { + sync_resource rsrc; + auto view = cuda::pmr::view_resource(&rsrc); + assert(view->extra_sync() == 42); + void *ptr = view->allocate(23, 32); + assert(ptr == view->allocated_pointer); + assert(23 == view->allocated_size); + assert(32 == view->allocated_alignment); + view->deallocate(ptr, 23, 32); + assert(nullptr == view->allocated_pointer); + assert(0 == view->allocated_size); + assert(0 == view->allocated_alignment); + + assert(view->is_equal(rsrc)); + assert(view->compared_resource == &rsrc); + } + { + sync_resource rsrc; + cuda::pmr::resource_ptr view = &rsrc; + void *ptr = view->allocate(23, 32); + assert(ptr == rsrc.allocated_pointer); + assert(23 == rsrc.allocated_size); + assert(32 == rsrc.allocated_alignment); + view->deallocate(ptr, 23, 32); + assert(nullptr == rsrc.allocated_pointer); + assert(0 == rsrc.allocated_size); + assert(0 == rsrc.allocated_alignment); + } + // stream-ordered resource + { + cuda::stream_view stream((cudaStream_t)0x1234); + async_resource rsrc; + auto view = cuda::pmr::view_resource(&rsrc); + + assert(view->extra_async() == 42); + + void *ptr = view->allocate(23, 32); + assert(ptr == view->allocated_pointer); + assert(23 == view->allocated_size); + assert(32 == view->allocated_alignment); + + view->deallocate(ptr, 23, 32); + assert(nullptr == view->allocated_pointer); + assert(0 == view->allocated_size); + assert(0 == view->allocated_alignment); + assert(0 == view->allocated_alignment); + + ptr = view->allocate_async(42, 64, stream); + assert(ptr == view->allocated_pointer); + assert(42 == view->allocated_size); + assert(64 == view->allocated_alignment); + assert(stream.get() == view->allocation_stream.get()); + + view->deallocate_async(ptr, 42, 64, stream); + assert(nullptr == view->allocated_pointer); + assert(0 == view->allocated_size); + assert(0 == view->allocated_alignment); + assert(0 == view->allocation_stream.get()); + + assert(view->is_equal(rsrc)); + assert(view->compared_resource == &rsrc); + } + { + cuda::stream_view stream((cudaStream_t)0x1234); + async_resource rsrc; + cuda::pmr::stream_ordered_resource_ptr view = &rsrc; + + void *ptr = view->allocate(23, 32); + assert(ptr == rsrc.allocated_pointer); + assert(23 == rsrc.allocated_size); + assert(32 == rsrc.allocated_alignment); + + view->deallocate(ptr, 23, 32); + assert(nullptr == rsrc.allocated_pointer); + assert(0 == rsrc.allocated_size); + assert(0 == rsrc.allocated_alignment); + assert(0 == rsrc.allocated_alignment); + + ptr = view->allocate_async(42, 64, stream); + assert(ptr == rsrc.allocated_pointer); + assert(42 == rsrc.allocated_size); + assert(64 == rsrc.allocated_alignment); + assert(stream.get() == rsrc.allocation_stream.get()); + + view->deallocate_async(ptr, 42, 64, stream); + assert(nullptr == rsrc.allocated_pointer); + assert(0 == rsrc.allocated_size); + assert(0 == rsrc.allocated_alignment); + assert(0 == rsrc.allocation_stream.get()); + } +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_pointer_incompatible.fail.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_pointer_incompatible.fail.cpp new file mode 100644 index 0000000000..221b7ab94c --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/basic_resource_ptr_pointer_incompatible.fail.cpp @@ -0,0 +1,28 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include +#include "resource_hierarchy.h" + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + { + view_D2 da2; + view_MA ma = da2; // cant't assign syncrhounous to asyncrhounous resource + } +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_hierarchy.h b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_hierarchy.h new file mode 100644 index 0000000000..9ef476188f --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_hierarchy.h @@ -0,0 +1,83 @@ +#pragma once + +#include + +class derived1 : public cuda::pmr::memory_resource { +}; + +class derived2 : public derived1 { +}; + + +class derived_async1 : public cuda::pmr::stream_ordered_resource { +}; + +class derived_async2 : public derived_async1 { +}; + +/* + D2 - derived2, host-accessible, oversubscribable, host-located + D1 - derived1, host-accessible, host-located + DA2 - derived_async2, host-accessible, oversubscribable, host-located + DA1 - derived_async1, oversubscribable, host-located + M - memory_resource, host-accessible, host-located + MA - stream_ordered_resource, host-accessible, host-located + B - memory_resource_base, host-located + BA - stream_ordered_resource_base, host-located + +The view compatibility is as follows: + +from---> DA2 DA1 D2 D1 MA M BA B +to--v + DA2 X + DA1 X X + D2 X + D1 X X + MA X X X + M X X X X X X + BA X X X X + B X X X X X X X X +*/ + +// NOTE - some views are parameterized with separate properties and some with property_list<...> + +using view_D2 = cuda::pmr::basic_resource_ptr< + derived2*, + cuda::pmr::memory_access::host, + cuda::pmr::oversubscribable, + cuda::pmr::memory_location::host>; + +using view_D1 = cuda::pmr::basic_resource_ptr< + derived1*, + cuda::pmr::memory_access::host, + cuda::pmr::memory_location::host>; + +using view_DA2 = cuda::pmr::basic_resource_ptr< + derived_async2*, + cuda::pmr::property_list< + cuda::pmr::memory_access::host, + cuda::pmr::oversubscribable, + cuda::pmr::memory_location::host>>; + +using view_DA1 = cuda::pmr::basic_resource_ptr< + derived_async1*, + cuda::pmr::property_list< + cuda::pmr::memory_access::host, + cuda::pmr::memory_location::host>>; + +using view_M = cuda::pmr::basic_resource_ptr< + cuda::pmr::memory_resource*, + cuda::pmr::property_list< + cuda::pmr::memory_access::host, + cuda::pmr::memory_location::host>>; + +using view_MA = cuda::pmr::basic_resource_ptr< + cuda::pmr::stream_ordered_resource*, + cuda::pmr::memory_access::host, + cuda::pmr::memory_location::host>; + +using view_B = cuda::pmr::resource_ptr< + cuda::pmr::memory_location::host>; + +using view_BA = cuda::pmr::stream_ordered_resource_ptr< + cuda::pmr::memory_location::host>; diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_pointer_compatibility.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_pointer_compatibility.pass.cpp new file mode 100644 index 0000000000..ea2ce122d3 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_pointer_compatibility.pass.cpp @@ -0,0 +1,63 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include +#include "resource_hierarchy.h" + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + static_assert(cuda::pmr::detail::is_resource_pointer_convertible::value, + "A pointer to a derived class should be convertible to a pointer of public a base class"); + + static_assert(cuda::pmr::detail::is_resource_pointer_convertible*>::value, + "A pointer to a derived class should be convertible to a pointer of public a base class"); + + static_assert(cuda::pmr::detail::is_resource_pointer_convertible::value, + "A pointer to a derived memory resource should be convertible to a pointer to the common base."); + + static_assert(cuda::pmr::detail::is_resource_pointer_convertible::value, + "A pointer to a derived class should be convertible to a pointer of public a base class"); + + static_assert(cuda::pmr::detail::is_resource_pointer_convertible*>::value, + "A pointer to a derived class should be convertible to a pointer of public a base class"); + + static_assert(cuda::pmr::detail::is_resource_pointer_convertible::value, + "A pointer to a derived memory resource should be convertible to a pointer to the common base."); + + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible::value, + "Pointers to unrelated classes should not be convertible"); + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible::value, + "Pointers to unrelated classes should not be convertible"); + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible::value, + "Conversion to a subclass pointer should not be possible"); + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible*>::value, + "A pointer to a synchronous resource should not be convertible to a stream ordered one"); + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible::value, + "A pointer to a synchronous resource should not be convertible to a stream ordered one"); + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible*>::value, + "A pointer to a commnon base class should not be convertible to a pointer to a kind-qualified resource."); + + static_assert(!cuda::pmr::detail::is_resource_pointer_convertible*>::value, + "A pointer to a commnon base class should not be convertible to a pointer to a kind-qualified resource."); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_comparison.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_comparison.pass.cpp new file mode 100644 index 0000000000..cef39ae6c1 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_comparison.pass.cpp @@ -0,0 +1,101 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + + +#include +#include +#include +#include +#include +#include +#include + + +template +class resource : public cuda::pmr::stream_ordered_resource { +public: + int value = 0; +private: + void *do_allocate(size_t, size_t) override { + return nullptr; + } + + void *do_allocate_async(size_t, size_t, cuda::stream_view) override { + return nullptr; + } + + void do_deallocate(void *, size_t, size_t) { + } + + void do_deallocate_async(void *, size_t, size_t, cuda::stream_view) override { + } + +#ifdef _LIBCUDACXX_EXT_RTTI_ENABLED + bool do_is_equal(const cuda::pmr::memory_resource &other) const noexcept override { + fprintf(stderr, "Comparison start: %p %p\n", this, &other); + if (auto *other_ptr = dynamic_cast(&other)) { + fprintf(stderr, "values: %d %d\n", value, other_ptr->value); + return value == other_ptr->value; + } else { + return false; + } + } +#endif +}; + +struct tag1; +struct tag2; + +int main(int argc, char **argv) { +#if !defined(__CUDA_ARCH__) && defined(_LIBCUDACXX_EXT_RTTI_ENABLED) + resource r1, r2, r3; + resource r4; + cuda::pmr::basic_resource_ptr*, cuda::pmr::kind_to_properties> v1_null; + cuda::pmr::resource_ptr v2_null; + assert(v1_null == v2_null); + r1.value = 42; + r2.value = 42; + r3.value = 99; + r4.value = 42; + using t1 = decltype(view_resource(&r1)); + using t2 = decltype(view_resource(&r2)); + using t4 = decltype(view_resource(&r4)); + assert(view_resource(&r1) == view_resource(&r2)); + assert(view_resource(&r1) != view_resource(&r3)); + assert(view_resource(&r4) == view_resource(&r4)); + cuda::pmr::resource_ptr v1 = &r1; + cuda::pmr::stream_ordered_resource_ptr v2 = &r2; + cuda::pmr::resource_ptr v3 = &r3; + cuda::pmr::resource_ptr v4 = &r4; + // compare views + assert(v1 == v2); + assert(v1 != v3); + assert(v1 != v4); + // assert(v2 != v3); - cannot compare - incompatible views + assert(v2 != v4); + assert(v3 != v4); + assert(v4 == v4); + + // compare views vs resources + assert(v1 == &r1); + assert(v1 == &r2); + assert(v1 != &r4); + assert(v2 == &r2); + assert(v2 != &r4); + assert(v4 == &r4); + + assert(&r1 == v1); + assert(&r2 == v1); + assert(&r4 != v1); + assert(&r2 == v2); + assert(&r4 != v2); + assert(&r4 == v4); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_compatibility.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_compatibility.pass.cpp new file mode 100644 index 0000000000..a474b3bf03 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_compatibility.pass.cpp @@ -0,0 +1,117 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include +#include "resource_hierarchy.h" + +int main(int argc, char **argv) { + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a self."); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to self."); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to self."); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to self."); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to self."); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of a superclass with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to self."); + + static_assert(cuda::pmr::is_view_convertible::value, + "A resource view should be convertible to a view of common base with a subset of properties"); + + + static_assert(!cuda::pmr::is_view_convertible::value, + "Views to unrelated types should not be compatible."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "Views to unrelated types should not be compatible."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view should not to a view with a superclass pointer."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view should not to a view with a superclass pointer."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view should not to a view with a superclass pointer."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view should not to a view with a superclass pointer."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view should not to a view with a superclass pointer."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view to a syncrhouns resource cannot be converted to a stream ordered one."); + + static_assert(!cuda::pmr::is_view_convertible::value, + "A resource view to a syncrhouns resource cannot be converted to a stream ordered one."); + + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind.pass.cpp new file mode 100644 index 0000000000..44645cc2bc --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind.pass.cpp @@ -0,0 +1,36 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + cuda::pmr::resource_ptr props_only; + + cuda::pmr::resource_ptr> kind_only; + + cuda::pmr::resource_ptr, + cuda::pmr::memory_access::host, + cuda::pmr::oversubscribable, + cuda::pmr::memory_location::host> props_and_kind; + + props_only = kind_only; // the source properties should be propagated form of_kind + props_and_kind = kind_only; // the additional properties should be propagated form of_kind +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind_mismatch.fail.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind_mismatch.fail.cpp new file mode 100644 index 0000000000..685729ecb0 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind_mismatch.fail.cpp @@ -0,0 +1,29 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + cuda::pmr::resource_ptr managed; + cuda::pmr::resource_ptr> host; + + // Despite managed having a superset of properties of host, + // this should fail because the kind is a property in itself. + host = managed; +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind_reverse.fail.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind_reverse.fail.cpp new file mode 100644 index 0000000000..db3537f88e --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_of_kind_reverse.fail.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + cuda::pmr::resource_ptr props_only; + cuda::pmr::resource_ptr> kind_only; + kind_only = props_only; // no conversion from a list of properties back to memory kind +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_prop_subset.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_prop_subset.pass.cpp new file mode 100644 index 0000000000..567599de77 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_prop_subset.pass.cpp @@ -0,0 +1,40 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + cuda::pmr::resource_ptr props_only; + + // test copy constructor + cuda::pmr::resource_ptr same_props = props_only; + + // test converting constructor + cuda::pmr::resource_ptr subset = props_only; + + // test asisgnment + same_props = props_only; + subset = props_only; +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_unsupported_prop.fail.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/resource_ptr_unsupported_prop.fail.cpp new file mode 100644 index 0000000000..e69de29bb2 diff --git a/.upstream-tests/test/cuda/memory_resource/resource_ptr/view_resource.pass.cpp b/.upstream-tests/test/cuda/memory_resource/resource_ptr/view_resource.pass.cpp new file mode 100644 index 0000000000..6e811232d7 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/resource_ptr/view_resource.pass.cpp @@ -0,0 +1,54 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include +#include "resource_hierarchy.h" + +int main(int argc, char **argv) { +#ifndef __CUDA_ARCH__ + derived_async2 *res = nullptr; + using expected_view_type = cuda::pmr::basic_resource_ptr>>; + auto view = cuda::pmr::view_resource>(res); + static_assert(std::is_base_of::value, + "Unexpected return type of `view_resource` - expected a basic_resource_ptr with is_kind and concrete pointer type"); + view_DA2 vda2(res); + vda2 = view; + view_DA1 vda1(res); + vda1 = view; + vda1 = vda2; + view_MA vma(res); + vma = view; + vma = vda2; + vma = vda1; + view_BA vba(res); + vba = view; + vba = vda1; + vba = vma; + vba = vba; + view_M vm(res); + vm = view; + vm = vda1; + vm = vma; + vm = vm; + view_B vb(res); + vb = view; + vb = vda1; + vb = vma; + vb = vba; + vb = vm; +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/stream_ordered/stream_memres_derived.pass.cpp b/.upstream-tests/test/cuda/memory_resource/stream_ordered/stream_memres_derived.pass.cpp new file mode 100644 index 0000000000..093695d588 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/stream_ordered/stream_memres_derived.pass.cpp @@ -0,0 +1,130 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include +#include +#include + +struct event { + enum action { ALLOCATE, DEALLOCATE }; + action act; + std::uintptr_t pointer; + cuda::std::size_t bytes; + cuda::std::size_t alignment; + cuda::stream_view stream; +}; + +bool operator==(event const &lhs, event const &rhs) { + return std::tie(lhs.act, lhs.pointer, lhs.bytes, lhs.alignment, lhs.stream) == + std::tie(rhs.act, rhs.pointer, rhs.bytes, rhs.alignment, rhs.stream); +} + +template +class derived_resource : public cuda::pmr::stream_ordered_resource { +public: + std::vector &events() { return events_; } + +private: + void *do_allocate_async(cuda::std::size_t bytes, cuda::std::size_t alignment, + cuda::stream_view stream) override { + auto p = 0xDEADBEEF; + events().push_back(event{event::ALLOCATE, p, bytes, alignment, stream}); + return reinterpret_cast(p); + } + + void do_deallocate_async(void *p, cuda::std::size_t bytes, + cuda::std::size_t alignment, + cuda::stream_view stream) override { + events().push_back(event{event::DEALLOCATE, + reinterpret_cast(p), bytes, + alignment, stream}); + } + + std::vector events_; +}; + +template void test_derived_resource() { + using derived = derived_resource; + using base = cuda::pmr::stream_ordered_resource; + + derived d; + base *b = &d; + + assert(b->is_equal(*b)); + assert(b->is_equal(d)); + + cuda::stream_view default_stream; + + auto p0 = b->allocate(100); + assert(d.events().size() == 1); + assert((d.events().back() == event{event::ALLOCATE, + reinterpret_cast(p0), 100, + derived::default_alignment, default_stream})); + + auto p1 = b->allocate(42, 32); + assert(d.events().size() == 2); + assert((d.events().back() == event{event::ALLOCATE, + reinterpret_cast(p1), 42, + 32, default_stream})); + + b->deallocate(p0, 100); + assert(d.events().size() == 3); + assert((d.events().back() == event{event::DEALLOCATE, + reinterpret_cast(p0), 100, + derived::default_alignment, default_stream})); + + b->deallocate(p1, 42, 32); + assert(d.events().size() == 4); + assert((d.events().back() == event{event::DEALLOCATE, + reinterpret_cast(p1), 42, + 32, default_stream})); + + cuda::stream_view s = reinterpret_cast(13); + + auto p2 = b->allocate_async(123, s); + assert(d.events().size() == 5); + assert((d.events().back() == event{event::ALLOCATE, + reinterpret_cast(p2), 123, + derived::default_alignment, s})); + + auto p3 = b->allocate_async(42, 64, s); + assert(d.events().size() == 6); + assert((d.events().back() == event{event::ALLOCATE, + reinterpret_cast(p3), 42, + 64, s})); + + b->deallocate_async(p2, 123, s); + assert(d.events().size() == 7); + assert((d.events().back() == event{event::DEALLOCATE, + reinterpret_cast(p2), 123, + derived::default_alignment, s})); + + b->deallocate_async(p3, 42, 64, s); + assert(d.events().size() == 8); + assert((d.events().back() == event{event::DEALLOCATE, + reinterpret_cast(p3), 42, + 64, s})); +} + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ + test_derived_resource(); + test_derived_resource(); + test_derived_resource(); + test_derived_resource(); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/stream_ordered/stream_memres_members.pass.cpp b/.upstream-tests/test/cuda/memory_resource/stream_ordered/stream_memres_members.pass.cpp new file mode 100644 index 0000000000..350c99f93a --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/stream_ordered/stream_memres_members.pass.cpp @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include + +template constexpr bool test_memory_kind() { + using mr = cuda::pmr::stream_ordered_resource; + return std::is_same::value; +} + +template +constexpr bool test_alignment() { + using mr = cuda::pmr::stream_ordered_resource; + return mr::default_alignment == Alignment; +} + +int main(int argc, char **argv) { + +#ifndef __CUDA_ARCH__ + namespace memory_kind = cuda::pmr::memory_kind; + static_assert(test_memory_kind(), ""); + static_assert(test_memory_kind(), ""); + static_assert(test_memory_kind(), ""); + static_assert(test_memory_kind(), ""); + + using mr = cuda::pmr::stream_ordered_resource; + + static_assert(test_alignment(), ""); + static_assert(test_alignment(), ""); + static_assert(test_alignment(), ""); + static_assert(test_alignment(), ""); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_convert_int.fail.cpp b/.upstream-tests/test/cuda/stream_view/stream_convert_int.fail.cpp new file mode 100644 index 0000000000..7b4e9a26c7 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_convert_int.fail.cpp @@ -0,0 +1,22 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +void foo(cuda::stream_view){} + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + foo(0); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_convert_nullptr.fail.cpp b/.upstream-tests/test/cuda/stream_view/stream_convert_nullptr.fail.cpp new file mode 100644 index 0000000000..5ad1519c85 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_convert_nullptr.fail.cpp @@ -0,0 +1,22 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +void foo(cuda::stream_view){} + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + foo(nullptr); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_convert_stream_handle.pass.cpp b/.upstream-tests/test/cuda/stream_view/stream_convert_stream_handle.pass.cpp new file mode 100644 index 0000000000..abe984adb1 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_convert_stream_handle.pass.cpp @@ -0,0 +1,23 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +void foo(cuda::stream_view){} + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + cudaStream_t s{}; + foo(s); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_default.pass.cpp b/.upstream-tests/test/cuda/stream_view/stream_default.pass.cpp new file mode 100644 index 0000000000..846d7c4f9f --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_default.pass.cpp @@ -0,0 +1,24 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + static_assert( + std::is_same::value, ""); + cuda::stream_view s; + assert(s.get() == cudaStream_t{0}); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_equality.pass.cpp b/.upstream-tests/test/cuda/stream_view/stream_equality.pass.cpp new file mode 100644 index 0000000000..52d861796c --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_equality.pass.cpp @@ -0,0 +1,26 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + cudaStream_t s = reinterpret_cast(42); + cuda::stream_view sv0{s}; + cuda::stream_view sv1{s}; + cuda::stream_view sv2{}; + assert(sv0 == sv0); + assert(sv0 == sv1); + assert(sv0 != sv2); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_get.pass.cpp b/.upstream-tests/test/cuda/stream_view/stream_get.pass.cpp new file mode 100644 index 0000000000..4a786be438 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_get.pass.cpp @@ -0,0 +1,22 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + cudaStream_t s = reinterpret_cast(42); + cuda::stream_view sv{s}; + assert(sv.get() == s); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_ready.pass.cpp b/.upstream-tests/test/cuda/stream_view/stream_ready.pass.cpp new file mode 100644 index 0000000000..3aec18a885 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_ready.pass.cpp @@ -0,0 +1,29 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include "../exception_helper.h" + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + cudaStream_t s; + cudaStreamCreate(&s); + cuda::stream_view sv{s}; + _LIBCUDACXX_TEST_TRY { + assert(sv.ready()); + } _LIBCUDACXX_TEST_CATCH(...) { + assert(false && "Should not have thrown"); + } + cudaStreamDestroy(s); +#endif + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_view/stream_wait.pass.cpp b/.upstream-tests/test/cuda/stream_view/stream_wait.pass.cpp new file mode 100644 index 0000000000..b9a5467f2b --- /dev/null +++ b/.upstream-tests/test/cuda/stream_view/stream_wait.pass.cpp @@ -0,0 +1,29 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include "../exception_helper.h" + +int main(int argc, char** argv){ + +#ifndef __CUDA_ARCH__ + cudaStream_t s; + cudaStreamCreate(&s); + cuda::stream_view sv{s}; + _LIBCUDACXX_TEST_TRY { + sv.wait(); + } _LIBCUDACXX_TEST_CATCH(...) { + assert(false && "Should not have thrown"); + } + cudaStreamDestroy(s); +#endif + + return 0; +} diff --git a/include/cuda/detail/__cuda_util b/include/cuda/detail/__cuda_util new file mode 100644 index 0000000000..d3b2a78477 --- /dev/null +++ b/include/cuda/detail/__cuda_util @@ -0,0 +1,53 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX_CUDA_UTIL +#define _LIBCUDACXX_CUDA_UTIL + +#include "cuda/std/detail/__config" +#include "cuda/std/detail/__pragma_push" + +#include // for string forward decl + +#include + +#ifdef _LIBCUDACXX_NO_EXCEPTIONS +#include +#endif +#include "cuda/std/detail/__pragma_push" + +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA + +namespace detail { +inline _LIBCUDACXX_INLINE_VISIBILITY +void __throw_on_cuda_error(::cudaError_t __error, char const *__msg) { + if (::cudaSuccess != __error) { + ::cudaGetLastError(); // Clear CUDA error state +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + throw cuda::cuda_error{__error, __msg}; +#else + ::std::abort(); +#endif + } +} + +inline _LIBCUDACXX_INLINE_VISIBILITY +void __throw_on_cuda_error(::cudaError_t __error) { + __throw_on_cuda_error(__error, ""); +} + +} // namespace detail + +_LIBCUDACXX_END_NAMESPACE_CUDA + +#include "../std/detail/__pragma_pop" + +#endif // _LIBCUDACXX_CUDA_UTIL diff --git a/include/cuda/exception b/include/cuda/exception new file mode 100644 index 0000000000..fdc9951ae3 --- /dev/null +++ b/include/cuda/exception @@ -0,0 +1,48 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_EXCEPTION +#define _CUDA_EXCEPTION + +#include "std/detail/__config" +#include "std/detail/__pragma_push" +#include "std/version" + +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA + +/** + * \brief Exception thrown when a CUDA runtime error is encountered. + * + */ +struct cuda_error : ::std::runtime_error { + + cuda_error(cudaError_t __error, ::std::string const &__msg) + : ::std::runtime_error{__msg + ": " + cudaGetErrorName(__error) + ": " + + cudaGetErrorString(__error)}, + __error_code_{__error} {} + + explicit cuda_error(cudaError_t __error) + : ::std::runtime_error{::std::string{cudaGetErrorName(__error)} + ": " + + cudaGetErrorString(__error)}, + __error_code_{__error} {} + + cudaError_t code() const noexcept { return __error_code_; } + +private: + cudaError_t __error_code_; +}; + +_LIBCUDACXX_END_NAMESPACE_CUDA + +#include "std/detail/__pragma_pop" + +#endif // _CUDA_EXCEPTION diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource new file mode 100644 index 0000000000..3cac9292ad --- /dev/null +++ b/include/cuda/memory_resource @@ -0,0 +1,1077 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_MEMORY_RESOURCE +#define _CUDA_MEMORY_RESOURCE + +#include +#include "std/cstddef" +#include "std/detail/__config" +#include "std/detail/__pragma_push" +#include "std/type_traits" +#include "std/utility" +#include "std/version" +#include "stream_view" + +#ifdef _LIBCUDACXX_EXT_RTTI_ENABLED +#include +#endif + +#if _LIBCUDACXX_STD_VER > 14 +#if __has_include() +#include +#define _LIBCUDACXX_STD_PMR_NS ::std::pmr +#elif __has_include() +#include +#define _LIBCUDACXX_STD_PMR_NS ::std::experimental::pmr +#endif // __has_include() +#endif // _LIBCUDACXX_STD_VER > 14 + +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA +namespace pmr { + +/*! + * \brief Groups the tag types denoting the kind of memory of an allocation. + * + * Memory allocation kind determines where memory can be accessed and the + * performance characteristics of accesses. + * + * This is not a closed set, the user code can define custom memory kinds. + */ +namespace memory_kind { + /*! + * \brief Ordinary host memory + */ + struct host; + + /*! + * \brief Device memory, as allocated by cudaMalloc. + */ + struct device; + + /*! + * \brief Device-accessible host memory. + */ + struct pinned; + + /*! + * \brief Virtual memory that is automatically migrated between the host and devices. + */ + struct managed; +} // namespace memory_kind + +namespace detail { + +template +struct __type_pack {}; + +namespace __fallback_typeid { + +template +struct _LIBCUDACXX_TEMPLATE_VIS __unique_typeinfo { static constexpr int __id = 0; }; +template constexpr int __unique_typeinfo<_Tp>::__id; + +template +inline _LIBCUDACXX_INLINE_VISIBILITY +constexpr const void* __get_fallback_typeid() { + return &__unique_typeinfo>::__id; +} + +template +const ::std::type_info *__get_typeid() { +#ifdef _LIBCUDACXX_EXT_RTTI_ENABLED + return &typeid(_Tp); +#else + return nullptr; +#endif +} + +inline bool __compare_type(const ::std::type_info *__ti1, const void *__fallback_ti1, + const ::std::type_info *__ti2, const void *__fallback_ti2) { +#ifdef _LIBCUDACXX_EXT_RTTI_ENABLED + if (__ti1 && __ti2 && *__ti1 == *__ti2) + return true; +#endif + return __fallback_ti1 == __fallback_ti2; +} + +template +bool __is_type(const ::std::type_info *__ti1, const void *__fallback_ti1) { + return __compare_type(__ti1, __fallback_ti1, __get_typeid<_Tp>(), __get_fallback_typeid<_Tp>()); +} + +} // namespace __fallback_typeid + +} // namespace detail + +/*! + * \brief Groups the tag types denoting the execution environment in which the memory can be accessed + * + * This is not a closed set, the user code can define custom accessibility. + */ +namespace memory_access { + struct host; + struct device; +} // namespace memory_access + +/*! + * \brief A memory property tag type indicating that the memory can be oversubscribed. + * + * Oversubscribable memory doesn't need to have backing physical storage at all times. + */ +struct oversubscribable; + +/*! + * \brief A memory property tag type indicating that the memory has a backing physical + * storage in the target location at all times. + */ +struct resident; + +/*! + * \brief Groups the tag types that denote the actual location of the physical storage + * + * Memory kinds which can be migrated between locations can define multiple locations. + */ +namespace memory_location { + /*! + * \brief A memory property tag type indicating that the memory is located on a device + */ + struct device; + + /*! + * \brief A memory property tag type indicating that the memory is located in the host memory + */ + struct host; +} // namespace memory_location + +template +class memory_resource; + +template +class basic_resource_ptr; + +namespace detail { + +class memory_resource_base { +public: + static constexpr _CUDA_VSTD::size_t default_alignment = alignof(_CUDA_VSTD::max_align_t); + + /*! + * \brief Allocates storage of size at least `__bytes` bytes. + * + * The returned storage is aligned to the specified `__alignment` if such + * alignment is supported. Otherwise throws. + * + * Storage may be accessed immediately within the execution contexts that + * can access the memory. + * + * \throws If storage of the requested size and alignment cannot be obtained. + * + * \param __bytes The size in bytes of the allocation + * \param __alignment The alignment of the allocation + * \return Pointer to the requested storage + */ + void *allocate(size_t __bytes, size_t __alignment = default_alignment) { + return do_allocate(__bytes, __alignment); + } + + /*! + * \brief Deallocates the storage pointed to by `__p`. + * + * `__p` must have been returned by a prior call to `allocate(__bytes, + * __alignment)` on a `memory_resource` that compares equal to `*this`, and + * the storage it points to must not yet have been deallocated, otherwise + * behavior is undefined. + * + * \throws Nothing. + * + * \param __p Pointer to storage to be deallocated + * \param __bytes The size in bytes of the allocation. This must be equal to + * the value of `__bytes` that was specified to the `allocate` call that + * returned `__p`. + * \param __alignment The alignment of the allocation. This must be equal to + * the value of `__alignment` that was specified to the `allocate` call that + * returned `__p`. + */ + void deallocate(void *__mem, size_t __bytes, size_t __alignment = default_alignment) { + do_deallocate(__mem, __bytes, __alignment); + } + + /*! + * \brief Tries to cast the resource to a resource of given kind + */ + template + memory_resource<_Kind> *as_kind() noexcept { + using __tag = detail::__type_pack<_Kind>; + return static_cast *>( + __do_as_kind(detail::__fallback_typeid::__get_typeid<__tag>(), + detail::__fallback_typeid::__get_fallback_typeid<__tag>())); + } + + /*! + * \brief Tries to cast the resource to a resource of given kind + */ + template + const memory_resource<_Kind> *as_kind() const noexcept { + using __tag = detail::__type_pack<_Kind>; + return static_cast *>( + __do_as_kind(detail::__fallback_typeid::__get_typeid<__tag>(), + detail::__fallback_typeid::__get_fallback_typeid<__tag>())); + } + +protected: + virtual void *do_allocate(size_t __bytes, size_t __alignment) = 0; + virtual void do_deallocate(void *__mem, size_t __bytes, size_t __alignment) = 0; + + virtual bool is_equal_base(const memory_resource_base &other) const noexcept = 0; + + bool is_equal(const memory_resource_base &other) const noexcept { + return is_equal_base(other); + } + + template + friend class cuda::pmr::basic_resource_ptr; + + virtual void *__do_as_kind(const ::std::type_info *__tag_type_id, const void *__tag_type_fallback_id) const noexcept = 0; +}; + +class stream_ordered_resource_base : public virtual memory_resource_base { +public: + /*! + * \brief Allocates storage of size at least `__bytes` bytes in stream order + * on `__stream`. + * + * The returned storage is aligned to `default_alignment`. + * + * The returned storage may be used immediately only on `__stream`. Accessing + * it on any other stream (or the host) requires first synchronizing with + * `__stream`. + * + * \throws If the storage of the requested size and `default_alignment` cannot + * be obtained. + * + * \param __bytes The size in bytes of the allocation. + * \param __stream The stream on which to perform the allocation. + * \return Pointer to the requested storage. + */ + void *allocate_async(size_t bytes, stream_view stream) { + return allocate_async(bytes, default_alignment, stream); + } + /*! + * \brief Allocates storage of size at least `__bytes` bytes in stream order + * on `__stream`. + * + * The returned storage is aligned to the specified `__alignment` if such + * alignment is supported. + * + * The returned storage may be used immediately only on `__stream`. Using it + * on any other stream (or the host) requires first synchronizing with + * `__stream`. + * + * \throws If the storage of the requested size and alignment cannot be + * obtained. + * + * \param __bytes The size in bytes of the allocation. + * \param __alignment The alignment of the allocation + * \param __stream The stream on which to perform the allocation. + * \return Pointer to the requested storage. + */ + void *allocate_async(size_t bytes, size_t alignment, stream_view stream) { + return do_allocate_async(bytes, alignment, stream); + } + + /*! + * \brief Deallocates the storage pointed to by `__p` in stream order on + * `__stream`. + * + * `__p` must have been returned by a prior call to + * `allocate_async(__bytes, default_alignment)` or `allocate(__bytes, + * default_alignment)` on a `stream_ordered_resource` that compares + * equal to `*this`, and the storage it points to must not yet have been + * deallocated, otherwise behavior is undefined. + * + * Asynchronous, stream-ordered operations on `__stream` initiated before + * `deallocate_async(__p, __bytes, __stream)` may still access the storage + * pointed to by `__p` after `deallocate_async` returns. + * + * Storage deallocated on `__stream` may be reused by a future + * call to `allocate_async` on the same stream without synchronizing + * `__stream`. Therefore, `__stream` is typically the last stream on which + * `__p` was last used. It is the caller's responsibility to ensure the + * storage pointed to by `__p` is not in use on any other stream (or the + * host), or behavior is undefined. + * + * \param __p Pointer to storage to be deallocated. + * \param __bytes The size in bytes of the allocation. This must be equal to + * the value of `__bytes` that was specified to the `allocate` or + * `allocate_async` call that returned `__p`. + * \param __stream The stream on which to perform the deallocation. + */ + void deallocate_async(void *__mem, size_t __bytes, stream_view __stream) { + deallocate_async(__mem, __bytes, default_alignment, __stream); + } + + /*! + * \brief Deallocates the storage pointed to by `__p` in stream order on + * `__stream`. + * + * `__p` must have been returned by a prior call to + * `allocate_async(__bytes, __alignment)` or `allocate(__bytes, + * __alignment)` on a `stream_ordered_resource` that compares + * equal to `*this`, and the storage it points to must not yet have been + * deallocated, otherwise behavior is undefined. + * + * Asynchronous, stream-ordered operations on `__stream` initiated before + * `deallocate_async(__p, __bytes, __stream)` may still access the storage + * pointed to by `__p` after `deallocate_async` returns. + * + * Storage deallocated on `__stream` may be reused by a future + * call to `allocate_async` on the same stream without synchronizing + * `__stream`. Therefore, `__stream` is typically the last stream on which + * `__p` was last used. It is the caller's responsibility to ensure the + * storage pointed to by `__p` is not in use on any other stream (or the + * host), or behavior is undefined. + * + * \param __p Pointer to storage to be deallocated. + * \param __bytes The size in bytes of the allocation. This must be equal to + * the value of `__bytes` that was specified to the `allocate` or + * `allocate_async` call that returned `__p`. + * \param __alignment The alignment of the allocation. This must be equal to + * the value of `__alignment` that was specified to the `allocate` or + * `allocate_async` call that returned `__p`. + * \param __stream The stream on which to perform the deallocation. + */ + void deallocate_async(void *__mem, size_t __bytes, size_t __alignment, stream_view __stream) { + do_deallocate_async(__mem, __bytes, __alignment, __stream); + } + +protected: + virtual void *do_allocate_async(size_t __bytes, size_t __alignment, stream_view __stream) = 0; + virtual void do_deallocate_async(void *__mem, size_t __bytes, size_t __alignment, stream_view __stream) = 0; + + template + friend class cuda::pmr::basic_resource_ptr; +}; + +} // namespace detail + +/*! + * \brief Abstract interface for memory allocation. + * + * \tparam _MemoryKind The kind of the allocated memory. + */ +template +class memory_resource : private virtual detail::memory_resource_base { +public: + using memory_kind = _MemoryKind; + static constexpr std::size_t default_alignment = memory_resource_base::default_alignment; + + virtual ~memory_resource() = default; + + /*! + * \brief Allocates storage of size at least `__bytes` bytes. + * + * The returned storage is aligned to the specified `__alignment` if such + * alignment is supported. Otherwise throws. + * + * Storage may be accessed immediately within the execution contexts that + * can access the memory. + * + * \throws If storage of the requested size and alignment cannot be obtained. + * + * \param __bytes The size in bytes of the allocation + * \param __alignment The alignment of the allocation + * \return Pointer to the requested storage + */ + void *allocate(std::size_t __bytes, + std::size_t __alignment = default_alignment) { + return do_allocate(__bytes, __alignment); + } + + /*! + * \brief Deallocates the storage pointed to by `__p`. + * + * `__p` must have been returned by a prior call to `allocate(__bytes, + * __alignment)` on a `memory_resource` that compares equal to `*this`, and + * the storage it points to must not yet have been deallocated, otherwise + * behavior is undefined. + * + * \throws Nothing. + * + * \param __p Pointer to storage to be deallocated + * \param __bytes The size in bytes of the allocation. This must be equal to + * the value of `__bytes` that was specified to the `allocate` call that + * returned `__p`. + * \param __alignment The alignment of the allocation. This must be equal to + * the value of `__alignment` that was specified to the `allocate` call that + * returned `__p`. + */ + void deallocate(void *__p, std::size_t __bytes, + std::size_t __alignment = default_alignment) { + do_deallocate(__p, __bytes, __alignment); + } + + /*! + * \brief Compare this resource to another. + * + * Two resources compare equal if and only if memory allocated from one + * resource can be deallocated from the other and vice versa. + * + * \param __other The other resource to compare against + */ + bool is_equal(memory_resource const& __other) const noexcept { + return do_is_equal(__other); + } + +private: + template + friend class basic_resource_ptr; + + void *do_allocate(std::size_t __bytes, std::size_t __alignment) override = 0; + + void do_deallocate(void *__p, std::size_t __bytes, + std::size_t __alignment) override = 0; + + // Default to identity comparison + virtual bool do_is_equal(memory_resource const &__other) const noexcept { + return this == &__other; + } + + void *__do_as_kind(const ::std::type_info *__tag_type_id, const void *__tag_type_fallback_id) const noexcept final { + using __tag = detail::__type_pack; + return detail::__fallback_typeid::__is_type<__tag>(__tag_type_id, __tag_type_fallback_id) + ? const_cast(this) : nullptr; + } + + bool is_equal_base(const detail::memory_resource_base &__other) const noexcept final { + if (auto *__other_res = __other.as_kind()) { + return do_is_equal(*__other_res); + } else { + return false; + } + } +}; + +template +inline _LIBCUDACXX_INLINE_VISIBILITY +bool operator==(const memory_resource<_Kind> &__a, const memory_resource<_Kind> &__b) { + return __a.is_equal(__b); +} + +template +inline _LIBCUDACXX_INLINE_VISIBILITY +bool operator!=(const memory_resource<_Kind> &__a, const memory_resource<_Kind> &__b) { + return !__a.is_equal(__b); +} + +/*! + * \brief Abstract interface for CUDA stream-ordered memory allocation. + * + * "Stream-ordered memory allocation" extends the CUDA programming model to + * include memory allocation as stream-ordered operations. + * + * All asynchronous accesses of the allocation must happen between the stream + * execution of the allocation and the free. If storage is accessed outside of + * the promised stream order, a use before allocation / use after free error + * will cause undefined behavior. + * + * Allocating on stream `s0` returns memory that is valid to access immediately + * only on `s0`. Accessing it on any other stream (or the host) first requires + * synchronization with `s0`, otherwise behavior is undefined. + * + * Deallocating memory on stream `s1` indicates that it is valid to reuse the + * deallocated memory immediately for another allocation on `s1`. + * + * Asynchronous, stream-ordered operations ordered before deallocation on `s1` + * may still access the storage after deallocation completes. + * + * Memory may be allocated and deallocated on different streams, `s0` and `s1` + * respectively, but requires synchronization between `s0` and `s1` before the + * deallocation occurs. + * + * \tparam _MemoryKind The kind of the allocated memory. + */ +template +class stream_ordered_resource : public virtual memory_resource<_MemoryKind>, + private virtual detail::stream_ordered_resource_base { +public: + using memory_kind = _MemoryKind; + static constexpr _CUDA_VSTD::size_t default_alignment = memory_resource<_MemoryKind>::default_alignment; + + /*! + * \brief Allocates storage of size at least `__bytes` bytes in stream order + * on `__stream`. + * + * The returned storage is aligned to `default_alignment`. + * + * The returned storage may be used immediately only on `__stream`. Accessing + * it on any other stream (or the host) requires first synchronizing with + * `__stream`. + * + * \throws If the storage of the requested size and `default_alignment` cannot + * be obtained. + * + * \param __bytes The size in bytes of the allocation. + * \param __stream The stream on which to perform the allocation. + * \return Pointer to the requested storage. + */ + void *allocate_async(std::size_t __bytes, stream_view __stream) { + return do_allocate_async(__bytes, default_alignment, __stream); + } + + /*! + * \brief Allocates storage of size at least `__bytes` bytes in stream order + * on `__stream`. + * + * The returned storage is aligned to the specified `__alignment` if such + * alignment is supported. + * + * The returned storage may be used immediately only on `__stream`. Using it + * on any other stream (or the host) requires first synchronizing with + * `__stream`. + * + * \throws If the storage of the requested size and alignment cannot be + * obtained. + * + * \param __bytes The size in bytes of the allocation. + * \param __alignment The alignment of the allocation + * \param __stream The stream on which to perform the allocation. + * \return Pointer to the requested storage. + */ + void *allocate_async(std::size_t __bytes, std::size_t __alignment, + stream_view __stream) { + return do_allocate_async(__bytes, __alignment, __stream); + } + + /*! + * \brief Deallocates the storage pointed to by `__p` in stream order on + * `__stream`. + * + * `__p` must have been returned by a prior call to + * `allocate_async(__bytes, default_alignment)` or `allocate(__bytes, + * default_alignment)` on a `stream_ordered_resource` that compares + * equal to `*this`, and the storage it points to must not yet have been + * deallocated, otherwise behavior is undefined. + * + * Asynchronous, stream-ordered operations on `__stream` initiated before + * `deallocate_async(__p, __bytes, __stream)` may still access the storage + * pointed to by `__p` after `deallocate_async` returns. + * + * Storage deallocated on `__stream` may be reused by a future + * call to `allocate_async` on the same stream without synchronizing + * `__stream`. Therefore, `__stream` is typically the last stream on which + * `__p` was last used. It is the caller's responsibility to ensure the + * storage pointed to by `__p` is not in use on any other stream (or the + * host), or behavior is undefined. + * + * \param __p Pointer to storage to be deallocated. + * \param __bytes The size in bytes of the allocation. This must be equal to + * the value of `__bytes` that was specified to the `allocate` or + * `allocate_async` call that returned `__p`. + * \param __stream The stream on which to perform the deallocation. + */ + void deallocate_async(void *__p, std::size_t __bytes, stream_view __stream) { + do_deallocate_async(__p, __bytes, default_alignment, __stream); + } + + /*! + * \brief Deallocates the storage pointed to by `__p` in stream order on + * `__stream`. + * + * `__p` must have been returned by a prior call to + * `allocate_async(__bytes, __alignment)` or `allocate(__bytes, + * __alignment)` on a `stream_ordered_resource` that compares + * equal to `*this`, and the storage it points to must not yet have been + * deallocated, otherwise behavior is undefined. + * + * Asynchronous, stream-ordered operations on `__stream` initiated before + * `deallocate_async(__p, __bytes, __stream)` may still access the storage + * pointed to by `__p` after `deallocate_async` returns. + * + * Storage deallocated on `__stream` may be reused by a future + * call to `allocate_async` on the same stream without synchronizing + * `__stream`. Therefore, `__stream` is typically the last stream on which + * `__p` was last used. It is the caller's responsibility to ensure the + * storage pointed to by `__p` is not in use on any other stream (or the + * host), or behavior is undefined. + * + * \param __p Pointer to storage to be deallocated. + * \param __bytes The size in bytes of the allocation. This must be equal to + * the value of `__bytes` that was specified to the `allocate` or + * `allocate_async` call that returned `__p`. + * \param __alignment The alignment of the allocation. This must be equal to + * the value of `__alignment` that was specified to the `allocate` or + * `allocate_async` call that returned `__p`. + * \param __stream The stream on which to perform the deallocation. + */ + void deallocate_async(void *__p, std::size_t __bytes, std::size_t __alignment, + stream_view __stream) { + do_deallocate_async(__p, __bytes, __alignment, __stream); + } + +private: + template + friend class basic_resource_ptr; + + /// Default synchronous implementation of `memory_resource::do_allocate` + void *do_allocate(std::size_t __bytes, std::size_t __alignment) override { + auto const __default_stream = stream_view{}; + auto __p = do_allocate_async(__bytes, __alignment, __default_stream); + __default_stream.wait(); + return __p; + } + + /// Default synchronous implementation of `memory_resource::do_deallocate` + void do_deallocate(void *__p, std::size_t __bytes, + std::size_t __alignment) override { + auto const __default_stream = stream_view{}; + __default_stream.wait(); + do_deallocate_async(__p, __bytes, __alignment, __default_stream); + } + + void *do_allocate_async(std::size_t __bytes, std::size_t __alignment, + stream_view __stream) override = 0; + + void do_deallocate_async(void *__p, std::size_t __bytes, + std::size_t __alignment, + stream_view __stream) override = 0; +}; + + +/*! + * \brief Indicates whether a memory kind `_MemoryKind` has a property `__property`. + */ +template +struct kind_has_property : std::false_type {}; + +/*! + * \brief A special property telling that given resource/resource view allocates + * memory of specific kind. + * + * When a view defines this property, it implicitly has all properties of this + * memory kind. + * This property is also a property in itself and views defining properties of + * the underlying memory kind cannot be converted to a view defining this property. + * This allows for future extension of the set of properties. + */ +template +struct of_kind; + +template +struct kind_has_property<_MemoryKind, of_kind<_MemoryKind>> : _CUDA_VSTD::true_type {}; + + +template +struct property_list {}; + +namespace detail { + template + struct __kind_from_properties_helper { + using type = void; + }; + + template + using __kind_from_properties_helper_t = typename __kind_from_properties_helper<_Properties...>::type; + + template + struct __kind_from_properties_helper> { + using type = _Kind; + }; + + template + struct __kind_from_properties_helper<_FirstProp, _Tail...> { + using type = __kind_from_properties_helper_t<_Tail...>; + }; + + /*! + * \brief Evaluaes to the a memory kind tag form _Properties (if listed as of_kind) + * or to a list of properties otherwise. + */ + template + struct kind_from_properties { + using type = _CUDA_VSTD::conditional_t< + _CUDA_VSTD::is_same, void>::value, + property_list<_Properties...>, + detail::__kind_from_properties_helper_t<_Properties...> + >; + }; + + template + struct kind_to_properties { + using type = property_list>; + }; + + template + struct kind_to_properties> { + using type = property_list<_Properties...>; + }; +} + + +namespace memory_kind { + /*! + * \brief A pseudokind + */ + template + using with_properties = typename detail::kind_from_properties<_Properties...>::type; +}; + +template +using kind_to_properties = typename detail::kind_to_properties<_Kind>::type; + +#define _LIBCUDACXX_MEMORY_KIND_PROPERTY(__kind, __property)\ +template <> struct kind_has_property<__kind, __property> : _CUDA_VSTD::true_type {}; + + +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::host, memory_access::host); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::host, oversubscribable); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::host, memory_location::host); + +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::pinned, memory_access::host); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::pinned, memory_access::device); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::pinned, resident); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::pinned, memory_location::host); + +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::device, memory_access::device); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::device, resident); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::device, memory_location::device); + +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::managed, memory_access::host); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::managed, memory_access::device); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::managed, oversubscribable); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::managed, memory_location::host); +_LIBCUDACXX_MEMORY_KIND_PROPERTY(memory_kind::managed, memory_location::device); + +namespace detail { + template + _CUDA_VSTD::false_type Has_Property(...); + template + kind_has_property<_MemoryKind, _Property> Has_Property(const memory_resource<_MemoryKind> *); + template + kind_has_property<_MemoryKind, _Property> Has_Property(const stream_ordered_resource<_MemoryKind> *); +} // namespace detail + +template +struct has_property : decltype(detail::Has_Property<_Property>(_CUDA_VSTD::declval<_CUDA_VSTD::remove_reference_t<_Target>*>())) {}; + +template +struct has_properties : _CUDA_VSTD::conjunction...> {}; + +template +struct has_properties<_Target, property_list<_Properties...>> : has_properties<_Target, _Properties...> {}; + +template +struct has_property<_Target, property_list<_Properties...>> : has_properties<_Target, _Properties...> {}; + +namespace detail { +template +struct is_property_in : _CUDA_VSTD::false_type {}; + +template +struct is_property_in<_Property, property_list<_Properties...>> : is_property_in<_Property, _Properties...> {}; + +template +struct is_property_in<_Property, _Mismatch, _Tail...> : is_property_in<_Property, _Tail...> {}; + +template +struct is_property_in<_Property, _Property, _Tail...> : _CUDA_VSTD::true_type {}; + +template +struct is_property_in<_Property, of_kind<_MemoryKind>, Tail...> : kind_has_property<_MemoryKind, _Property> {}; + +template +struct is_property_in, of_kind<_MemoryKind>, Tail...> : _CUDA_VSTD::true_type {}; + +template +struct is_resource_pointer_convertible : _CUDA_VSTD::is_convertible<_FromPointer, _ToPointer> {}; +// Private inheritance from (stream_ordered_)memory_resource_base* requires explicit partial specializations as `is_convertible` will return false + +template +struct is_resource_pointer_convertible<_FromPointer, detail::memory_resource_base*> + : _CUDA_VSTD::conjunction<_CUDA_VSTD::is_pointer<_FromPointer>, + _CUDA_VSTD::is_base_of::element_type>> {}; + +template +struct is_resource_pointer_convertible<_FromPointer, detail::stream_ordered_resource_base*> + : _CUDA_VSTD::conjunction<_CUDA_VSTD::is_pointer<_FromPointer>, + _CUDA_VSTD::is_base_of::element_type>> {}; +} // namespace detail + +template +struct kind_has_property, _Property> : detail::is_property_in<_Property, _Properties...> {}; + +template +struct has_property, _Property> : detail::is_property_in<_Property, _Properties...> {}; + +template +struct is_view_convertible; + +template +struct is_view_convertible< + basic_resource_ptr<_FromPointer, _FromProperties...>, + basic_resource_ptr<_ToPointer, _ToProperties...>> + : _CUDA_VSTD::conjunction, + has_properties, _ToProperties...>> {}; + +/*! + * \brief A pointer-like object to a memory resource based on resource. + * + * Resource view is an object that acts as a memory resource pointer, but provides + * enhanced implicit conversions. The idea behind this type is that a user of + * a memory resource may be interested in many kinds of resources as long as they + * have certain properties. For example, a function may work with any resource + * that can provide host-accessible memory, regardless of whether it is plain host + * memory, pinned memory, managed memory, or some,yet-to-be-defined future kind + * of memory. + * + * A resource view can be created from a memory resource pointer or from another + * resource view that defines a superset of the target properties. + * + * The resource view exposes the underlying resource's interface via `operator->`. + * + * The `basic_resource_ptr` class can be parameterized with the resource pointer type, + * which can be either one of the base resource classes or a concrete resource type. + * + * \tparam _ResourcePointer a pointer-like object to the underlying memory resource + * \tparam _Properies properties of a memory resource required by resource view + */ +template +class basic_resource_ptr<_ResourcePointer, property_list<_Properties...>> { +public: + static_assert( + _CUDA_VSTD::is_base_of::element_type>::value || + _CUDA_VSTD::is_base_of::element_type>::value, + "ResourcePointer must be a pointer to a memory_resource_base, stream_ordered_resource_base or a derived class"); + + basic_resource_ptr() = default; + + basic_resource_ptr(int) = delete; + + basic_resource_ptr(std::nullptr_t) {} + + basic_resource_ptr(const basic_resource_ptr &) = default; + + basic_resource_ptr(basic_resource_ptr &&) = default; + + using properties = property_list<_Properties...>; + + /*! + * \brief Constructs a resource view from a compatible memory resource pointer. + * + * The memory resource is considered compatible if a pointer to it can be converted to + * `_ResourcePointer` and the resource type has the required properties listed + * in `_Properties`. + * + * \tparam _Resource Type of a mmeory resource object. + * \param __p pointer to a memory resource object. + */ + template ::value && + _CUDA_VSTD::conjunction...>::value + >> + basic_resource_ptr(_Resource *__p) : __pointer(__p) {} + + /*! + * \brief Constructs a resource view by copying the resource pointer from a compatible resource view. + * + * A resource view is considered compatible if it defines all properties required by this + * view in `_Properties`. + * + * \tparam _OtherPointer The resource pointer type of the source resource view + * \tparam _OtherProperties The properties defined byt the source resource view + */ + template , basic_resource_ptr>::value + >> + basic_resource_ptr(basic_resource_ptr<_OtherPointer, _OtherProperties...> v) : __pointer(v.__pointer) {} + + basic_resource_ptr &operator=(const basic_resource_ptr &)=default; + basic_resource_ptr &operator=(basic_resource_ptr &&)=default; + + template , basic_resource_ptr>::value + >> + basic_resource_ptr &operator=(basic_resource_ptr<_OtherPointer, _OtherProperties...> v) { + __pointer = std::move(v.__pointer); + return *this; + } + + + /*! + * \brief Exposes the interface of the underlying memory resource. + * + * \note This method should not be used to obtain the pointer to the memory resource. + */ + _ResourcePointer operator->() const noexcept { return __pointer; } + + template + bool operator==(const cuda::pmr::basic_resource_ptr<_Ptr2, _Props2...> &__v2) const noexcept { + if (__pointer == nullptr || __v2.__pointer == nullptr) + return __pointer == nullptr && __v2.__pointer == nullptr; + return static_cast(__pointer)->is_equal(*__v2.__pointer); + } + + template + bool operator!=(const cuda::pmr::basic_resource_ptr<_Ptr2, _Props2...> &__v2) const noexcept { + return !(*this == __v2); + } + + /*! + * \brief Returns true if the underlying pointer is not null. + */ + constexpr explicit operator bool() const noexcept { + return !!__pointer; + } + +private: + template + friend class basic_resource_ptr; + + _ResourcePointer __pointer{}; +}; + +template +class basic_resource_ptr : public basic_resource_ptr<_ResourcePointer, property_list<_Properties...>> { +public: + using __base = basic_resource_ptr<_ResourcePointer, property_list<_Properties...>>; + using __base::__base; + using __base::operator=; + using __base::operator==; + using __base::operator!=; + using __base::operator->; + using __base::operator bool; + using properties = property_list<_Properties...>; + + /*! + * \brief Constructs a resource view from a compatible memory resource pointer. + * + * The memory resource is considered compatible if a pointer to it can be converted to + * `_ResourcePointer` and the resource type has the required properties listed + * in `_Properties`. + * + * \tparam _Resource Type of a mmeory resource object. + * \param __p pointer to a memory resource object. + */ + template ::value && + _CUDA_VSTD::conjunction...>::value + >> + basic_resource_ptr(_Resource *__p) : __base(__p) {} + + /*! + * \brief Constructs a resource view by copying the resource pointer from a compatible resource view. + * + * A resource view is considered compatible if it defines all properties required by this + * view in `_Properties`. + * + * \tparam _OtherPointer The resource pointer type of the source resource view + * \tparam _OtherProperties The properties defined byt the source resource view + */ + template , basic_resource_ptr>::value + >> + basic_resource_ptr(basic_resource_ptr<_OtherPointer, _OtherProperties...> v) : __base(v) {} +}; + +template +basic_resource_ptr<_ResourcePointer, _FirstProperty, _Properties...> +view_resource(_ResourcePointer __rsrc_ptr) { + return __rsrc_ptr; +} + +template +basic_resource_ptr<_ResourcePointer, kind_to_properties::memory_kind>> +view_resource(_ResourcePointer __rsrc_ptr) { + return __rsrc_ptr; +} + + +template +bool operator==(const basic_resource_ptr<_ResourcePointer, _Properties...> &__view, const memory_resource<_Kind> *__mr) { + return __view == view_resource(__mr); +} + +template +bool operator!=(const basic_resource_ptr<_ResourcePointer, _Properties...> &__view, const memory_resource<_Kind> *__mr) { + return __view != view_resource(__mr); +} + +template +bool operator==(const memory_resource<_Kind> *__mr, const basic_resource_ptr<_ResourcePointer, _Properties...> &__view) { + return view_resource(__mr) == __view; +} + +template +bool operator!=(const memory_resource<_Kind> *__mr, const basic_resource_ptr<_ResourcePointer, _Properties...> &__view) { + return view_resource(__mr) != __view; +} + +template +using resource_ptr = basic_resource_ptr; + +template +using stream_ordered_resource_ptr = basic_resource_ptr; + +#if _LIBCUDACXX_STD_VER > 14 + +#if defined(_LIBCUDACXX_STD_PMR_NS) + +namespace detail{ +class __pmr_adaptor_base : public _LIBCUDACXX_STD_PMR_NS::memory_resource { +public: + virtual cuda::pmr::memory_resource* resource() const noexcept = 0; +}; +} + +template +class pmr_adaptor final : public detail::__pmr_adaptor_base { + + using resource_type = _CUDA_VSTD::remove_reference_t())>; + + static constexpr bool __is_host_accessible_resource = has_property::value; + static_assert( + __is_host_accessible_resource, + "Pointer must be a pointer-like type to a type that allocates host-accessible memory."); + +public: + pmr_adaptor(_Pointer __mr) : __mr_{std::move(__mr)} {} + + using raw_pointer = _CUDA_VSTD::remove_reference_t())>; + + raw_pointer resource() const noexcept override { return &*__mr_; } + +private: + void *do_allocate(std::size_t __bytes, std::size_t __alignment) override { + return __mr_->allocate(__bytes, __alignment); + } + + void do_deallocate(void *__p, std::size_t __bytes, + std::size_t __alignment) override { + return __mr_->deallocate(__p, __bytes, __alignment); + } + + bool do_is_equal(_LIBCUDACXX_STD_PMR_NS::memory_resource const &__other) const noexcept override { + auto __other_p = dynamic_cast(&__other); + return __other_p and (__other_p->resource() == resource() or + __other_p->resource()->is_equal(*resource())); + } + + _Pointer __mr_; +}; +#endif // defined(_LIBCUDACXX_STD_PMR_NS) +#endif // _LIBCUDACXX_STD_VER > 14 + +} // namespace pmr +_LIBCUDACXX_END_NAMESPACE_CUDA + +#include "std/detail/__pragma_pop" + +#endif //_CUDA_MEMORY_RESOURCE diff --git a/include/cuda/std/detail/libcxx/include/__config b/include/cuda/std/detail/libcxx/include/__config index 5b998eec87..8909fd7738 100644 --- a/include/cuda/std/detail/libcxx/include/__config +++ b/include/cuda/std/detail/libcxx/include/__config @@ -443,11 +443,19 @@ typedef __char16_t char16_t; typedef __char32_t char32_t; #endif -#if !(__has_feature(cxx_exceptions)) && !defined(_LIBCUDACXX_NO_EXCEPTIONS) +#if __has_feature(cxx_exceptions) +#define _LIBCUDACXX_EXT_EXCEPTIONS_ENABLED +#endif + +#if !defined(_LIBCUDACXX_EXT_EXCEPTIONS_ENABLED) && !defined(_LIBCUDACXX_NO_EXCEPTIONS) #define _LIBCUDACXX_NO_EXCEPTIONS #endif -#if !(__has_feature(cxx_rtti)) && !defined(_LIBCUDACXX_NO_RTTI) +#if __has_feature(cxx_rtti) +#define _LIBCUDACXX_EXT_RTTI_ENABLED +#endif + +#if !defined(_LIBCUDACXX_EXT_RTTI_ENABLED) && !defined(_LIBCUDACXX_NO_RTTI) #define _LIBCUDACXX_NO_RTTI #endif @@ -744,7 +752,12 @@ typedef __char32_t char32_t; // #define _LIBCUDACXX_IS_SAME(...) __is_same_as(__VA_ARGS__) #endif -#if !__EXCEPTIONS && !defined(_LIBCUDACXX_NO_EXCEPTIONS) +// GCC defines __EXCEPTIONS while MSVC uses _CPPUNWIND +#if defined(__EXCEPTIONS) || defined(_CPPUNWIND) +#define _LIBCUDACXX_EXT_EXCEPTIONS_ENABLED +#endif + +#if !defined(_LIBCUDACXX_EXT_EXCEPTIONS_ENABLED) && !defined(_LIBCUDACXX_NO_EXCEPTIONS) #define _LIBCUDACXX_NO_EXCEPTIONS #endif @@ -1478,16 +1491,19 @@ extern "C" _LIBCUDACXX_FUNC_VIS void __sanitizer_annotate_contiguous_container( const void *, const void *, const void *, const void *); #endif -// Try to find out if RTTI is disabled. + +// Try to find out if RTTI is enabled. // g++ and cl.exe have RTTI on by default and define a macro when it is. -// g++ only defines the macro in 4.3.2 and onwards. +// g++ only defines the macro in 4.3.2 and onwards, but we only care about 4.8 and up +#if defined(__GNUC__) && defined(__GXX_RTTI) +# define _LIBCUDACXX_EXT_RTTI_ENABLED +#elif defined(_LIBCUDACXX_COMPILER_MSVC) && defined(_CPPRTTI) +# define _LIBCUDACXX_EXT_RTTI_ENABLED +#endif + +// Disable RTTI if it wasn't requested earlier #if !defined(_LIBCUDACXX_NO_RTTI) -# if defined(__GNUC__) && \ - ((__GNUC__ >= 5) || \ - (__GNUC__ == 4 && (__GNUC_MINOR__ >= 3 || __GNUC_PATCHLEVEL__ >= 2))) && \ - !defined(__GXX_RTTI) -# define _LIBCUDACXX_NO_RTTI -# elif defined(_LIBCUDACXX_COMPILER_MSVC) && !defined(_CPPRTTI) +# if !defined(_LIBCUDACXX_EXT_RTTI_ENABLED) # define _LIBCUDACXX_NO_RTTI # endif #endif diff --git a/include/cuda/stream_view b/include/cuda/stream_view new file mode 100644 index 0000000000..ed1aa2a77e --- /dev/null +++ b/include/cuda/stream_view @@ -0,0 +1,131 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_STREAM +#define _CUDA_STREAM + +#include "std/cstddef" +#include "std/detail/__config" +#include "std/detail/__pragma_push" +#include "std/version" +#include "detail/__cuda_util" + +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA + +/** + * \brief A non-owning wrapper for a `cudaStream_t`. + * + * `stream_view` is a non-owning "view" type similar to `std::span` or `std::string_view`. + * \see https://en.cppreference.com/w/cpp/container/span and + * \see https://en.cppreference.com/w/cpp/string/basic_string_view + * + */ +class stream_view { +public: + + using value_type = ::cudaStream_t; + + /** + * \brief Constructs a `stream_view` of the "default" CUDA stream. + * + * For behavior of the default stream, + * \see https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html + * + */ + constexpr stream_view() noexcept = default; + + /** + * \brief Constructs a `stream_view` from a `cudaStream_t` handle. + * + * This constructor provides implicit conversion from `cudaStream_t`. + * + * \note: It is the callers responsibilty to ensure the `stream_view` does not + * outlive the stream identified by the `cudaStream_t` handle. + * + */ + constexpr stream_view(value_type stream) : __stream{stream} {} + + /// Disallow construction from an `int`, e.g., `0`. + stream_view(int) = delete; + + /// Disallow construction from `nullptr`. + stream_view(std::nullptr_t) = delete; + + /// Returns the wrapped `cudaStream_t` handle. + constexpr value_type get() const noexcept { return __stream; } + + /** + * \brief Synchronizes the wrapped stream. + * + * \throws cuda::cuda_error if synchronization fails. + * + */ + void wait() const { + detail::__throw_on_cuda_error(::cudaStreamSynchronize(get()), + "Failed to synchronize stream."); + } + + /** + * \brief Queries if all operations on the wrapped stream have completed. + * + * \throws cuda::cuda_error if the query fails. + * + * \return `true` if all operations have completed, or `false` if not. + */ + bool ready() const{ + auto const __result = ::cudaStreamQuery(get()); + if(__result == ::cudaSuccess){ + return true; + } else if (__result == ::cudaErrorNotReady){ + return false; + } + detail::__throw_on_cuda_error(__result); + return false; + } + +private: + value_type __stream{0}; ///< Handle of the viewed stream +}; + +/** + * \brief Compares two `stream_view`s for equality + * + * \note Allows comparison with `cudaStream_t` due to implicit conversion to + * `stream_view`. + * + * \param lhs The first `stream_view` to compare + * \param rhs The second `stream_view` to compare + * \return true if equal, false if unequal + */ +inline constexpr bool operator==(stream_view __lhs, stream_view __rhs) { + return __lhs.get() == __rhs.get(); +} + +/** + * \brief Compares two `stream_view`s for inequality + * + * \note Allows comparison with `cudaStream_t` due to implicit conversion to + * `stream_view`. + * + * \param lhs The first `stream_view` to compare + * \param rhs The second `stream_view` to compare + * \return true if unequal, false if equal + */ +inline constexpr bool operator!=(stream_view __lhs, stream_view __rhs) { + return not(__lhs == __rhs); +} + +_LIBCUDACXX_END_NAMESPACE_CUDA + +#include "std/detail/__pragma_pop" + +#endif //_CUDA_STREAM