From 5e6cdb1deb9da77f434744d3dc22977bc59ec839 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 1/9] Implement `cuda::stream_ref` --- .../stream_ref.constructor.pass.cpp | 46 +++++ .../stream_ref/stream_ref.equality.pass.cpp | 26 +++ .../cuda/stream_ref/stream_ref.get.pass.cpp | 21 ++ .../cuda/stream_ref/stream_ref.ready.fail.cpp | 32 +++ .../cuda/stream_ref/stream_ref.ready.pass.cpp | 33 ++++ .../cuda/stream_ref/stream_ref.wait.fail.cpp | 32 +++ .../cuda/stream_ref/stream_ref.wait.pass.cpp | 33 ++++ include/cuda/stream_ref | 185 ++++++++++++++++++ 8 files changed, 408 insertions(+) create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp create mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp create mode 100644 include/cuda/stream_ref diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp new file mode 100644 index 0000000000..0c6a899f7a --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp @@ -0,0 +1,46 @@ +//===----------------------------------------------------------------------===// +// +// 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-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include +#include + +static_assert(cuda::std::is_default_constructible::value); +static_assert(!cuda::std::is_constructible::value); +static_assert(!cuda::std::is_constructible::value); + +template +using void_t = void; + +template +constexpr bool has_value_type = false; + +template +constexpr bool has_value_type > = true; + +static_assert(has_value_type, ""); + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + { // default construction + cuda::stream_ref ref; + static_assert(noexcept(cuda::stream_ref{}), ""); + assert(ref.get() == reinterpret_cast(0)); + } + + { // from stream + cudaStream_t stream = reinterpret_cast(42); + cuda::stream_ref ref{stream}; + static_assert(noexcept(cuda::stream_ref{stream}), ""); + assert(ref.get() == reinterpret_cast(42)); + } +#endif // __CUDA_ARCH__ + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp new file mode 100644 index 0000000000..b749db58d3 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + cuda::stream_ref left{reinterpret_cast(42)}; + cuda::stream_ref right{reinterpret_cast(1337)}; + static_assert(noexcept(left == right), ""); + static_assert(noexcept(left != right), ""); + + assert(left == left); + assert(left != right); +#endif // __CUDA_ARCH__ + + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp new file mode 100644 index 0000000000..dcf8c1647c --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp @@ -0,0 +1,21 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + cudaStream_t stream = reinterpret_cast(42); + cuda::stream_ref ref{stream}; + assert(ref.get() == stream); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp new file mode 100644 index 0000000000..fa8bcb8891 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp @@ -0,0 +1,32 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + cudaStream_t stream = reinterpret_cast(42); + cuda::stream_ref ref{stream}; +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + try { + assert(ref.ready()); + } catch (const cuda::cuda_error& e) { + assert(cudaGetErrorString(e) == ""); + } catch (...) { + assert(false && "Should have thrown"); + } + assert(false); +#else + assert(ref.ready()); +#endif // _LIBCUDACXX_NO_EXCEPTIONS +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp new file mode 100644 index 0000000000..6df56cf894 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp @@ -0,0 +1,33 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + { // passing case + cudaStream_t stream; + cudaStreamCreate(&stream); + cuda::stream_ref ref{stream}; +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + try { + assert(ref.ready()); + } catch (...) { + assert(false && "Should not have thrown"); + } +#else + assert(ref.ready()); +#endif // _LIBCUDACXX_NO_EXCEPTIONS + cudaStreamDestroy(stream); + } +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp new file mode 100644 index 0000000000..837d03d289 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp @@ -0,0 +1,32 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + cudaStream_t stream = reinterpret_cast(42); + cuda::stream_ref ref{stream}; +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + try { + ref.wait(); + } catch (const cuda::cuda_error& e) { + assert(cudaGetErrorString(e) == "Failed to synchronize stream."); + } catch (...) { + assert(false && "Should have thrown"); + } + assert(false); +#else + ref.wait(); +#endif // _LIBCUDACXX_NO_EXCEPTIONS +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp new file mode 100644 index 0000000000..26c03ca7f5 --- /dev/null +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp @@ -0,0 +1,33 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main(int argc, char** argv) { +#ifndef __CUDA_ARCH__ + { // passing case + cudaStream_t stream; + cudaStreamCreate(&stream); + cuda::stream_ref ref{stream}; +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + try { + ref.wait(); + } catch (...) { + assert(false && "Should not have thrown"); + } +#else + ref.wait(); +#endif // _LIBCUDACXX_NO_EXCEPTIONS + cudaStreamDestroy(stream); + } +#endif + return 0; +} diff --git a/include/cuda/stream_ref b/include/cuda/stream_ref new file mode 100644 index 0000000000..2d69a0bc08 --- /dev/null +++ b/include/cuda/stream_ref @@ -0,0 +1,185 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_STREAM_REF +#define _CUDA_STREAM_REF + +// clang-format off +/* + stream_ref synopsis +namespace cuda { +class stream_ref { + using value_type = cudaStream_t; + + stream_ref() = default; + stream_ref(cudaStream_t stream_) noexcept : stream(stream_) {} + + stream_ref(int) = delete; + stream_ref(nullptr_t) = delete; + + [[nodiscard]] value_type get() const noexcept; + + void wait() const; + + [[nodiscard]] bool ready() const; + + [[nodiscard]] friend bool operator==(stream_ref, stream_ref); + [[nodiscard]] friend bool operator!=(stream_ref, stream_ref); + +private: + cudaStream_t stream = 0; // exposition only +}; +} // cuda +*/ + +#include // cuda_runtime_api needs to come first +// clang-format on + +#include +#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_ref +{ +private: + ::cudaStream_t __stream{0}; + +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 + * + */ + stream_ref() = 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_ref(value_type __stream_) noexcept + : __stream{__stream_} + {} + + /// Disallow construction from an `int`, e.g., `0`. + stream_ref(int) = delete; + + /// Disallow construction from `nullptr`. + stream_ref(_CUDA_VSTD::nullptr_t) = delete; + + /** + * \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 + */ + _LIBCUDACXX_NODISCARD_ATTRIBUTE friend bool constexpr operator==(stream_ref __lhs, stream_ref __rhs) noexcept + { + return __lhs.__stream == __rhs.__stream; + } + + /** + * \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 + */ + _LIBCUDACXX_NODISCARD_ATTRIBUTE friend bool constexpr operator!=(stream_ref __lhs, stream_ref __rhs) noexcept + { + return __lhs.__stream != __rhs.__stream; + } + + /// Returns the wrapped `cudaStream_t` handle. + _LIBCUDACXX_NODISCARD_ATTRIBUTE constexpr value_type get() const noexcept { return __stream; } + + /** + * \brief Synchronizes the wrapped stream. + * + * \throws cuda::cuda_error if synchronization fails. + * + */ + void wait() const + { + const auto __result = ::cudaStreamQuery(get()); + switch (__result) + { + case ::cudaSuccess: + return; + default: + ::cudaGetLastError(); // Clear CUDA error state +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + throw cuda::cuda_error{__result, "Failed to synchronize stream."}; +#else + // _LIBCUDACXX_UNREACHABLE; +#endif + } + } + + /** + * \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. + */ + _LIBCUDACXX_NODISCARD_ATTRIBUTE bool ready() const + { + const auto __result = ::cudaStreamQuery(get()); + switch (__result) + { + case ::cudaSuccess: + return true; + case ::cudaErrorNotReady: + return false; + default: + ::cudaGetLastError(); // Clear CUDA error state +#ifndef _LIBCUDACXX_NO_EXCEPTIONS + throw cuda::cuda_error{__result, ""}; +#else + // _LIBCUDACXX_UNREACHABLE; +#endif + return false; + } + } +}; + +_LIBCUDACXX_END_NAMESPACE_CUDA + +#include + +#endif //_CUDA_STREAM_REF From 6d5268e607ee5539ba3084cb4f7e5e95d6c0c474 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 2/9] Implement concepts for This includes: * cuda::mr::has_property * cuda::mr::has_property_with * cuda::mr::resource * cuda::mr::resource_with * cuda::mr::async_resource * cuda::mr::async_resource_with --- .../async_resource.pass.cpp | 121 ++++++++++ .../async_resource_with.pass.cpp | 84 +++++++ .../has_property.pass.cpp | 60 +++++ .../resource.pass.cpp | 75 +++++++ .../resource_with.pass.cpp | 71 ++++++ include/cuda/memory_resource | 209 ++++++++++++++++++ 6 files changed, 620 insertions(+) create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp create mode 100644 include/cuda/memory_resource diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp new file mode 100644 index 0000000000..1551e4dc44 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp @@ -0,0 +1,121 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource +#include + +#include + +#include "test_macros.h" + +struct invalid_argument {}; + +struct valid_resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const valid_resource&) const { return true; } + bool operator!=(const valid_resource&) const { return false; } +}; +static_assert(cuda::mr::async_resource, ""); + +struct invalid_allocate_missing { + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const invalid_allocate_missing&) const { return true; } + bool operator!=(const invalid_allocate_missing&) const { return false; } +}; +static_assert(!cuda::mr::async_resource, ""); + +struct invalid_deallocate_missing { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const invalid_allocate_missing&) const { return true; } + bool operator!=(const invalid_allocate_missing&) const { return false; } +}; +static_assert(!cuda::mr::async_resource, ""); + +struct invalid_allocate_async_argument { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(invalid_argument, std::size_t) { return nullptr; } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const invalid_allocate_async_argument&) const { return true; } + bool operator!=(const invalid_allocate_async_argument&) const { return false; } +}; +static_assert(!cuda::mr::async_resource, ""); + +struct invalid_allocate_async_return { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + int allocate_async(std::size_t, std::size_t, cuda::stream_ref) { return 42; } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const invalid_allocate_async_return&) const { return true; } + bool operator!=(const invalid_allocate_async_return&) const { return false; } +}; +static_assert(!cuda::mr::async_resource, ""); + +struct invalid_deallocate_async_argument { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, invalid_argument, std::size_t) {} + bool operator==(const invalid_deallocate_async_argument&) const { return true; } + bool operator!=(const invalid_deallocate_async_argument&) const { return false; } +}; +static_assert(!cuda::mr::async_resource, ""); + +struct non_comparable { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} +}; +static_assert(!cuda::mr::async_resource, ""); + +struct non_eq_comparable { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator!=(const non_eq_comparable&) const { return false; } +}; +static_assert(!cuda::mr::async_resource, ""); + +#if TEST_STD_VER < 20 +struct non_neq_comparable { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const non_neq_comparable&) const { return true; } +}; +static_assert(!cuda::mr::async_resource, ""); +#endif // TEST_STD_VER < 20 + +int main(int, char**) { return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp new file mode 100644 index 0000000000..967f701715 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp @@ -0,0 +1,84 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource_with +#include + +#include + +struct prop_with_value {}; +struct prop {}; + +struct valid_resource_with_property { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const valid_resource_with_property&) const { return true; } + bool operator!=(const valid_resource_with_property&) const { return false; } + friend void get_property(const valid_resource_with_property&, + prop_with_value) {} +}; +static_assert(cuda::mr::async_resource_with, + ""); + +struct valid_resource_without_property { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const valid_resource_without_property&) const { return true; } + bool operator!=(const valid_resource_without_property&) const { return false; } +}; +static_assert(!cuda::mr::async_resource_with, + ""); + +struct invalid_resource_with_property { + friend void get_property(const invalid_resource_with_property&, + prop_with_value) {} +}; +static_assert(!cuda::mr::async_resource_with, + ""); + +struct resource_with_many_properties { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + void deallocate_async(void*, std::size_t, std::size_t, cuda::stream_ref) {} + bool operator==(const resource_with_many_properties&) const { return true; } + bool operator!=(const resource_with_many_properties&) const { return false; } + friend void get_property(const resource_with_many_properties&, + prop_with_value) {} + friend void get_property(const resource_with_many_properties&, prop) {} +}; +static_assert(cuda::mr::async_resource_with, + ""); +static_assert(!cuda::mr::async_resource_with, + ""); + +struct derived_with_property : public valid_resource_without_property { + friend void get_property(const derived_with_property&, prop_with_value) {} +}; +static_assert( + cuda::mr::async_resource_with, ""); + +int main(int, char**) { return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp new file mode 100644 index 0000000000..9a727de082 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp @@ -0,0 +1,60 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::has_property, cuda::mr::has_property_with +#include + +struct prop_with_value { + using value_type = int; +}; +struct prop {}; + +static_assert(cuda::mr::property_with_value); +static_assert(!cuda::mr::property_with_value); + +struct valid_property { + friend void get_property(const valid_property&, prop) {} +}; +static_assert(!cuda::mr::has_property, ""); +static_assert(cuda::mr::has_property, ""); +static_assert(!cuda::mr::has_property_with, ""); + +struct valid_property_with_value { + friend int get_property(const valid_property_with_value&, prop_with_value) { + return 42; + } +}; +static_assert( + cuda::mr::has_property, ""); +static_assert(!cuda::mr::has_property, ""); +static_assert(cuda::mr::has_property_with, + ""); +static_assert(!cuda::mr::has_property_with, + ""); + +struct derived_from_property : public valid_property { + friend int get_property(const derived_from_property&, prop_with_value) { + return 42; + } +}; +static_assert(cuda::mr::has_property, + ""); +static_assert(cuda::mr::has_property, ""); +static_assert( + cuda::mr::has_property_with, + ""); +static_assert(!cuda::mr::has_property_with, + ""); + +int main(int, char**) { return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp new file mode 100644 index 0000000000..bc2d7a86f7 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp @@ -0,0 +1,75 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource +#include + +#include + +#include "test_macros.h" + +struct invalid_argument {}; + +struct valid_resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const valid_resource&) const { return true; } + bool operator!=(const valid_resource&) const { return false; } +}; +static_assert(cuda::mr::resource, ""); + +struct invalid_allocate_argument { + void* allocate(invalid_argument, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const invalid_allocate_argument&) { return true; } + bool operator!=(const invalid_allocate_argument&) { return false; } +}; +static_assert(!cuda::mr::resource, ""); + +struct invalid_allocate_return { + int allocate(std::size_t, std::size_t) { return 42; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const invalid_allocate_return&) { return true; } + bool operator!=(const invalid_allocate_return&) { return false; } +}; +static_assert(!cuda::mr::resource, ""); + +struct invalid_deallocate_argument { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, invalid_argument, std::size_t) {} + bool operator==(const invalid_deallocate_argument&) { return true; } + bool operator!=(const invalid_deallocate_argument&) { return false; } +}; +static_assert(!cuda::mr::resource, ""); + +struct non_comparable { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} +}; +static_assert(!cuda::mr::resource, ""); + +struct non_eq_comparable { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator!=(const non_eq_comparable&) { return false; } +}; +static_assert(!cuda::mr::resource, ""); + +#if TEST_STD_VER < 20 +struct non_neq_comparable { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const non_neq_comparable&) { return true; } +}; +static_assert(!cuda::mr::resource, ""); +#endif // TEST_STD_VER <20 + +int main(int, char**) { return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp new file mode 100644 index 0000000000..75020d7899 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp @@ -0,0 +1,71 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_with +#include + +#include + +struct prop_with_value {}; +struct prop {}; + +struct valid_resource_with_property { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const valid_resource_with_property&) const { return true; } + bool operator!=(const valid_resource_with_property&) const { return false; } + friend void get_property(const valid_resource_with_property&, + prop_with_value) {} +}; +static_assert( + cuda::mr::resource_with, ""); + +struct valid_resource_without_property { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const valid_resource_without_property&) const { return true; } + bool operator!=(const valid_resource_without_property&) const { return false; } +}; +static_assert( + !cuda::mr::resource_with, + ""); + +struct invalid_resource_with_property { + friend void get_property(const invalid_resource_with_property&, + prop_with_value) {} +}; +static_assert( + !cuda::mr::resource_with, + ""); + +struct resource_with_many_properties { + void* allocate(std::size_t, std::size_t) { return nullptr; } + void deallocate(void*, std::size_t, std::size_t) {} + bool operator==(const resource_with_many_properties&) const { return true; } + bool operator!=(const resource_with_many_properties&) const { return false; } + friend void get_property(const resource_with_many_properties&, + prop_with_value) {} + friend void get_property(const resource_with_many_properties&, prop) {} +}; +static_assert(cuda::mr::resource_with, + ""); +static_assert(!cuda::mr::resource_with, + ""); + +struct derived_with_property : public valid_resource_without_property { + friend void get_property(const derived_with_property&, prop_with_value) {} +}; +static_assert(cuda::mr::resource_with, + ""); + +int main(int, char**) { return 0; } diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource new file mode 100644 index 0000000000..5e86efabc1 --- /dev/null +++ b/include/cuda/memory_resource @@ -0,0 +1,209 @@ +//===----------------------------------------------------------------------===// +// +// Part of the CUDA Toolkit, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_MEMORY_RESOURCE +#define _CUDA_MEMORY_RESOURCE + +// clang-format off +/* + memory_resource synopsis +namespace cuda { +namespace mr { +template +concept resource = equality_comparable + && requires(Resource& res, void* ptr, size_t size, size_t alignment) { + { res.allocate(size, alignment) } -> same_as; + { res.deallocate(ptr, size, alignment) } -> same_as; +}; + +template +concept async_resource = resource + && requires(Resource& res, void* ptr, size_t size, size_t alignment, cuda_stream_ref stream) { + { res.allocate_async(size, alignment, stream) } -> same_as; + { res.deallocate_async(ptr, size, alignment, stream) } -> same_as; +}; + +template +concept has_property = resource && requires(const Resource& res, Property prop) { + get_property(res, prop); +}; + +template +concept property_with_value = requires { + typename Property::value_type; +}; + +template +concept has_property_with = resource + && property_with_value + && same_as + && requires(const Resource& res, Property prop) { + get_property(res, prop) -> Return; +}; + +template +concept resource_with = resource && (has_property && ...); + +template +concept async_resource_with = async_resource && (has_property && ...); + +template +class resource_ref { + template Resource> + resource_ref(Resource&) noexcept; + + void* allocate(size_t size, size_t alignment); + void deallocate(void* ptr, size_t size, size_t alignment); + + template + requires resource_with + && resource_with, Properties...> + friend bool operator==(const resource_ref& left, const resource_ref& right); + + template + requires has_property + friend typename Property::value_type get_property(const resource_ref& ref, Property) noexcept; + + template + requires (has_property && !property_with_value) + friend void get_property(const resource_ref& ref, Property) noexcept; +}; + +} // mr +} // cuda +*/ +// clang-format on +#include + +#include +#include + +#include +#include + +#if _LIBCUDACXX_STD_VER > 11 +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA +namespace mr +{ + +/////////////////////////////////////////////////////////////////////////////// +// memory_resource + +/// \concept resource +/// \brief The \c resource concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_member_allocate_, + requires(_Resource& __res, size_t __bytes, size_t __alignment) // + (_CUDA_VSTD::_Returns_exactly(__res.allocate(__bytes, __alignment)))); +template +_LIBCUDACXX_CONCEPT _Has_member_allocate = _LIBCUDACXX_FRAGMENT(_Has_member_allocate_, _Resource); + +template +_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_member_deallocate_, + requires(_Resource& __res, + void* __ptr, + size_t __bytes, + size_t __alignment) // + (__res.deallocate(__ptr, __bytes, __alignment))); +template +_LIBCUDACXX_CONCEPT _Has_member_deallocate = _LIBCUDACXX_FRAGMENT(_Has_member_deallocate_, _Resource); + +template +_LIBCUDACXX_CONCEPT resource = + _Has_member_allocate<_Resource>&& _Has_member_deallocate<_Resource>&& _CUDA_VSTD::equality_comparable<_Resource>; + +/// \concept async_resource +/// \brief The \c async_resource concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT( + _Has_member_allocate_async_, + requires(_Resource& __res, + size_t __bytes, + size_t __alignment, + cuda::stream_ref __stream) // + (_CUDA_VSTD::_Returns_exactly(__res.allocate_async(__bytes, __alignment, __stream)))); +template +_LIBCUDACXX_CONCEPT _Has_member_allocate_async = _LIBCUDACXX_FRAGMENT(_Has_member_allocate_async_, _Resource); + +template +_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_member_deallocate_async_, + requires(_Resource& __res, + void* __ptr, + size_t __bytes, + size_t __alignment, + cuda::stream_ref __stream) // + (__res.deallocate_async(__ptr, __bytes, __alignment, __stream))); +template +_LIBCUDACXX_CONCEPT _Has_member_deallocate_async = _LIBCUDACXX_FRAGMENT(_Has_member_deallocate_async_, _Resource); + +/// \concept async_resource +/// \brief The \c async_resource concept +template +_LIBCUDACXX_CONCEPT async_resource = + resource<_Resource>&& _Has_member_allocate_async<_Resource>&& _Has_member_deallocate_async<_Resource>; + +/// \concept has_property +/// \brief The \c has_property concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_property_, + requires(const _Resource& __res) // + (get_property(__res, _Property{}))); +template +_LIBCUDACXX_CONCEPT has_property = _LIBCUDACXX_FRAGMENT(_Has_property_, _Resource, _Property); + +template +using _Has_value_types = typename T::value_type; + +/// \concept property_with_value +/// \brief The \c property_with_value concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT(_Property_with_value_, + requires() // + (typename(_Has_value_types<_Property>))); +template +_LIBCUDACXX_CONCEPT property_with_value = _LIBCUDACXX_FRAGMENT(_Property_with_value_, _Property); + +/// \concept has_property_with +/// \brief The \c has_property_with concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_property_with_, + requires(const _Resource& __res) // + (_CUDA_VSTD::_Returns_exactly<_Return>(get_property(__res, _Property{})))); +template +_LIBCUDACXX_CONCEPT has_property_with = property_with_value<_Property>&& + _LIBCUDACXX_FRAGMENT(_Has_property_with_, _Resource, _Property, _Return); + +/// \concept resource_with +/// \brief The \c resource_with concept +template +#if _LIBCUDACXX_STD_VER < 17 +_LIBCUDACXX_CONCEPT resource_with = + resource<_Resource>&& _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant>...>; +#else +_LIBCUDACXX_CONCEPT resource_with = resource<_Resource> && (has_property<_Resource, _Properties> && ...); +#endif + +/// \concept async_resource_with +/// \brief The \c async_resource_with concept +template +#if _LIBCUDACXX_STD_VER < 17 +_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource>&& + _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant>...>; +#else +_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> // + && (has_property<_Resource, _Properties> && ...); +#endif + +} // namespace mr +_LIBCUDACXX_END_NAMESPACE_CUDA +#endif // _LIBCUDACXX_STD_VER > 11 + +#include + +#endif //_LIBCUDACXX_BEGIN_NAMESPACE_CUDA From 4a2656a25546f880bbaf9eec1ae540093710f4de Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 3/9] Implement resource wrappers for This includes * `cuda::mr::resource_ref` * `cuda::mr::async_resource_ref` --- .../async_resource_ref.allocate.pass.cpp | 104 ++++ .../async_resource_ref.construction.pass.cpp | 115 +++++ .../async_resource_ref.conversion.pass.cpp | 123 +++++ .../async_resource_ref.equality.fail.cpp | 92 ++++ .../async_resource_ref.equality.pass.cpp | 111 +++++ .../async_resource_ref.inheritance.pass.cpp | 171 +++++++ .../async_resource_ref.properties.pass.cpp | 190 ++++++++ .../resource_ref.allocate.pass.cpp | 63 +++ .../resource_ref.construction.pass.cpp | 102 ++++ .../resource_ref.conversion.pass.cpp | 171 +++++++ .../resource_ref.equality.fail.cpp | 80 ++++ .../resource_ref.equality.pass.cpp | 99 ++++ .../resource_ref.inheritance.pass.cpp | 146 ++++++ .../resource_ref.properties.pass.cpp | 175 +++++++ include/cuda/memory_resource | 452 +++++++++++++++--- 15 files changed, 2134 insertions(+), 60 deletions(-) create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp create mode 100644 .upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp new file mode 100644 index 0000000000..8a013ddf80 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp @@ -0,0 +1,104 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource_ref properties +#include + +#include +#include + +struct async_resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const async_resource& other) const { + return _val == other._val; + } + bool operator!=(const async_resource& other) const { + return _val != other._val; + } + + int _val = 0; +}; + +void test_allocate() { + { // allocate(size) + async_resource input{42}; + cuda::mr::async_resource_ref<> ref{input}; + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0)); + + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0); + assert(input._val == expected_after_deallocate); + } + + { // allocate(size, alignment) + async_resource input{42}; + cuda::mr::async_resource_ref<> ref{input}; + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0, 0)); + + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0, 0); + assert(input._val == expected_after_deallocate); + } +} + +void test_allocate_async() { + { // allocate(size) + async_resource input{42}; + cuda::mr::async_resource_ref<> ref{input}; + + // Ensure that we properly pass on the allocate function + assert(input.allocate_async(0, 0, {}) == ref.allocate_async(0, {})); + + int expected_after_deallocate = 1337; + ref.deallocate_async(static_cast(&expected_after_deallocate), 0, {}); + assert(input._val == expected_after_deallocate); + } + + { // allocate(size, alignment) + async_resource input{42}; + cuda::mr::async_resource_ref<> ref{input}; + + // Ensure that we properly pass on the allocate function + assert(input.allocate_async(0, 0, {}) == ref.allocate_async(0, 0, {})); + + int expected_after_deallocate = 1337; + ref.deallocate_async(static_cast(&expected_after_deallocate), 0, 0, + {}); + assert(input._val == expected_after_deallocate); + } +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_allocate(); + test_allocate_async(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp new file mode 100644 index 0000000000..e3dd1feb6f --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -0,0 +1,115 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource_ref construction +#include + +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct async_resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + + void deallocate(void* ptr, std::size_t, std::size_t) {} + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const async_resource& other) const { + return _val == other._val; + } + bool operator!=(const async_resource& other) const { + return _val != other._val; + } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const async_resource& res, Property) noexcept { + return res._val; + } +}; + +namespace constructible { +using ref = cuda::mr::async_resource_ref, + property_with_value, + property_without_value >; + +using matching_properties = async_resource, + property_without_value, + property_with_value >; + +using missing_stateful_property = + async_resource, + property_without_value >; +using missing_stateless_property = + async_resource, property_with_value >; + +using cuda::std::is_constructible; +static_assert(is_constructible::value, ""); +static_assert(!is_constructible::value, ""); +static_assert(!is_constructible::value, ""); + +static_assert(is_constructible::value, ""); + +// Ensure we require a mutable valid reference and do not bind against rvalues +static_assert(!is_constructible::value, ""); +static_assert(!is_constructible::value, ""); + +static_assert(cuda::std::is_copy_constructible::value, ""); +static_assert(cuda::std::is_move_constructible::value, ""); +} // namespace constructible + +namespace assignable { +using ref = cuda::mr::async_resource_ref, + property_with_value, + property_without_value >; + +using res = + async_resource, property_with_value, + property_without_value >; + +using other_res = + async_resource, property_with_value, + property_with_value, + property_without_value >; + +using cuda::std::is_assignable; +static_assert(cuda::std::is_assignable::value, ""); +static_assert(cuda::std::is_assignable::value, ""); + +static_assert(cuda::std::is_copy_assignable::value, ""); +static_assert(cuda::std::is_move_assignable::value, ""); +} // namespace assignable + +int main(int, char**) { return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp new file mode 100644 index 0000000000..392272e6b3 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp @@ -0,0 +1,123 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref properties +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +struct Fake_alloc_base { + void* object = nullptr; + const cuda::mr::_Async_alloc_vtable* static_vtable = nullptr; +}; + +template +struct resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const resource& other) const { return _val == other._val; } + bool operator!=(const resource& other) const { return _val != other._val; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +template +void test_conversion_from_async_resource_ref() { + resource input{42}; + cuda::mr::async_resource_ref ref_input{input}; + + { // lvalue + cuda::mr::async_resource_ref ref{ref_input}; + + // Ensure that we properly "punch through" the resource ref + const auto fake_orig = reinterpret_cast(&ref_input); + const auto fake_conv = reinterpret_cast(&ref); + assert(fake_orig->object == fake_conv->object); + assert(fake_orig->static_vtable == fake_conv->static_vtable); + + // Ensure that we properly pass on the allocate function + assert(input.allocate_async(0, 0, {}) == ref.allocate_async(0, 0, {})); + + // Ensure we are deallocating properly + int expected_after_deallocate = 1337; + ref.deallocate_async(static_cast(&expected_after_deallocate), 0, 0, + {}); + assert(input._val == expected_after_deallocate); + } + + { // prvalue + cuda::mr::async_resource_ref ref{ + cuda::mr::async_resource_ref{input}}; + + // Ensure that we properly "punch through" the resource ref + const auto fake_orig = reinterpret_cast(&ref_input); + const auto fake_conv = reinterpret_cast(&ref); + assert(fake_orig->object == fake_conv->object); + assert(fake_orig->static_vtable == fake_conv->static_vtable); + + // Ensure that we properly pass on the allocate function + assert(input.allocate_async(0, 0, {}) == ref.allocate_async(0, 0, {})); + + // Ensure we are deallocating properly + int expected_after_deallocate = 1337; + ref.deallocate_async(static_cast(&expected_after_deallocate), 0, 0, + {}); + assert(input._val == expected_after_deallocate); + } +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_conversion_from_async_resource_ref, + property_with_value >(); + test_conversion_from_async_resource_ref, + property_without_value >(); + test_conversion_from_async_resource_ref, + property_without_value >(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp new file mode 100644 index 0000000000..b8d9ac5fc9 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -0,0 +1,92 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource_ref equality +#include +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct async_resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + + void deallocate(void* ptr, std::size_t, std::size_t) {} + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const async_resource& other) const { + return _val == other._val; + } + bool operator!=(const async_resource& other) const { + return _val != other._val; + } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const async_resource& res, Property) noexcept { + return res._val; + } +}; + +using ref = cuda::mr::async_resource_ref, + property_with_value, + property_without_value >; +using different_properties = + cuda::mr::async_resource_ref, + property_with_value, + property_without_value >; + +using res = + async_resource, property_with_value, + property_without_value >; + +void test_equality() { + res input{42}; + res with_equal_value{42}; + res with_different_value{1337}; + + // Requires matching properties + assert(ref{input} == different_properties{with_equal_value}); + assert(ref{input} != different_properties{with_different_value}); +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_equality(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp new file mode 100644 index 0000000000..8991aeb2da --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -0,0 +1,111 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource_ref equality +#include +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct async_resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + + void deallocate(void* ptr, std::size_t, std::size_t) {} + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const async_resource& other) const { + return _val == other._val; + } + bool operator!=(const async_resource& other) const { + return _val != other._val; + } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const async_resource& res, Property) noexcept { + return res._val; + } +}; + +using ref = cuda::mr::async_resource_ref, + property_with_value, + property_without_value >; + +using pertubed_properties = + cuda::mr::async_resource_ref, + property_with_value, + property_without_value >; + +using res = + async_resource, property_with_value, + property_without_value >; +using other_res = + async_resource, property_with_value, + property_without_value >; + +void test_equality() { + res input{42}; + res with_equal_value{42}; + res with_different_value{1337}; + + assert(input == with_equal_value); + assert(input != with_different_value); + + assert(ref{input} == ref{with_equal_value}); + assert(ref{input} != ref{with_different_value}); + + // Should ignore pertubed properties + assert(ref{input} == pertubed_properties{with_equal_value}); + assert(ref{input} != pertubed_properties{with_different_value}); + + // Should reject different resources + other_res other_with_matching_value{42}; + other_res other_with_different_value{1337}; + assert(ref{input} != ref{other_with_matching_value}); + assert(ref{input} != ref{other_with_different_value}); + + assert(ref{input} != pertubed_properties{other_with_matching_value}); + assert(ref{input} != pertubed_properties{other_with_matching_value}); +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_equality(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp new file mode 100644 index 0000000000..6d3966baec --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp @@ -0,0 +1,171 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref properties +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct async_resource_base { + virtual void* allocate(std::size_t, std::size_t) = 0; + + virtual void deallocate(void* ptr, std::size_t, std::size_t) = 0; + + virtual void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) = 0; + + virtual void deallocate_async(void* ptr, std::size_t, std::size_t, + cuda::stream_ref) = 0; + + bool operator==(const async_resource_base& other) const { return true; } + bool operator!=(const async_resource_base& other) const { return false; } + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource_base&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const async_resource_base& res, Property) noexcept { + return 42; + } +}; + +template +struct async_resource_derived_first + : public async_resource_base { + using super_t = async_resource_base; + using super_t::operator==; + + async_resource_derived_first(const int val) : _val(val) {} + + void* allocate(std::size_t, std::size_t) override { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) override {} + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) override { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, + cuda::stream_ref) override {} + + int _val = 0; +}; +static_assert(cuda::mr::async_resource >); + +struct some_data { + int _val; +}; + +template +struct async_resource_derived_second + : public async_resource_base { + using super_t = async_resource_base; + using super_t::operator==; + + async_resource_derived_second(some_data* val) : _val(val) {} + + void* allocate(std::size_t, std::size_t) override { return &_val->_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) override {} + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) override { + return &_val->_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, + cuda::stream_ref) override {} + + some_data* _val = 0; +}; + +template +void test_async_resource_ref() { + some_data input{1337}; + async_resource_derived_first first{42}; + async_resource_derived_second second{&input}; + + cuda::mr::async_resource_ref ref_first{first}; + cuda::mr::async_resource_ref ref_second{second}; + + // Ensure that we properly pass on the allocate function + assert(ref_first.allocate_async(0, 0, {}) == first.allocate_async(0, 0, {})); + assert(ref_second.allocate_async(0, 0, {}) == + second.allocate_async(0, 0, {})); + + // Ensure that assignment still works + ref_second = ref_first; + assert(ref_second.allocate_async(0, 0, {}) == first.allocate_async(0, 0, {})); +} + +template +cuda::mr::async_resource_ref +indirection(async_resource_base* res) { + return {*res}; +} + +template +void test_async_resource_ref_from_pointer() { + some_data input{1337}; + async_resource_derived_first first{42}; + async_resource_derived_second second{&input}; + + cuda::mr::async_resource_ref ref_first = indirection(&first); + cuda::mr::async_resource_ref ref_second = indirection(&second); + + // Ensure that we properly pass on the allocate function + assert(ref_first.allocate_async(0, 0, {}) == first.allocate_async(0, 0, {})); + assert(ref_second.allocate_async(0, 0, {}) == + second.allocate_async(0, 0, {})); + + // Ensure that assignment still works + ref_second = ref_first; + assert(ref_second.allocate_async(0, 0, {}) == first.allocate_async(0, 0, {})); +} + +// clang complains about pure virtual functions being called, so ensure that we properly crash if so +extern "C" void __cxa_pure_virtual() { + while (1) + ; +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + // Test some basic combinations of properties w/o state + test_async_resource_ref, + property_with_value >(); + test_async_resource_ref, + property_without_value >(); + test_async_resource_ref, + property_without_value >(); + + test_async_resource_ref_from_pointer, + property_with_value >(); + test_async_resource_ref_from_pointer, + property_without_value >(); + test_async_resource_ref_from_pointer, + property_without_value >(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp new file mode 100644 index 0000000000..9d637db781 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -0,0 +1,190 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::async_resource_ref properties +#include +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +namespace properties_test { +static_assert(cuda::mr::property_with_value >, ""); +static_assert( + cuda::mr::property_with_value >, ""); + +static_assert(!cuda::mr::property_with_value >, ""); +static_assert( + !cuda::mr::property_with_value >, + ""); +} // namespace properties_test + +namespace resource_test { + +template +struct async_resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + + void deallocate(void* ptr, std::size_t, std::size_t) {} + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return nullptr; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + } + + bool operator==(const async_resource& other) const { return true; } + bool operator!=(const async_resource& other) const { return false; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const async_resource& res, Property) noexcept { + return res._val; + } +}; + +// Ensure we have the right size +static_assert(sizeof(cuda::mr::async_resource_ref, + property_with_value >) == + (4 * sizeof(void*))); +static_assert( + sizeof(cuda::mr::async_resource_ref, + property_without_value >) == + (3 * sizeof(void*))); +static_assert(sizeof(cuda::mr::async_resource_ref, + property_with_value >) == + (3 * sizeof(void*))); +static_assert( + sizeof(cuda::mr::async_resource_ref, + property_without_value >) == + (2 * sizeof(void*))); + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires !cuda::mr::property_with_value) // + int InvokeIfWithValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires cuda::mr::property_with_value) // + typename Property::value_type InvokeIfWithValue(const Ref& ref) { + return get_property(ref, Property{}); +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires cuda::mr::property_with_value) // + int InvokeIfWithoutValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires !cuda::mr::property_with_value) // + int InvokeIfWithoutValue(const Ref& ref) { + get_property(ref, Property{}); + return 1; +} + +template +void test_async_resource_ref() { + constexpr int expected_initially = 42; + async_resource input{expected_initially}; + cuda::mr::async_resource_ref ref{input}; + + // Check all the potentially stateful properties + const int properties_with_value[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value[] = { + ((cuda::mr::property_with_value) ? expected_initially + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value[i] == expected_with_value[i]); + } + + const int properties_without_value[] = { + InvokeIfWithoutValue(ref)...}; + const int expected_without_value[] = { + ((cuda::mr::property_with_value) ? -1 : 1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_without_value[i] == expected_without_value[i]); + } + + constexpr int expected_after_change = 1337; + input._val = expected_after_change; + + // Check whether we truly get the right value + const int properties_with_value2[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value2[] = { + ((cuda::mr::property_with_value) ? expected_after_change + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value2[i] == expected_with_value2[i]); + } +} + +void test_property_forwarding() { + using res = + async_resource, property_with_value >; + using ref = cuda::mr::async_resource_ref >; + + static_assert(cuda::mr::async_resource_with, + property_with_value >); + static_assert(!cuda::mr::async_resource_with, + property_with_value >); + + static_assert( + cuda::mr::async_resource_with >); +} + +void test_async_resource_ref() { + // Test some basic combinations of properties w/o state + test_async_resource_ref, + property_with_value >(); + test_async_resource_ref, + property_without_value >(); + test_async_resource_ref, + property_without_value >(); + + // Test duplicated properties + test_async_resource_ref, property_with_value, + property_with_value >(); + + test_async_resource_ref, + property_without_value, + property_without_value >(); + + // Ensure we only forward requested properties + test_property_forwarding(); +} +} // namespace resource_test + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + resource_test::test_async_resource_ref(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp new file mode 100644 index 0000000000..27a9439357 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp @@ -0,0 +1,63 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref properties +#include + +#include +#include + +struct resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const resource& other) const { return _val == other._val; } + bool operator!=(const resource& other) const { return _val != other._val; } + + int _val = 0; +}; + +void test_allocate() { + { // allocate(size) + resource input{42}; + cuda::mr::resource_ref<> ref{input}; + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0)); + + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0); + assert(input._val == expected_after_deallocate); + } + + { // allocate(size, alignment) + resource input{42}; + cuda::mr::resource_ref<> ref{input}; + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0, 0)); + + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0, 0); + assert(input._val == expected_after_deallocate); + } +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_allocate(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp new file mode 100644 index 0000000000..774973a397 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp @@ -0,0 +1,102 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref construction +#include + +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const resource& other) const { return _val == other._val; } + bool operator!=(const resource& other) const { return _val != other._val; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +namespace constructible { +using ref = cuda::mr::resource_ref, + property_with_value, + property_without_value >; + +using matching_properties = + resource, property_without_value, + property_with_value >; + +using missing_stateful_property = + resource, property_without_value >; +using missing_stateless_property = + resource, property_with_value >; + +using cuda::std::is_constructible; +static_assert(is_constructible::value, ""); +static_assert(!is_constructible::value, ""); +static_assert(!is_constructible::value, ""); + +static_assert(is_constructible::value, ""); + +// Ensure we require a mutable valid reference and do not bind against rvalues +static_assert(!is_constructible::value, ""); +static_assert(!is_constructible::value, ""); + +static_assert(cuda::std::is_copy_constructible::value, ""); +static_assert(cuda::std::is_move_constructible::value, ""); +} // namespace constructible + +namespace assignable { +using ref = cuda::mr::resource_ref, + property_with_value, + property_without_value >; + +using res = resource, property_with_value, + property_without_value >; + +using other_res = + resource, property_with_value, + property_with_value, property_without_value >; + +using cuda::std::is_assignable; +static_assert(cuda::std::is_assignable::value, ""); +static_assert(cuda::std::is_assignable::value, ""); + +static_assert(cuda::std::is_copy_assignable::value, ""); +static_assert(cuda::std::is_move_assignable::value, ""); +} // namespace assignable + +int main(int, char**) { return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp new file mode 100644 index 0000000000..eb633032e2 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp @@ -0,0 +1,171 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref properties +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +struct Fake_alloc_base { + void* object = nullptr; + const cuda::mr::_Async_alloc_vtable* static_vtable = nullptr; +}; + +template +struct resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) { + return &_val; + } + + void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const resource& other) const { return _val == other._val; } + bool operator!=(const resource& other) const { return _val != other._val; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +template +void test_conversion_from_resource_ref() { + resource input{42}; + cuda::mr::resource_ref ref_input{input}; + + { // lvalue + cuda::mr::resource_ref ref{ref_input}; + + // Ensure that we properly "punch through" the resource ref + const auto fake_orig = *reinterpret_cast(&ref_input); + const auto fake_conv = *reinterpret_cast(&ref); + assert(fake_orig.object == fake_conv.object); + assert(fake_orig.static_vtable == fake_conv.static_vtable); + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0, 0)); + + // Ensure we are deallocating properly + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0, 0); + assert(input._val == expected_after_deallocate); + } + + { // prvalue + cuda::mr::resource_ref ref{ + cuda::mr::resource_ref{input}}; + + // Ensure that we properly "punch through" the resource ref + const auto fake_orig = *reinterpret_cast(&ref_input); + const auto fake_conv = *reinterpret_cast(&ref); + assert(fake_orig.object == fake_conv.object); + assert(fake_orig.static_vtable == fake_conv.static_vtable); + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0, 0)); + + // Ensure we are deallocating properly + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0, 0); + assert(input._val == expected_after_deallocate); + } +} + +template +void test_conversion_from_async_resource_ref() { + resource input{42}; + cuda::mr::async_resource_ref ref_input{input}; + + { // lvalue + cuda::mr::resource_ref ref{ref_input}; + + // Ensure that we properly "punch through" the resource ref + const auto fake_orig = reinterpret_cast(&ref_input); + const auto fake_conv = reinterpret_cast(&ref); + assert(fake_orig->object == fake_conv->object); + assert(fake_orig->static_vtable == fake_conv->static_vtable); + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0, 0)); + + // Ensure we are deallocating properly + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0, 0); + assert(input._val == expected_after_deallocate); + } + + { // prvalue + cuda::mr::resource_ref ref{ + cuda::mr::async_resource_ref{input}}; + + // Ensure that we properly "punch through" the resource ref + const auto fake_orig = reinterpret_cast(&ref_input); + const auto fake_conv = reinterpret_cast(&ref); + assert(fake_orig->object == fake_conv->object); + assert(fake_orig->static_vtable == fake_conv->static_vtable); + + // Ensure that we properly pass on the allocate function + assert(input.allocate(0, 0) == ref.allocate(0, 0)); + + // Ensure we are deallocating properly + int expected_after_deallocate = 1337; + ref.deallocate(static_cast(&expected_after_deallocate), 0, 0); + assert(input._val == expected_after_deallocate); + } +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_conversion_from_resource_ref, + property_with_value >(); + test_conversion_from_resource_ref, + property_without_value >(); + test_conversion_from_resource_ref, + property_without_value >(); + + test_conversion_from_async_resource_ref, + property_with_value >(); + test_conversion_from_async_resource_ref, + property_without_value >(); + test_conversion_from_async_resource_ref, + property_without_value >(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp new file mode 100644 index 0000000000..af9be2f75c --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -0,0 +1,80 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref equality +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const resource& other) const { return _val == other._val; } + bool operator!=(const resource& other) const { return _val != other._val; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +using ref = cuda::mr::resource_ref, + property_with_value, + property_without_value >; + +using different_properties = + cuda::mr::resource_ref, property_with_value, + property_without_value >; + +using res = resource, property_with_value, + property_without_value >; + +void test_equality() { + res input{42}; + res with_equal_value{42}; + res with_different_value{1337}; + + // Requires matching properties + assert(ref{input} == different_properties{with_equal_value}); + assert(ref{input} != different_properties{with_different_value}); +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_equality(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp new file mode 100644 index 0000000000..f32f1abbb3 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -0,0 +1,99 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref equality +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct resource { + void* allocate(std::size_t, std::size_t) { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + bool operator==(const resource& other) const { return _val == other._val; } + bool operator!=(const resource& other) const { return _val != other._val; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +using ref = cuda::mr::resource_ref, + property_with_value, + property_without_value >; + +using pertubed_properties = + cuda::mr::resource_ref, + property_with_value, + property_without_value >; + +using res = resource, property_with_value, + property_without_value >; +using other_res = + resource, property_with_value, + property_without_value >; + +void test_equality() { + res input{42}; + res with_equal_value{42}; + res with_different_value{1337}; + + assert(input == with_equal_value); + assert(input != with_different_value); + + assert(ref{input} == ref{with_equal_value}); + assert(ref{input} != ref{with_different_value}); + + // Should ignore pertubed properties + assert(ref{input} == pertubed_properties{with_equal_value}); + assert(ref{input} != pertubed_properties{with_different_value}); + + // Should reject different resources + other_res other_with_matching_value{42}; + other_res other_with_different_value{1337}; + assert(ref{input} != ref{other_with_matching_value}); + assert(ref{input} != ref{other_with_different_value}); + + assert(ref{input} != pertubed_properties{other_with_matching_value}); + assert(ref{input} != pertubed_properties{other_with_matching_value}); +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + test_equality(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp new file mode 100644 index 0000000000..e398ad313a --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp @@ -0,0 +1,146 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref properties +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct resource_base { + virtual void* allocate(std::size_t, std::size_t) = 0; + + virtual void deallocate(void* ptr, std::size_t, std::size_t) = 0; + + bool operator==(const resource_base& other) const { return true; } + bool operator!=(const resource_base& other) const { return false; } + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource_base&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource_base& res, Property) noexcept { + return 42; + } +}; + +template +struct resource_derived_first : public resource_base { + using super_t = resource_base; + using super_t::operator==; + + resource_derived_first(const int val) : _val(val) {} + + void* allocate(std::size_t, std::size_t) override { return &_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) override {} + + int _val = 0; +}; +static_assert(cuda::mr::resource >); + +struct some_data { + int _val; +}; + +template +struct resource_derived_second : public resource_base { + using super_t = resource_base; + using super_t::operator==; + + resource_derived_second(some_data* val) : _val(val) {} + + void* allocate(std::size_t, std::size_t) override { return &_val->_val; } + + void deallocate(void* ptr, std::size_t, std::size_t) override {} + + some_data* _val = 0; +}; + +template +void test_resource_ref() { + some_data input{1337}; + resource_derived_first first{42}; + resource_derived_second second{&input}; + + cuda::mr::resource_ref ref_first{first}; + cuda::mr::resource_ref ref_second{second}; + + // Ensure that we properly pass on the allocate function + assert(ref_first.allocate(0, 0) == first.allocate(0, 0)); + assert(ref_second.allocate(0, 0) == second.allocate(0, 0)); + + // Ensure that assignment still works + ref_second = ref_first; + assert(ref_second.allocate(0, 0) == first.allocate(0, 0)); +} + +template +cuda::mr::resource_ref +indirection(resource_base* res) { + return {*res}; +} + +template +void test_resource_ref_from_pointer() { + some_data input{1337}; + resource_derived_first first{42}; + resource_derived_second second{&input}; + + cuda::mr::resource_ref ref_first = indirection(&first); + cuda::mr::resource_ref ref_second = indirection(&second); + + // Ensure that we properly pass on the allocate function + assert(ref_first.allocate(0, 0) == first.allocate(0, 0)); + assert(ref_second.allocate(0, 0) == second.allocate(0, 0)); + + // Ensure that assignment still works + ref_second = ref_first; + assert(ref_second.allocate(0, 0) == first.allocate(0, 0)); +} + +// clang complains about pure virtual functions being called, so ensure that we properly crash if so +extern "C" void __cxa_pure_virtual() { + while (1) + ; +} + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + // Test some basic combinations of properties w/o state + test_resource_ref, property_with_value >(); + test_resource_ref, property_without_value >(); + test_resource_ref, + property_without_value >(); + + test_resource_ref_from_pointer, + property_with_value >(); + test_resource_ref_from_pointer, + property_without_value >(); + test_resource_ref_from_pointer, + property_without_value >(); +#endif + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp new file mode 100644 index 0000000000..759b1ed1c8 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -0,0 +1,175 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::mr::resource_ref properties +#include + +#include +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +namespace properties_test { +static_assert(cuda::mr::property_with_value >, ""); +static_assert( + cuda::mr::property_with_value >, ""); + +static_assert(!cuda::mr::property_with_value >, ""); +static_assert( + !cuda::mr::property_with_value >, + ""); +} // namespace properties_test + +namespace resource_test { +template +struct resource { + void* allocate(std::size_t, std::size_t) { return nullptr; } + + void deallocate(void* ptr, std::size_t, std::size_t) {} + + bool operator==(const resource& other) const { return true; } + bool operator!=(const resource& other) const { return false; } + + int _val = 0; + + _LIBCUDACXX_TEMPLATE(class Property) + (requires !cuda::mr::property_with_value && + _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +// Ensure we have the right size +static_assert(sizeof(cuda::mr::resource_ref, + property_with_value >) == + (4 * sizeof(void*))); +static_assert(sizeof(cuda::mr::resource_ref, + property_without_value >) == + (3 * sizeof(void*))); +static_assert(sizeof(cuda::mr::resource_ref, + property_with_value >) == + (3 * sizeof(void*))); +static_assert(sizeof(cuda::mr::resource_ref, + property_without_value >) == + (2 * sizeof(void*))); + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires !cuda::mr::property_with_value) // + int InvokeIfWithValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires cuda::mr::property_with_value) // + typename Property::value_type InvokeIfWithValue(const Ref& ref) { + return get_property(ref, Property{}); +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires cuda::mr::property_with_value) // + int InvokeIfWithoutValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires !cuda::mr::property_with_value) // + int InvokeIfWithoutValue(const Ref& ref) { + get_property(ref, Property{}); + return 1; +} + +template +void test_resource_ref() { + constexpr int expected_initially = 42; + resource input{expected_initially}; + cuda::mr::resource_ref ref{input}; + + // Check all the potentially stateful properties + const int properties_with_value[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value[] = { + ((cuda::mr::property_with_value) ? expected_initially + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value[i] == expected_with_value[i]); + } + + const int properties_without_value[] = { + InvokeIfWithoutValue(ref)...}; + const int expected_without_value[] = { + ((cuda::mr::property_with_value) ? -1 : 1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_without_value[i] == expected_without_value[i]); + } + + constexpr int expected_after_change = 1337; + input._val = expected_after_change; + + // Check whether we truly get the right value + const int properties_with_value2[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value2[] = { + ((cuda::mr::property_with_value) ? expected_after_change + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value2[i] == expected_with_value2[i]); + } +} + +void test_property_forwarding() { + using res = resource, property_with_value >; + using ref = cuda::mr::resource_ref >; + + static_assert(cuda::mr::resource_with, + property_with_value >); + static_assert(!cuda::mr::resource_with, + property_with_value >); + + static_assert(cuda::mr::resource_with >); +} + +void test_resource_ref() { + // Test some basic combinations of properties w/o state + test_resource_ref, property_with_value >(); + test_resource_ref, property_without_value >(); + test_resource_ref, + property_without_value >(); + + // Test duplicated properties + test_resource_ref, property_with_value, + property_with_value >(); + + test_resource_ref, property_without_value, + property_without_value >(); + + // Ensure we only forward requested properties + test_property_forwarding(); +} + +} // namespace resource_test + +int main(int, char**) { +#ifndef __CUDA_ARCH__ + resource_test::test_resource_ref(); +#endif + return 0; +} diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index 5e86efabc1..9faf03f233 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -98,86 +98,65 @@ namespace mr /// \concept resource /// \brief The \c resource concept template -_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_member_allocate_, - requires(_Resource& __res, size_t __bytes, size_t __alignment) // - (_CUDA_VSTD::_Returns_exactly(__res.allocate(__bytes, __alignment)))); -template -_LIBCUDACXX_CONCEPT _Has_member_allocate = _LIBCUDACXX_FRAGMENT(_Has_member_allocate_, _Resource); - -template -_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_member_deallocate_, - requires(_Resource& __res, - void* __ptr, - size_t __bytes, - size_t __alignment) // - (__res.deallocate(__ptr, __bytes, __alignment))); -template -_LIBCUDACXX_CONCEPT _Has_member_deallocate = _LIBCUDACXX_FRAGMENT(_Has_member_deallocate_, _Resource); +_LIBCUDACXX_CONCEPT_FRAGMENT( + __resource_, + requires(_Resource& __res, void* __ptr, size_t __bytes, size_t __alignment) ( + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::equality_comparable<_Resource>) + )); template -_LIBCUDACXX_CONCEPT resource = - _Has_member_allocate<_Resource>&& _Has_member_deallocate<_Resource>&& _CUDA_VSTD::equality_comparable<_Resource>; +_LIBCUDACXX_CONCEPT resource = _LIBCUDACXX_FRAGMENT(__resource_, _Resource); /// \concept async_resource /// \brief The \c async_resource concept template _LIBCUDACXX_CONCEPT_FRAGMENT( - _Has_member_allocate_async_, - requires(_Resource& __res, - size_t __bytes, - size_t __alignment, - cuda::stream_ref __stream) // - (_CUDA_VSTD::_Returns_exactly(__res.allocate_async(__bytes, __alignment, __stream)))); -template -_LIBCUDACXX_CONCEPT _Has_member_allocate_async = _LIBCUDACXX_FRAGMENT(_Has_member_allocate_async_, _Resource); - -template -_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_member_deallocate_async_, - requires(_Resource& __res, - void* __ptr, - size_t __bytes, - size_t __alignment, - cuda::stream_ref __stream) // - (__res.deallocate_async(__ptr, __bytes, __alignment, __stream))); -template -_LIBCUDACXX_CONCEPT _Has_member_deallocate_async = _LIBCUDACXX_FRAGMENT(_Has_member_deallocate_async_, _Resource); + __async_resource_, + requires(_Resource& __res, void* __ptr, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) ( + requires(resource<_Resource>), + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::equality_comparable<_Resource>) + )); -/// \concept async_resource -/// \brief The \c async_resource concept template -_LIBCUDACXX_CONCEPT async_resource = - resource<_Resource>&& _Has_member_allocate_async<_Resource>&& _Has_member_deallocate_async<_Resource>; +_LIBCUDACXX_CONCEPT async_resource = _LIBCUDACXX_FRAGMENT(__async_resource_, _Resource); /// \concept has_property /// \brief The \c has_property concept template -_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_property_, - requires(const _Resource& __res) // - (get_property(__res, _Property{}))); +_LIBCUDACXX_CONCEPT_FRAGMENT( + __has_property_, + requires(const _Resource& __res)( + get_property(__res, _Property{}) + )); template -_LIBCUDACXX_CONCEPT has_property = _LIBCUDACXX_FRAGMENT(_Has_property_, _Resource, _Property); - -template -using _Has_value_types = typename T::value_type; +_LIBCUDACXX_CONCEPT has_property = _LIBCUDACXX_FRAGMENT(__has_property_, _Resource, _Property); /// \concept property_with_value /// \brief The \c property_with_value concept template -_LIBCUDACXX_CONCEPT_FRAGMENT(_Property_with_value_, - requires() // - (typename(_Has_value_types<_Property>))); +_LIBCUDACXX_CONCEPT_FRAGMENT( + __property_with_value_, + requires()( + typename(typename _Property::value_type) + )); template -_LIBCUDACXX_CONCEPT property_with_value = _LIBCUDACXX_FRAGMENT(_Property_with_value_, _Property); +_LIBCUDACXX_CONCEPT property_with_value = _LIBCUDACXX_FRAGMENT(__property_with_value_, _Property); /// \concept has_property_with /// \brief The \c has_property_with concept template -_LIBCUDACXX_CONCEPT_FRAGMENT(_Has_property_with_, - requires(const _Resource& __res) // - (_CUDA_VSTD::_Returns_exactly<_Return>(get_property(__res, _Property{})))); +_LIBCUDACXX_CONCEPT_FRAGMENT( + __has_property_with_, + requires(const _Resource& __res)( + requires(property_with_value<_Property>), + requires(_CUDA_VSTD::same_as<_Return, decltype(get_property(__res, _Property{}))>) + )); template -_LIBCUDACXX_CONCEPT has_property_with = property_with_value<_Property>&& - _LIBCUDACXX_FRAGMENT(_Has_property_with_, _Resource, _Property, _Return); +_LIBCUDACXX_CONCEPT has_property_with = _LIBCUDACXX_FRAGMENT(__has_property_with_, _Resource, _Property, _Return); /// \concept resource_with /// \brief The \c resource_with concept @@ -186,20 +165,373 @@ template _LIBCUDACXX_CONCEPT resource_with = resource<_Resource>&& _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant>...>; #else -_LIBCUDACXX_CONCEPT resource_with = resource<_Resource> && (has_property<_Resource, _Properties> && ...); +_LIBCUDACXX_CONCEPT resource_with = resource<_Resource> && (has_property<_Resource, _Properties> && ...); #endif /// \concept async_resource_with /// \brief The \c async_resource_with concept template #if _LIBCUDACXX_STD_VER < 17 -_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource>&& +_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant>...>; #else -_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> // - && (has_property<_Resource, _Properties> && ...); +_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> && + (has_property<_Resource, _Properties> && ...); #endif +/////////////////////////////////////////////////////////////////////////////// +/// class resource_ref +/// class async_resource_ref +enum class _AllocType +{ + _Default, + _Async, +}; + +struct _Alloc_vtable +{ + using _AllocFn = void* (*)(void*, size_t, size_t); + using _DeallocFn = void (*)(void*, void*, size_t, size_t); + using _EqualFn = bool (*)(void*, void*); + + _AllocFn __alloc_fn; + _DeallocFn __dealloc_fn; + _EqualFn __equal_fn; + + constexpr _Alloc_vtable(_AllocFn __alloc_fn_, _DeallocFn __dealloc_fn_, _EqualFn __equal_fn_) noexcept + : __alloc_fn(__alloc_fn_) + , __dealloc_fn(__dealloc_fn_) + , __equal_fn(__equal_fn_) + {} +}; + +struct _Async_alloc_vtable : public _Alloc_vtable +{ + using _AsyncAllocFn = void* (*)(void*, size_t, size_t, cuda::stream_ref); + using _AsyncDeallocFn = void (*)(void*, void*, size_t, size_t, cuda::stream_ref); + + _AsyncAllocFn __async_alloc_fn; + _AsyncDeallocFn __async_dealloc_fn; + + constexpr _Async_alloc_vtable(_Alloc_vtable::_AllocFn __alloc_fn_, + _Alloc_vtable::_DeallocFn __dealloc_fn_, + _Alloc_vtable::_EqualFn __equal_fn_, + _AsyncAllocFn __async_alloc_fn_, + _AsyncDeallocFn __async_dealloc_fn_) noexcept + : _Alloc_vtable(__alloc_fn_, __dealloc_fn_, __equal_fn_) + , __async_alloc_fn(__async_alloc_fn_) + , __async_dealloc_fn(__async_dealloc_fn_) + {} +}; + +// clang-format off +struct _Resource_vtable_builder +{ + template + static typename _Property::value_type _Get_property(void* __res) noexcept { + return get_property(*static_cast(__res), _Property{}); + } + + template + static void* _Alloc(void* __object, size_t __bytes, size_t __alignment) { + return static_cast<_Resource *>(__object)->allocate(__bytes, __alignment); + } + + template + static void _Dealloc(void* __object, void* __ptr, size_t __bytes, size_t __alignment) { + return static_cast<_Resource *>(__object)->deallocate(__ptr, __bytes, __alignment); + } + + template + static void* _Alloc_async(void* __object, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) { + return static_cast<_Resource *>(__object)->allocate_async(__bytes, __alignment, __stream); + } + + template + static void _Dealloc_async(void* __object, void* __ptr, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) { + return static_cast<_Resource *>(__object)->deallocate_async(__ptr, __bytes, __alignment, __stream); + } + + template + static bool _Equal(void* __left, void* __right) { + return *static_cast<_Resource *>(__left) == *static_cast<_Resource *>(__right); + } + + _LIBCUDACXX_TEMPLATE(class _Resource, _AllocType _Alloc_type) + (requires(_Alloc_type == _AllocType::_Default)) // + static constexpr _Alloc_vtable _Create() noexcept + { + return {&_Resource_vtable_builder::_Alloc<_Resource>, + &_Resource_vtable_builder::_Dealloc<_Resource>, + &_Resource_vtable_builder::_Equal<_Resource>}; + } + + _LIBCUDACXX_TEMPLATE(class _Resource, _AllocType _Alloc_type) + (requires(_Alloc_type == _AllocType::_Async)) // + static constexpr _Async_alloc_vtable _Create() noexcept + { + return {&_Resource_vtable_builder::_Alloc<_Resource>, + &_Resource_vtable_builder::_Dealloc<_Resource>, + &_Resource_vtable_builder::_Equal<_Resource>, + &_Resource_vtable_builder::_Alloc_async<_Resource>, + &_Resource_vtable_builder::_Dealloc_async<_Resource>}; + } +}; +// clang-format on + +template +struct _Property_vtable +{ + using _PropertyFn = typename _Property::value_type (*)(void*); + _PropertyFn __property_fn = nullptr; + + constexpr _Property_vtable(_PropertyFn __property_fn_) noexcept + : __property_fn(__property_fn_) + {} +}; + +template <_AllocType _Alloc_type, class... _Properties> // +class basic_resource_ref; + +template +struct _Resource_vtable : public _Property_vtable<_Properties>... +{ + template + constexpr _Resource_vtable(_PropertyFns... __property_fn_) noexcept + : _Property_vtable<_Properties>(__property_fn_)... + {} + + template <_AllocType _Alloc_type, class... _OtherProperties> + constexpr _Resource_vtable(basic_resource_ref<_Alloc_type, _OtherProperties...>& __ref) noexcept + : _Property_vtable<_Properties>(__ref._Property_vtable<_Properties>::__property_fn)... + {} + + template + static constexpr _Resource_vtable _Create() noexcept + { + return {&_Resource_vtable_builder::_Get_property<_Resource, _Properties>...}; + } +}; + +template +struct _Filtered; + +template +struct _Property_filter +{ + template + using _Filtered_properties = + typename _Filtered<_Properties...>::_Filtered_vtable::template _Append_property<_Property>; +}; +template <> +struct _Property_filter +{ + template + using _Filtered_properties = typename _Filtered<_Properties...>::_Filtered_vtable; +}; + +template +struct _Filtered<_Property, _Properties...> +{ + using _Filtered_vtable = + typename _Property_filter && !_CUDA_VSTD::_One_of<_Property, _Properties...>>:: + template _Filtered_properties<_Property, _Properties...>; + + template + using _Append_property = _Filtered<_OtherPropery, _Property, _Properties...>; + + using _Vtable = _Resource_vtable<_Property, _Properties...>; +}; + +template <> +struct _Filtered<> +{ + using _Filtered_vtable = _Filtered<>; + + template + using _Append_property = _Filtered<_OtherPropery>; + + using _Vtable = _Resource_vtable<>; +}; + +template +using _Filtered_vtable = typename _Filtered<_Properties...>::_Filtered_vtable::_Vtable; + +template +struct _Alloc_base +{ + static_assert(cuda::std::is_base_of_v<_Alloc_vtable, _Vtable>); + + _Alloc_base(void* __object_, const _Vtable* __static_vtabl_) noexcept + : __object(__object_) + , __static_vtable(__static_vtabl_) + {} + + void* allocate(size_t __bytes, size_t __alignment = alignof(max_align_t)) + { + return __static_vtable->__alloc_fn(__object, __bytes, __alignment); + } + + void deallocate(void* _Ptr, size_t __bytes, size_t __alignment = alignof(max_align_t)) + { + __static_vtable->__dealloc_fn(__object, _Ptr, __bytes, __alignment); + } + +protected: + void* __object = nullptr; + const _Vtable* __static_vtable = nullptr; +}; + +template +struct _Async_alloc_base : public _Alloc_base<_Vtable> +{ + static_assert(cuda::std::is_base_of_v<_Async_alloc_vtable, _Vtable>); + + _Async_alloc_base(void* __object_, const _Vtable* __static_vtabl_) noexcept + : _Alloc_base<_Vtable>(__object_, __static_vtabl_) + {} + + void* allocate_async(size_t __bytes, size_t __alignment, cuda::stream_ref __stream) + { + return this->__static_vtable->__async_alloc_fn(this->__object, __bytes, __alignment, __stream); + } + + void* allocate_async(size_t __bytes, cuda::stream_ref __stream) + { + return this->__static_vtable->__async_alloc_fn(this->__object, __bytes, alignof(max_align_t), __stream); + } + + void deallocate_async(void* _Ptr, size_t __bytes, cuda::stream_ref __stream) + { + this->__static_vtable->__async_dealloc_fn(this->__object, _Ptr, __bytes, alignof(max_align_t), __stream); + } + + void deallocate_async(void* _Ptr, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) + { + this->__static_vtable->__async_dealloc_fn(this->__object, _Ptr, __bytes, __alignment, __stream); + } +}; + +template <_AllocType _Alloc_type> +using _Resource_ref_base = _CUDA_VSTD:: + _If<_Alloc_type == _AllocType::_Default, _Alloc_base<_Alloc_vtable>, _Async_alloc_base<_Async_alloc_vtable>>; + +template <_AllocType _Alloc_type> +using _Vtable_store = _CUDA_VSTD::_If<_Alloc_type == _AllocType::_Default, _Alloc_vtable, _Async_alloc_vtable>; + +template <_AllocType _Alloc_type, class _Resource> +_LIBCUDACXX_INLINE_VAR constexpr _Vtable_store<_Alloc_type> + __alloc_vtable = _Resource_vtable_builder::template _Create<_Resource, _Alloc_type>(); + +template +_LIBCUDACXX_INLINE_VAR constexpr bool _Is_basic_resource_ref = false; + +template <_AllocType _Alloc_type, class... _Properties> // +class basic_resource_ref + : public _Resource_ref_base<_Alloc_type> + , private _Filtered_vtable<_Properties...> +{ +private: + template <_AllocType, class...> + friend class basic_resource_ref; + + template + friend struct _Resource_vtable; + +public: + // clang-format off + _LIBCUDACXX_TEMPLATE(class _Resource) + (requires (!_Is_basic_resource_ref<_Resource> + && (((_Alloc_type == _AllocType::_Default) && resource_with<_Resource, _Properties...>) // + ||((_Alloc_type == _AllocType::_Async) && async_resource_with<_Resource, _Properties...>)))) // + basic_resource_ref(_Resource& __res) noexcept + : _Resource_ref_base<_Alloc_type>(&__res, &__alloc_vtable<_Alloc_type, _Resource>) + , _Filtered_vtable<_Properties...>(_Filtered_vtable<_Properties...>::template _Create<_Resource>()) + {} + + #if _LIBCUDACXX_STD_VER > 14 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #endif + basic_resource_ref( + basic_resource_ref<_Alloc_type, _OtherProperties...> __ref) noexcept + : _Resource_ref_base<_Alloc_type>(__ref.__object, __ref.__static_vtable) + , _Filtered_vtable<_Properties...>(__ref) + {} + + #if _LIBCUDACXX_STD_VER > 14 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (_Alloc_type == _AllocType::_Default) + && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (_Alloc_type == _AllocType::_Default) + && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #endif + basic_resource_ref( + basic_resource_ref<_AllocType::_Async, _OtherProperties...> __ref) noexcept + : _Resource_ref_base<_Alloc_type>(__ref.__object, __ref.__static_vtable) + , _Filtered_vtable<_Properties...>(__ref) + {} + + #if _LIBCUDACXX_STD_VER > 14 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires(sizeof...(_Properties) == sizeof...(_OtherProperties)) + && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #endif + bool operator==( + const basic_resource_ref<_Alloc_type, _OtherProperties...> &__right) const { + return (this->__static_vtable->__equal_fn == __right.__static_vtable->__equal_fn) // + && this->__static_vtable->__equal_fn(this->__object, __right.__object); + } + + #if _LIBCUDACXX_STD_VER > 14 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #endif + bool operator!=( + const basic_resource_ref<_Alloc_type, _OtherProperties...> &__right) const { + return !(*this == __right); + } + + _LIBCUDACXX_TEMPLATE(class _Property) + (requires !property_with_value<_Property> && _CUDA_VSTD::_One_of<_Property, _Properties...>) // + friend void get_property(const basic_resource_ref &, _Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class _Property) + (requires property_with_value<_Property> &&_CUDA_VSTD::_One_of<_Property, _Properties...>) // + friend typename _Property::value_type get_property( + const basic_resource_ref &__res, _Property) noexcept { + return __res._Property_vtable<_Property>::__property_fn(__res.__object); + } + // clang-format on +}; + +template <_AllocType _Alloc_type, class... _Properties> +_LIBCUDACXX_INLINE_VAR constexpr bool _Is_basic_resource_ref> = true; + +template // +using resource_ref = basic_resource_ref<_AllocType::_Default, _Properties...>; + +template // +using async_resource_ref = basic_resource_ref<_AllocType::_Async, _Properties...>; + } // namespace mr _LIBCUDACXX_END_NAMESPACE_CUDA #endif // _LIBCUDACXX_STD_VER > 11 From 8befe8da3d7661a0b18a657b25585fe2d6f0c82b Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 4/9] Add `device_accessible` and `host_accessible` properties --- include/cuda/memory_resource | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index 9faf03f233..45a6a1cfc8 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -532,6 +532,14 @@ using resource_ref = basic_resource_ref<_AllocType::_Default, _Properties...>; template // using async_resource_ref = basic_resource_ref<_AllocType::_Async, _Properties...>; +/// \struct device_accessible +/// \brief The \c device_accessible property signals that the allocated memory is device accessible +struct device_accessible{}; + +/// \struct host_accessible +/// \brief The \c host_accessible property signals that the allocated memory is host accessible +struct host_accessible{}; + } // namespace mr _LIBCUDACXX_END_NAMESPACE_CUDA #endif // _LIBCUDACXX_STD_VER > 11 From f8b8d0a68cd003d6d1a073788d075a7d115e2fda Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 5/9] Move properties out of namespace mr and add `forward_property` helper --- .../get_property/forward_property.pass.cpp | 75 +++++++++++++++ .../has_property.pass.cpp | 28 +++--- .../async_resource_ref.construction.pass.cpp | 4 +- .../async_resource_ref.conversion.pass.cpp | 4 +- .../async_resource_ref.equality.fail.cpp | 4 +- .../async_resource_ref.equality.pass.cpp | 4 +- .../async_resource_ref.inheritance.pass.cpp | 4 +- .../async_resource_ref.properties.pass.cpp | 26 ++--- .../resource_ref.construction.pass.cpp | 4 +- .../resource_ref.conversion.pass.cpp | 4 +- .../resource_ref.equality.fail.cpp | 4 +- .../resource_ref.equality.pass.cpp | 4 +- .../resource_ref.inheritance.pass.cpp | 4 +- .../resource_ref.properties.pass.cpp | 26 ++--- include/cuda/memory_resource | 96 ++++++++++++------- 15 files changed, 195 insertions(+), 96 deletions(-) create mode 100644 .upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp rename .upstream-tests/test/cuda/memory_resource/{memory_resource.concepts => get_property}/has_property.pass.cpp (58%) diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp new file mode 100644 index 0000000000..c6f8597f19 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp @@ -0,0 +1,75 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::forward_property +#include +#include + +struct prop_with_value { + using value_type = int; +}; +struct prop {}; + +template +struct derived_plain : public cuda::forward_property, Upstream> +{ + __host__ __device__ constexpr Upstream upstream_resource() const noexcept { return Upstream{}; } +}; + +struct upstream_with_valueless_property { + __host__ __device__ friend constexpr void get_property(const upstream_with_valueless_property&, prop) {} +}; +static_assert( cuda::has_property, prop>, ""); +static_assert(!cuda::has_property, prop_with_value>, ""); + +struct upstream_with_stateful_property { + __host__ __device__ friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) { + return 42; + } +}; +static_assert(!cuda::has_property, prop>, ""); +static_assert( cuda::has_property, prop_with_value>, ""); + +struct upstream_with_both_properties { + __host__ __device__ friend constexpr void get_property(const upstream_with_both_properties&, prop) {} + __host__ __device__ friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) { + return 42; + } +}; +static_assert( cuda::has_property, prop>, ""); +static_assert( cuda::has_property, prop_with_value>, ""); + +struct derived_override : public cuda::forward_property +{ + __host__ __device__ constexpr upstream_with_both_properties upstream_resource() const noexcept { + return upstream_with_both_properties{}; + } + __host__ __device__ friend constexpr int get_property(const derived_override&, prop_with_value) { + return 1337; + } +}; + +__host__ __device__ constexpr bool test_stateful() { + using derived_no_override = derived_plain; + const derived_no_override without_override{}; + assert(get_property(without_override, prop_with_value{}) == 42); + + const derived_override with_override{}; + assert(get_property(with_override, prop_with_value{}) == 1337); + + return true; +} + +int main(int, char**) { + test_stateful(); + static_assert(test_stateful(), ""); + return 0; +} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp similarity index 58% rename from .upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp rename to .upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp index 9a727de082..6090c88f35 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/has_property.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp @@ -9,7 +9,7 @@ // UNSUPPORTED: c++03, c++11 -// cuda::mr::has_property, cuda::mr::has_property_with +// cuda::has_property, cuda::has_property_with #include struct prop_with_value { @@ -17,15 +17,15 @@ struct prop_with_value { }; struct prop {}; -static_assert(cuda::mr::property_with_value); -static_assert(!cuda::mr::property_with_value); +static_assert(cuda::property_with_value); +static_assert(!cuda::property_with_value); struct valid_property { friend void get_property(const valid_property&, prop) {} }; -static_assert(!cuda::mr::has_property, ""); -static_assert(cuda::mr::has_property, ""); -static_assert(!cuda::mr::has_property_with, ""); +static_assert(!cuda::has_property, ""); +static_assert(cuda::has_property, ""); +static_assert(!cuda::has_property_with, ""); struct valid_property_with_value { friend int get_property(const valid_property_with_value&, prop_with_value) { @@ -33,12 +33,12 @@ struct valid_property_with_value { } }; static_assert( - cuda::mr::has_property, ""); -static_assert(!cuda::mr::has_property, ""); -static_assert(cuda::mr::has_property_with, ""); +static_assert(!cuda::has_property, ""); +static_assert(cuda::has_property_with, ""); -static_assert(!cuda::mr::has_property_with, ""); @@ -47,13 +47,13 @@ struct derived_from_property : public valid_property { return 42; } }; -static_assert(cuda::mr::has_property, +static_assert(cuda::has_property, ""); -static_assert(cuda::mr::has_property, ""); +static_assert(cuda::has_property, ""); static_assert( - cuda::mr::has_property_with, + cuda::has_property_with, ""); -static_assert(!cuda::mr::has_property_with, ""); diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp index e3dd1feb6f..7d773f0377 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -47,12 +47,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp index 392272e6b3..c1da2c5a09 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp @@ -52,12 +52,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp index b8d9ac5fc9..77b39d0cb7 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -49,12 +49,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp index 8991aeb2da..e5ef35e98f 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -49,12 +49,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp index 6d3966baec..4ecbda80b8 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp @@ -38,12 +38,12 @@ struct async_resource_base { bool operator!=(const async_resource_base& other) const { return false; } _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const async_resource_base&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const async_resource_base& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp index 9d637db781..d1c24b7267 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -25,13 +25,13 @@ template struct property_without_value {}; namespace properties_test { -static_assert(cuda::mr::property_with_value >, ""); +static_assert(cuda::property_with_value >, ""); static_assert( - cuda::mr::property_with_value >, ""); + cuda::property_with_value >, ""); -static_assert(!cuda::mr::property_with_value >, ""); +static_assert(!cuda::property_with_value >, ""); static_assert( - !cuda::mr::property_with_value >, + !cuda::property_with_value >, ""); } // namespace properties_test @@ -56,12 +56,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { @@ -86,25 +86,25 @@ static_assert( (2 * sizeof(void*))); _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::mr::property_with_value) // +(requires !cuda::property_with_value) // int InvokeIfWithValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::mr::property_with_value) // +(requires cuda::property_with_value) // typename Property::value_type InvokeIfWithValue(const Ref& ref) { return get_property(ref, Property{}); } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::mr::property_with_value) // +(requires cuda::property_with_value) // int InvokeIfWithoutValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::mr::property_with_value) // +(requires !cuda::property_with_value) // int InvokeIfWithoutValue(const Ref& ref) { get_property(ref, Property{}); return 1; @@ -119,7 +119,7 @@ void test_async_resource_ref() { // Check all the potentially stateful properties const int properties_with_value[] = {InvokeIfWithValue(ref)...}; const int expected_with_value[] = { - ((cuda::mr::property_with_value) ? expected_initially + ((cuda::property_with_value) ? expected_initially : -1)...}; for (std::size_t i = 0; i < sizeof...(Properties); ++i) { assert(properties_with_value[i] == expected_with_value[i]); @@ -128,7 +128,7 @@ void test_async_resource_ref() { const int properties_without_value[] = { InvokeIfWithoutValue(ref)...}; const int expected_without_value[] = { - ((cuda::mr::property_with_value) ? -1 : 1)...}; + ((cuda::property_with_value) ? -1 : 1)...}; for (std::size_t i = 0; i < sizeof...(Properties); ++i) { assert(properties_without_value[i] == expected_without_value[i]); } @@ -139,7 +139,7 @@ void test_async_resource_ref() { // Check whether we truly get the right value const int properties_with_value2[] = {InvokeIfWithValue(ref)...}; const int expected_with_value2[] = { - ((cuda::mr::property_with_value) ? expected_after_change + ((cuda::property_with_value) ? expected_after_change : -1)...}; for (std::size_t i = 0; i < sizeof...(Properties); ++i) { assert(properties_with_value2[i] == expected_with_value2[i]); diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp index 774973a397..20e2b6e4a5 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp @@ -37,12 +37,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp index eb633032e2..75dff8f2d5 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp @@ -52,12 +52,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp index af9be2f75c..0cf95c89be 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -38,12 +38,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp index f32f1abbb3..6049a90e62 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -38,12 +38,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp index e398ad313a..dd598d662b 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp @@ -33,12 +33,12 @@ struct resource_base { bool operator!=(const resource_base& other) const { return false; } _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource_base&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource_base& res, Property) noexcept { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp index 759b1ed1c8..26fddacf8a 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -24,13 +24,13 @@ template struct property_without_value {}; namespace properties_test { -static_assert(cuda::mr::property_with_value >, ""); +static_assert(cuda::property_with_value >, ""); static_assert( - cuda::mr::property_with_value >, ""); + cuda::property_with_value >, ""); -static_assert(!cuda::mr::property_with_value >, ""); +static_assert(!cuda::property_with_value >, ""); static_assert( - !cuda::mr::property_with_value >, + !cuda::property_with_value >, ""); } // namespace properties_test @@ -47,12 +47,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::mr::property_with_value && + (requires !cuda::property_with_value && _CUDA_VSTD::_One_of) // friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::mr::property_with_value&& + (requires cuda::property_with_value&& _CUDA_VSTD::_One_of) // friend typename Property::value_type get_property(const resource& res, Property) noexcept { @@ -75,25 +75,25 @@ static_assert(sizeof(cuda::mr::resource_ref, (2 * sizeof(void*))); _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::mr::property_with_value) // +(requires !cuda::property_with_value) // int InvokeIfWithValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::mr::property_with_value) // +(requires cuda::property_with_value) // typename Property::value_type InvokeIfWithValue(const Ref& ref) { return get_property(ref, Property{}); } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::mr::property_with_value) // +(requires cuda::property_with_value) // int InvokeIfWithoutValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::mr::property_with_value) // +(requires !cuda::property_with_value) // int InvokeIfWithoutValue(const Ref& ref) { get_property(ref, Property{}); return 1; @@ -108,7 +108,7 @@ void test_resource_ref() { // Check all the potentially stateful properties const int properties_with_value[] = {InvokeIfWithValue(ref)...}; const int expected_with_value[] = { - ((cuda::mr::property_with_value) ? expected_initially + ((cuda::property_with_value) ? expected_initially : -1)...}; for (std::size_t i = 0; i < sizeof...(Properties); ++i) { assert(properties_with_value[i] == expected_with_value[i]); @@ -117,7 +117,7 @@ void test_resource_ref() { const int properties_without_value[] = { InvokeIfWithoutValue(ref)...}; const int expected_without_value[] = { - ((cuda::mr::property_with_value) ? -1 : 1)...}; + ((cuda::property_with_value) ? -1 : 1)...}; for (std::size_t i = 0; i < sizeof...(Properties); ++i) { assert(properties_without_value[i] == expected_without_value[i]); } @@ -128,7 +128,7 @@ void test_resource_ref() { // Check whether we truly get the right value const int properties_with_value2[] = {InvokeIfWithValue(ref)...}; const int expected_with_value2[] = { - ((cuda::mr::property_with_value) ? expected_after_change + ((cuda::property_with_value) ? expected_after_change : -1)...}; for (std::size_t i = 0; i < sizeof...(Properties); ++i) { assert(properties_with_value2[i] == expected_with_value2[i]); diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index 45a6a1cfc8..af480f22f5 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -89,40 +89,9 @@ class resource_ref { #if _LIBCUDACXX_STD_VER > 11 _LIBCUDACXX_BEGIN_NAMESPACE_CUDA -namespace mr -{ /////////////////////////////////////////////////////////////////////////////// -// memory_resource - -/// \concept resource -/// \brief The \c resource concept -template -_LIBCUDACXX_CONCEPT_FRAGMENT( - __resource_, - requires(_Resource& __res, void* __ptr, size_t __bytes, size_t __alignment) ( - requires(_CUDA_VSTD::same_as), - requires(_CUDA_VSTD::same_as), - requires(_CUDA_VSTD::equality_comparable<_Resource>) - )); - -template -_LIBCUDACXX_CONCEPT resource = _LIBCUDACXX_FRAGMENT(__resource_, _Resource); - -/// \concept async_resource -/// \brief The \c async_resource concept -template -_LIBCUDACXX_CONCEPT_FRAGMENT( - __async_resource_, - requires(_Resource& __res, void* __ptr, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) ( - requires(resource<_Resource>), - requires(_CUDA_VSTD::same_as), - requires(_CUDA_VSTD::same_as), - requires(_CUDA_VSTD::equality_comparable<_Resource>) - )); - -template -_LIBCUDACXX_CONCEPT async_resource = _LIBCUDACXX_FRAGMENT(__async_resource_, _Resource); +// properties /// \concept has_property /// \brief The \c has_property concept @@ -137,11 +106,14 @@ _LIBCUDACXX_CONCEPT has_property = _LIBCUDACXX_FRAGMENT(__has_property_, _Resour /// \concept property_with_value /// \brief The \c property_with_value concept +template +using __property_value_t = typename _Property::value_type; + template _LIBCUDACXX_CONCEPT_FRAGMENT( __property_with_value_, requires()( - typename(typename _Property::value_type) + typename(__property_value_t<_Property>) )); template _LIBCUDACXX_CONCEPT property_with_value = _LIBCUDACXX_FRAGMENT(__property_with_value_, _Property); @@ -158,6 +130,58 @@ _LIBCUDACXX_CONCEPT_FRAGMENT( template _LIBCUDACXX_CONCEPT has_property_with = _LIBCUDACXX_FRAGMENT(__has_property_with_, _Resource, _Property, _Return); +/// class forward_property +/// \brief The \c forward_property crtp temaplate simplifies the user facing side of forwarding properties +/// We can always tell people to just derive from it to properly forward all properties +template +struct forward_property { + _LIBCUDACXX_TEMPLATE(class _Property) + (requires (!property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) + _LIBCUDACXX_INLINE_VISIBILITY friend constexpr void get_property(const _Derived&, _Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class _Property) + (requires (property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) + _LIBCUDACXX_INLINE_VISIBILITY friend constexpr __property_value_t<_Property> get_property( + const _Derived& __res, _Property __prop) { + return get_property(__res.upstream_resource(), __prop); + } +}; + +namespace mr +{ + +/////////////////////////////////////////////////////////////////////////////// +// memory_resource + +/// \concept resource +/// \brief The \c resource concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT( + __resource_, + requires(_Resource& __res, void* __ptr, size_t __bytes, size_t __alignment) ( + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::equality_comparable<_Resource>) + )); + +template +_LIBCUDACXX_CONCEPT resource = _LIBCUDACXX_FRAGMENT(__resource_, _Resource); + +/// \concept async_resource +/// \brief The \c async_resource concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT( + __async_resource_, + requires(_Resource& __res, void* __ptr, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) ( + requires(resource<_Resource>), + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::same_as), + requires(_CUDA_VSTD::equality_comparable<_Resource>) + )); + +template +_LIBCUDACXX_CONCEPT async_resource = _LIBCUDACXX_FRAGMENT(__async_resource_, _Resource); + /// \concept resource_with /// \brief The \c resource_with concept template @@ -228,7 +252,7 @@ struct _Async_alloc_vtable : public _Alloc_vtable struct _Resource_vtable_builder { template - static typename _Property::value_type _Get_property(void* __res) noexcept { + static __property_value_t<_Property> _Get_property(void* __res) noexcept { return get_property(*static_cast(__res), _Property{}); } @@ -282,7 +306,7 @@ struct _Resource_vtable_builder template struct _Property_vtable { - using _PropertyFn = typename _Property::value_type (*)(void*); + using _PropertyFn = __property_value_t<_Property> (*)(void*); _PropertyFn __property_fn = nullptr; constexpr _Property_vtable(_PropertyFn __property_fn_) noexcept @@ -516,7 +540,7 @@ public: _LIBCUDACXX_TEMPLATE(class _Property) (requires property_with_value<_Property> &&_CUDA_VSTD::_One_of<_Property, _Properties...>) // - friend typename _Property::value_type get_property( + friend __property_value_t<_Property> get_property( const basic_resource_ref &__res, _Property) noexcept { return __res._Property_vtable<_Property>::__property_fn(__res.__object); } From d5c53082dc9c9231ae9e1df1b1e31b232dd95e8f Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 6/9] Implement `cuda::get_property` niebloid --- .../get_property/get_property.pass.cpp | 61 +++++++++++++++++++ include/cuda/memory_resource | 28 +++++++++ 2 files changed, 89 insertions(+) create mode 100644 .upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp new file mode 100644 index 0000000000..fab8609cee --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp @@ -0,0 +1,61 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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 +// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11 + +// cuda::get_property +#include +#include + +struct prop_with_value { + using value_type = int; +}; +struct prop {}; + +struct upstream_with_valueless_property { + friend constexpr void get_property(const upstream_with_valueless_property&, prop) {} +}; +static_assert( cuda::std::invocable, ""); +static_assert(!cuda::std::invocable, ""); + +struct upstream_with_stateful_property { + friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) { + return 42; + } +}; +static_assert(!cuda::std::invocable, ""); +static_assert( cuda::std::invocable, ""); + +struct upstream_with_both_properties { + friend constexpr void get_property(const upstream_with_both_properties&, prop) {} + friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) { + return 42; + } +}; +static_assert( cuda::std::invocable, ""); +static_assert( cuda::std::invocable, ""); + +__host__ __device__ constexpr bool test() { + upstream_with_valueless_property with_valueless{}; + cuda::get_property(with_valueless, prop{}); + + upstream_with_stateful_property with_value{}; + assert(cuda::get_property(with_value, prop_with_value{}) == 42); + + upstream_with_both_properties with_both{}; + cuda::get_property(with_both, prop{}); + assert(cuda::get_property(with_both, prop_with_value{}) == 42); + return true; +} + +int main(int, char**) { + test(); + static_assert(test(), ""); + return 0; +} diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index af480f22f5..42ae360a7c 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -147,6 +147,34 @@ struct forward_property { } }; +/// class get_property +/// \brief The \c get_property crtp temaplate simplifies the user facing side of forwarding properties +/// We can always tell people to just derive from it to properly forward all properties +_LIBCUDACXX_BEGIN_NAMESPACE_CPO(__get_property) +struct __fn { + #if defined(__CUDACC__) && !defined(__NVCOMPILER) + #pragma nv_exec_check_disable + #endif + _LIBCUDACXX_TEMPLATE(class _Upstream, class _Property) + (requires (!property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) + _LIBCUDACXX_INLINE_VISIBILITY constexpr void operator()(const _Upstream&, _Property) const noexcept {} + + #if defined(__CUDACC__) && !defined(__NVCOMPILER) + #pragma nv_exec_check_disable + #endif + _LIBCUDACXX_TEMPLATE(class _Upstream, class _Property) + (requires (property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) + _LIBCUDACXX_INLINE_VISIBILITY constexpr __property_value_t<_Property> operator()( + const _Upstream& __res, _Property __prop) const { + return get_property(__res, __prop); + } +}; +_LIBCUDACXX_END_NAMESPACE_CPO + +inline namespace __cpo { + _LIBCUDACXX_CPO_ACCESSIBILITY _LIBCUDACXX_INLINE_VAR constexpr auto get_property = __get_property::__fn{}; +} // namespace __cpo + namespace mr { From b809e677609b32be787b0fb89111ce6b0a98c224 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Tue, 14 Feb 2023 09:49:41 +0100 Subject: [PATCH 7/9] Improve constraints on `forward_property` to require a `upstream_resource` method --- .../get_property/forward_property.pass.cpp | 159 ++++++++++++------ include/cuda/memory_resource | 27 ++- 2 files changed, 128 insertions(+), 58 deletions(-) diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp index c6f8597f19..fd7dcc167e 100644 --- a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp @@ -13,63 +13,114 @@ #include #include -struct prop_with_value { - using value_type = int; -}; -struct prop {}; - -template -struct derived_plain : public cuda::forward_property, Upstream> -{ - __host__ __device__ constexpr Upstream upstream_resource() const noexcept { return Upstream{}; } -}; - -struct upstream_with_valueless_property { - __host__ __device__ friend constexpr void get_property(const upstream_with_valueless_property&, prop) {} -}; -static_assert( cuda::has_property, prop>, ""); -static_assert(!cuda::has_property, prop_with_value>, ""); - -struct upstream_with_stateful_property { - __host__ __device__ friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) { - return 42; - } -}; -static_assert(!cuda::has_property, prop>, ""); -static_assert( cuda::has_property, prop_with_value>, ""); - -struct upstream_with_both_properties { - __host__ __device__ friend constexpr void get_property(const upstream_with_both_properties&, prop) {} - __host__ __device__ friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) { - return 42; - } -}; -static_assert( cuda::has_property, prop>, ""); -static_assert( cuda::has_property, prop_with_value>, ""); - -struct derived_override : public cuda::forward_property -{ - __host__ __device__ constexpr upstream_with_both_properties upstream_resource() const noexcept { - return upstream_with_both_properties{}; - } - __host__ __device__ friend constexpr int get_property(const derived_override&, prop_with_value) { - return 1337; - } -}; +namespace has_upstream_resource { + struct Upstream{}; -__host__ __device__ constexpr bool test_stateful() { - using derived_no_override = derived_plain; - const derived_no_override without_override{}; - assert(get_property(without_override, prop_with_value{}) == 42); - - const derived_override with_override{}; - assert(get_property(with_override, prop_with_value{}) == 1337); + __device__ Upstream upstream{}; + + struct with_reference { + Upstream& upstream_resource() const { return upstream; } + }; + static_assert(cuda::__has_upstream_resource, ""); + + struct with_const_reference { + const Upstream& upstream_resource() const { return upstream; } + }; + static_assert(cuda::__has_upstream_resource, ""); + + struct with_value { + Upstream upstream_resource() const { return Upstream{}; } + }; + static_assert(cuda::__has_upstream_resource, ""); + + struct with_const_value { + const Upstream upstream_resource() const { return Upstream{}; } + }; + static_assert(cuda::__has_upstream_resource, ""); + + struct Convertible { + operator Upstream() { return Upstream{}; } + }; + + struct with_conversion { + Convertible upstream_resource() const { return Convertible{}; } + }; + static_assert(!cuda::__has_upstream_resource, ""); +} // namespace has_upstream_resource + +namespace forward_property { + struct prop_with_value { using value_type = int; }; + struct prop {}; + + template + struct derived_plain : public cuda::forward_property, Upstream> + { + constexpr Upstream upstream_resource() const noexcept { return Upstream{}; } + }; + + struct upstream_with_valueless_property { + friend constexpr void get_property(const upstream_with_valueless_property&, prop) {} + }; + static_assert( cuda::has_property, prop>, ""); + static_assert(!cuda::has_property, prop_with_value>, ""); + + struct upstream_with_stateful_property { + friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) { + return 42; + } + }; + static_assert(!cuda::has_property, prop>, ""); + static_assert( cuda::has_property, prop_with_value>, ""); + + struct upstream_with_both_properties { + friend constexpr void get_property(const upstream_with_both_properties&, prop) {} + friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) { + return 42; + } + }; + static_assert( cuda::has_property, prop>, ""); + static_assert( cuda::has_property, prop_with_value>, ""); + + struct derived_override : public cuda::forward_property + { + constexpr upstream_with_both_properties upstream_resource() const noexcept { + return upstream_with_both_properties{}; + } + // Get called directly so needs to be annotated + __host__ __device__ friend constexpr int get_property(const derived_override&, prop_with_value) { + return 1337; + } + }; + + struct convertible_to_upstream { + operator upstream_with_both_properties() const noexcept { + return upstream_with_both_properties{}; + } + }; + + struct derived_with_converstin_upstream_resource : public cuda::forward_property + { + constexpr convertible_to_upstream upstream_resource() const noexcept { + return convertible_to_upstream{}; + } + }; + static_assert(!cuda::has_property); + + __host__ __device__ constexpr bool test_stateful() { + using derived_no_override = derived_plain; + const derived_no_override without_override{}; + assert(get_property(without_override, prop_with_value{}) == 42); + + const derived_override with_override{}; + assert(get_property(with_override, prop_with_value{}) == 1337); + + return true; + } +} // namespace forward_property - return true; -} int main(int, char**) { - test_stateful(); - static_assert(test_stateful(), ""); + forward_property::test_stateful(); + static_assert(forward_property::test_stateful(), ""); return 0; } diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index 42ae360a7c..31a869ef6b 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -130,17 +130,36 @@ _LIBCUDACXX_CONCEPT_FRAGMENT( template _LIBCUDACXX_CONCEPT has_property_with = _LIBCUDACXX_FRAGMENT(__has_property_with_, _Resource, _Property, _Return); +/// \concept __has_upstream_resource +/// \brief The \c __has_upstream_resource concept +template +_LIBCUDACXX_CONCEPT_FRAGMENT( + __has_upstream_resource_, + requires(const _Resource& __res)( + requires(_CUDA_VSTD::same_as<_CUDA_VSTD::__remove_const_ref_t, _Upstream>) + )); +template +_LIBCUDACXX_CONCEPT __has_upstream_resource = _LIBCUDACXX_FRAGMENT(__has_upstream_resource_, _Resource, _Upstream); + /// class forward_property -/// \brief The \c forward_property crtp temaplate simplifies the user facing side of forwarding properties -/// We can always tell people to just derive from it to properly forward all properties +/// \brief The \c forward_property crtp template simplifies the user facing side of forwarding properties +/// We can just derive from it to properly forward all properties template struct forward_property { + #if defined(__CUDACC__) && !defined(__NVCOMPILER) + #pragma nv_exec_check_disable + #endif _LIBCUDACXX_TEMPLATE(class _Property) (requires (!property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) _LIBCUDACXX_INLINE_VISIBILITY friend constexpr void get_property(const _Derived&, _Property) noexcept {} - _LIBCUDACXX_TEMPLATE(class _Property) - (requires (property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) + // The indirection is needed, otherwise the compiler might believe that _Derived is an incomplete type + #if defined(__CUDACC__) && !defined(__NVCOMPILER) + #pragma nv_exec_check_disable + #endif + _LIBCUDACXX_TEMPLATE(class _Property, class _Derived2 = _Derived) + (requires property_with_value<_Property> _LIBCUDACXX_AND has_property<_Upstream, _Property> _LIBCUDACXX_AND + __has_upstream_resource<_Derived2, _Upstream>) _LIBCUDACXX_INLINE_VISIBILITY friend constexpr __property_value_t<_Property> get_property( const _Derived& __res, _Property __prop) { return get_property(__res.upstream_resource(), __prop); From d0a55c51ef742bf19e705358e6fdfdd1b779d027 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 15 Feb 2023 15:34:14 +0100 Subject: [PATCH 8/9] Make `memory_resource` and `stream_ref` experimental --- .../get_property/forward_property.pass.cpp | 23 ++++++----- .../get_property/get_property.pass.cpp | 9 +++-- .../get_property/has_property.pass.cpp | 3 ++ .../async_resource_ref.allocate.pass.cpp | 3 ++ .../async_resource_ref.construction.pass.cpp | 3 ++ .../async_resource_ref.conversion.pass.cpp | 3 ++ .../async_resource_ref.equality.fail.cpp | 3 ++ .../async_resource_ref.equality.pass.cpp | 3 ++ .../async_resource_ref.inheritance.pass.cpp | 3 ++ .../async_resource_ref.properties.pass.cpp | 3 ++ .../async_resource.pass.cpp | 3 ++ .../async_resource_with.pass.cpp | 3 ++ .../resource.pass.cpp | 3 ++ .../resource_with.pass.cpp | 3 ++ .../resource_ref.allocate.pass.cpp | 3 ++ .../resource_ref.construction.pass.cpp | 3 ++ .../resource_ref.conversion.pass.cpp | 3 ++ .../resource_ref.equality.fail.cpp | 3 ++ .../resource_ref.equality.pass.cpp | 3 ++ .../resource_ref.inheritance.pass.cpp | 3 ++ .../resource_ref.properties.pass.cpp | 3 ++ .../stream_ref.constructor.pass.cpp | 2 + .../stream_ref/stream_ref.equality.pass.cpp | 2 + .../cuda/stream_ref/stream_ref.get.pass.cpp | 2 + .../cuda/stream_ref/stream_ref.ready.fail.cpp | 4 ++ .../cuda/stream_ref/stream_ref.ready.pass.cpp | 2 + .../cuda/stream_ref/stream_ref.wait.fail.cpp | 4 ++ .../cuda/stream_ref/stream_ref.wait.pass.cpp | 2 + include/cuda/memory_resource | 38 +++++++++++-------- include/cuda/stream_ref | 5 +++ 30 files changed, 121 insertions(+), 29 deletions(-) diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp index fd7dcc167e..fd76c205d8 100644 --- a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp @@ -10,10 +10,13 @@ // UNSUPPORTED: c++03, c++11 // cuda::forward_property + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include -namespace has_upstream_resource { +namespace has_upstream_resource { struct Upstream{}; __device__ Upstream upstream{}; @@ -53,7 +56,7 @@ namespace forward_property { struct prop {}; template - struct derived_plain : public cuda::forward_property, Upstream> + struct derived_plain : public cuda::forward_property, Upstream> { constexpr Upstream upstream_resource() const noexcept { return Upstream{}; } }; @@ -81,10 +84,10 @@ namespace forward_property { static_assert( cuda::has_property, prop>, ""); static_assert( cuda::has_property, prop_with_value>, ""); - struct derived_override : public cuda::forward_property + struct derived_override : public cuda::forward_property { - constexpr upstream_with_both_properties upstream_resource() const noexcept { - return upstream_with_both_properties{}; + constexpr upstream_with_both_properties upstream_resource() const noexcept { + return upstream_with_both_properties{}; } // Get called directly so needs to be annotated __host__ __device__ friend constexpr int get_property(const derived_override&, prop_with_value) { @@ -98,10 +101,10 @@ namespace forward_property { } }; - struct derived_with_converstin_upstream_resource : public cuda::forward_property + struct derived_with_converstin_upstream_resource : public cuda::forward_property { - constexpr convertible_to_upstream upstream_resource() const noexcept { - return convertible_to_upstream{}; + constexpr convertible_to_upstream upstream_resource() const noexcept { + return convertible_to_upstream{}; } }; static_assert(!cuda::has_property); @@ -110,7 +113,7 @@ namespace forward_property { using derived_no_override = derived_plain; const derived_no_override without_override{}; assert(get_property(without_override, prop_with_value{}) == 42); - + const derived_override with_override{}; assert(get_property(with_override, prop_with_value{}) == 1337); @@ -122,5 +125,5 @@ namespace forward_property { int main(int, char**) { forward_property::test_stateful(); static_assert(forward_property::test_stateful(), ""); - return 0; + return 0; } diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp index fab8609cee..5582faa653 100644 --- a/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::get_property + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include @@ -44,10 +47,10 @@ static_assert( cuda::std::invocable struct prop_with_value { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp index 8a013ddf80..ae565b9c57 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp index 7d773f0377..672a4fae96 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource_ref construction + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp index c1da2c5a09..15a04c2b71 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp index 77b39d0cb7..1116a13e11 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource_ref equality + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp index e5ef35e98f..a4feced151 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource_ref equality + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp index 4ecbda80b8..2f6f3bf883 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp index d1c24b7267..73ba361686 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp index 1551e4dc44..a6245bd012 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp index 967f701715..5f2210de28 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::async_resource_with + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp index bc2d7a86f7..6740350761 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp index 75020d7899..c63f0798c8 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_with + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp index 27a9439357..b90746c0dd 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp index 20e2b6e4a5..1374700a84 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref construction + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp index 75dff8f2d5..d93430935b 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp index 0cf95c89be..8c43807797 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref equality + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp index 6049a90e62..d4b8f4b3ea 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref equality + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp index dd598d662b..6add9c7ad9 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp index 26fddacf8a..b2641f5b44 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -10,6 +10,9 @@ // UNSUPPORTED: c++03, c++11 // cuda::mr::resource_ref properties + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp index 0c6a899f7a..7b77968255 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp @@ -7,6 +7,8 @@ // //===----------------------------------------------------------------------===// +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp index b749db58d3..a239eb060d 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.equality.pass.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp index dcf8c1647c..f7b9fb606d 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.get.pass.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp index fa8bcb8891..a298ac2ce5 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp @@ -8,6 +8,10 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: nvrtc + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp index 6df56cf894..e7fd5dc59a 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.pass.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp index 837d03d289..d3c9654260 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp @@ -8,6 +8,10 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: nvrtc + +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp index 26c03ca7f5..8d20775b74 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.pass.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + #include #include diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index 31a869ef6b..181770d0d3 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -31,7 +31,7 @@ concept async_resource = resource template concept has_property = resource && requires(const Resource& res, Property prop) { - get_property(res, prop); + get_property(res, prop); }; template @@ -40,11 +40,11 @@ concept property_with_value = requires { }; template -concept has_property_with = resource +concept has_property_with = resource && property_with_value && same_as && requires(const Resource& res, Property prop) { - get_property(res, prop) -> Return; + get_property(res, prop) -> Return; }; template @@ -79,12 +79,18 @@ class resource_ref { } // cuda */ // clang-format on + +#ifndef LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE +#error " is experimental and requires LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE to be defined" +#endif + #include #include #include #include + #include #if _LIBCUDACXX_STD_VER > 11 @@ -152,13 +158,13 @@ struct forward_property { _LIBCUDACXX_TEMPLATE(class _Property) (requires (!property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) _LIBCUDACXX_INLINE_VISIBILITY friend constexpr void get_property(const _Derived&, _Property) noexcept {} - + // The indirection is needed, otherwise the compiler might believe that _Derived is an incomplete type #if defined(__CUDACC__) && !defined(__NVCOMPILER) #pragma nv_exec_check_disable #endif _LIBCUDACXX_TEMPLATE(class _Property, class _Derived2 = _Derived) - (requires property_with_value<_Property> _LIBCUDACXX_AND has_property<_Upstream, _Property> _LIBCUDACXX_AND + (requires property_with_value<_Property> _LIBCUDACXX_AND has_property<_Upstream, _Property> _LIBCUDACXX_AND __has_upstream_resource<_Derived2, _Upstream>) _LIBCUDACXX_INLINE_VISIBILITY friend constexpr __property_value_t<_Property> get_property( const _Derived& __res, _Property __prop) { @@ -177,7 +183,7 @@ struct __fn { _LIBCUDACXX_TEMPLATE(class _Upstream, class _Property) (requires (!property_with_value<_Property>) _LIBCUDACXX_AND has_property<_Upstream, _Property>) _LIBCUDACXX_INLINE_VISIBILITY constexpr void operator()(const _Upstream&, _Property) const noexcept {} - + #if defined(__CUDACC__) && !defined(__NVCOMPILER) #pragma nv_exec_check_disable #endif @@ -246,7 +252,7 @@ template _LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant>...>; #else -_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> && +_LIBCUDACXX_CONCEPT async_resource_with = async_resource<_Resource> && (has_property<_Resource, _Properties> && ...); #endif @@ -512,14 +518,14 @@ private: public: // clang-format off _LIBCUDACXX_TEMPLATE(class _Resource) - (requires (!_Is_basic_resource_ref<_Resource> + (requires (!_Is_basic_resource_ref<_Resource> && (((_Alloc_type == _AllocType::_Default) && resource_with<_Resource, _Properties...>) // ||((_Alloc_type == _AllocType::_Async) && async_resource_with<_Resource, _Properties...>)))) // basic_resource_ref(_Resource& __res) noexcept : _Resource_ref_base<_Alloc_type>(&__res, &__alloc_vtable<_Alloc_type, _Resource>) , _Filtered_vtable<_Properties...>(_Filtered_vtable<_Properties...>::template _Create<_Resource>()) {} - + #if _LIBCUDACXX_STD_VER > 14 _LIBCUDACXX_TEMPLATE(class... _OtherProperties) (requires (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) @@ -533,14 +539,14 @@ public: : _Resource_ref_base<_Alloc_type>(__ref.__object, __ref.__static_vtable) , _Filtered_vtable<_Properties...>(__ref) {} - + #if _LIBCUDACXX_STD_VER > 14 _LIBCUDACXX_TEMPLATE(class... _OtherProperties) - (requires (_Alloc_type == _AllocType::_Default) + (requires (_Alloc_type == _AllocType::_Default) && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) #else _LIBCUDACXX_TEMPLATE(class... _OtherProperties) - (requires (_Alloc_type == _AllocType::_Default) + (requires (_Alloc_type == _AllocType::_Default) && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) #endif @@ -552,11 +558,11 @@ public: #if _LIBCUDACXX_STD_VER > 14 _LIBCUDACXX_TEMPLATE(class... _OtherProperties) - (requires(sizeof...(_Properties) == sizeof...(_OtherProperties)) + (requires(sizeof...(_Properties) == sizeof...(_OtherProperties)) && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) #else _LIBCUDACXX_TEMPLATE(class... _OtherProperties) - (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) #endif @@ -568,11 +574,11 @@ public: #if _LIBCUDACXX_STD_VER > 14 _LIBCUDACXX_TEMPLATE(class... _OtherProperties) - (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) #else _LIBCUDACXX_TEMPLATE(class... _OtherProperties) - (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) #endif diff --git a/include/cuda/stream_ref b/include/cuda/stream_ref index 2d69a0bc08..b4128421c9 100644 --- a/include/cuda/stream_ref +++ b/include/cuda/stream_ref @@ -39,10 +39,15 @@ private: } // cuda */ +#ifndef LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE +#error " is experimental and requires LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE to be defined" +#endif + #include // cuda_runtime_api needs to come first // clang-format on #include + #include #include From 572f73d281808e802d5e572ebb13952271515a33 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 15 Feb 2023 15:43:02 +0100 Subject: [PATCH 9/9] Make the tests pass --- .../get_property/forward_property.pass.cpp | 2 +- .../get_property/has_property.pass.cpp | 4 +-- .../async_resource_ref.construction.pass.cpp | 11 +++--- .../async_resource_ref.conversion.pass.cpp | 11 +++--- .../async_resource_ref.equality.fail.cpp | 1 + .../async_resource_ref.equality.pass.cpp | 11 +++--- .../async_resource_ref.inheritance.pass.cpp | 21 +++++------ .../async_resource_ref.properties.pass.cpp | 33 ++++++++--------- .../resource_ref.construction.pass.cpp | 11 +++--- .../resource_ref.conversion.pass.cpp | 11 +++--- .../resource_ref.equality.fail.cpp | 1 + .../resource_ref.equality.pass.cpp | 11 +++--- .../resource_ref.inheritance.pass.cpp | 21 +++++------ .../resource_ref.properties.pass.cpp | 33 ++++++++--------- .../stream_ref.constructor.pass.cpp | 18 +++++++--- .../cuda/stream_ref/stream_ref.ready.fail.cpp | 36 ------------------- .../cuda/stream_ref/stream_ref.wait.fail.cpp | 36 ------------------- include/cuda/memory_resource | 8 ++--- 18 files changed, 98 insertions(+), 182 deletions(-) delete mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp delete mode 100644 .upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp index fd76c205d8..234f4bc8cd 100644 --- a/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp @@ -107,7 +107,7 @@ namespace forward_property { return convertible_to_upstream{}; } }; - static_assert(!cuda::has_property); + static_assert(!cuda::has_property, ""); __host__ __device__ constexpr bool test_stateful() { using derived_no_override = derived_plain; diff --git a/.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp b/.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp index 6fda218306..c107aeb9d3 100644 --- a/.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp @@ -20,8 +20,8 @@ struct prop_with_value { }; struct prop {}; -static_assert(cuda::property_with_value); -static_assert(!cuda::property_with_value); +static_assert(cuda::property_with_value, ""); +static_assert(!cuda::property_with_value, ""); struct valid_property { friend void get_property(const valid_property&, prop) {} diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp index 672a4fae96..12b8e969de 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -50,15 +50,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const async_resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const async_resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { return res._val; } }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp index 15a04c2b71..e3e8bd8a2a 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp @@ -55,15 +55,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const resource& res, Property) noexcept { return res._val; } }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp index 1116a13e11..4db52db56b 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -8,6 +8,7 @@ //===----------------------------------------------------------------------===// // UNSUPPORTED: c++03, c++11 +// UNSUPPORTED: nvrtc // cuda::mr::async_resource_ref equality diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp index a4feced151..0d194714ea 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -52,15 +52,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const async_resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const async_resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { return res._val; } }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp index 2f6f3bf883..e24a6ca2bd 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp @@ -41,15 +41,12 @@ struct async_resource_base { bool operator!=(const async_resource_base& other) const { return false; } _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const async_resource_base&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource_base&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const async_resource_base& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const async_resource_base& res, Property) noexcept { return 42; } }; @@ -58,7 +55,6 @@ template struct async_resource_derived_first : public async_resource_base { using super_t = async_resource_base; - using super_t::operator==; async_resource_derived_first(const int val) : _val(val) {} @@ -73,9 +69,12 @@ struct async_resource_derived_first void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) override {} + bool operator==(const async_resource_derived_first& other) const { return true; } + bool operator!=(const async_resource_derived_first& other) const { return false; } + int _val = 0; }; -static_assert(cuda::mr::async_resource >); +static_assert(cuda::mr::async_resource >, ""); struct some_data { int _val; @@ -85,7 +84,6 @@ template struct async_resource_derived_second : public async_resource_base { using super_t = async_resource_base; - using super_t::operator==; async_resource_derived_second(some_data* val) : _val(val) {} @@ -100,6 +98,9 @@ struct async_resource_derived_second void deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) override {} + bool operator==(const async_resource_derived_second& other) const { return true; } + bool operator!=(const async_resource_derived_second& other) const { return false; } + some_data* _val = 0; }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp index 73ba361686..03c3b7ed08 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -59,15 +59,12 @@ struct async_resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const async_resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const async_resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const async_resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const async_resource& res, Property) noexcept { return res._val; } }; @@ -75,39 +72,39 @@ struct async_resource { // Ensure we have the right size static_assert(sizeof(cuda::mr::async_resource_ref, property_with_value >) == - (4 * sizeof(void*))); + (4 * sizeof(void*)), ""); static_assert( sizeof(cuda::mr::async_resource_ref, property_without_value >) == - (3 * sizeof(void*))); + (3 * sizeof(void*)), ""); static_assert(sizeof(cuda::mr::async_resource_ref, property_with_value >) == - (3 * sizeof(void*))); + (3 * sizeof(void*)), ""); static_assert( sizeof(cuda::mr::async_resource_ref, property_without_value >) == - (2 * sizeof(void*))); + (2 * sizeof(void*)), ""); _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::property_with_value) // + (requires (!cuda::property_with_value)) // int InvokeIfWithValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::property_with_value) // + (requires cuda::property_with_value) // typename Property::value_type InvokeIfWithValue(const Ref& ref) { return get_property(ref, Property{}); } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::property_with_value) // + (requires cuda::property_with_value) // int InvokeIfWithoutValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::property_with_value) // + (requires (!cuda::property_with_value)) // int InvokeIfWithoutValue(const Ref& ref) { get_property(ref, Property{}); return 1; @@ -155,12 +152,12 @@ void test_property_forwarding() { using ref = cuda::mr::async_resource_ref >; static_assert(cuda::mr::async_resource_with, - property_with_value >); + property_with_value >, ""); static_assert(!cuda::mr::async_resource_with, - property_with_value >); + property_with_value >, ""); static_assert( - cuda::mr::async_resource_with >); + cuda::mr::async_resource_with >, ""); } void test_async_resource_ref() { diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp index 1374700a84..e838a3a286 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp @@ -40,15 +40,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const resource& res, Property) noexcept { return res._val; } }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp index d93430935b..15f7331c74 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp @@ -55,15 +55,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const resource& res, Property) noexcept { return res._val; } }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp index 8c43807797..0efff529d3 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -8,6 +8,7 @@ //===----------------------------------------------------------------------===// // UNSUPPORTED: c++03, c++11 +// UNSUPPORTED: nvrtc // cuda::mr::resource_ref equality diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp index d4b8f4b3ea..40ce8ef9ca 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -41,15 +41,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const resource& res, Property) noexcept { return res._val; } }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp index 6add9c7ad9..547b254637 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp @@ -36,15 +36,12 @@ struct resource_base { bool operator!=(const resource_base& other) const { return false; } _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const resource_base&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const resource_base&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const resource_base& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const resource_base& res, Property) noexcept { return 42; } }; @@ -52,7 +49,6 @@ struct resource_base { template struct resource_derived_first : public resource_base { using super_t = resource_base; - using super_t::operator==; resource_derived_first(const int val) : _val(val) {} @@ -60,9 +56,12 @@ struct resource_derived_first : public resource_base { void deallocate(void* ptr, std::size_t, std::size_t) override {} + bool operator==(const resource_derived_first& other) const { return true; } + bool operator!=(const resource_derived_first& other) const { return false; } + int _val = 0; }; -static_assert(cuda::mr::resource >); +static_assert(cuda::mr::resource >, ""); struct some_data { int _val; @@ -71,7 +70,6 @@ struct some_data { template struct resource_derived_second : public resource_base { using super_t = resource_base; - using super_t::operator==; resource_derived_second(some_data* val) : _val(val) {} @@ -79,6 +77,9 @@ struct resource_derived_second : public resource_base { void deallocate(void* ptr, std::size_t, std::size_t) override {} + bool operator==(const resource_derived_second& other) const { return true; } + bool operator!=(const resource_derived_second& other) const { return false; } + some_data* _val = 0; }; diff --git a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp index b2641f5b44..a013d1e41e 100644 --- a/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -50,15 +50,12 @@ struct resource { int _val = 0; _LIBCUDACXX_TEMPLATE(class Property) - (requires !cuda::property_with_value && - _CUDA_VSTD::_One_of) // - friend void get_property(const resource&, Property) noexcept {} + (requires (!cuda::property_with_value) && _CUDA_VSTD::_One_of) // + friend void get_property(const resource&, Property) noexcept {} _LIBCUDACXX_TEMPLATE(class Property) - (requires cuda::property_with_value&& - _CUDA_VSTD::_One_of) // - friend typename Property::value_type - get_property(const resource& res, Property) noexcept { + (requires cuda::property_with_value && _CUDA_VSTD::_One_of) // + friend typename Property::value_type get_property(const resource& res, Property) noexcept { return res._val; } }; @@ -66,37 +63,37 @@ struct resource { // Ensure we have the right size static_assert(sizeof(cuda::mr::resource_ref, property_with_value >) == - (4 * sizeof(void*))); + (4 * sizeof(void*)), ""); static_assert(sizeof(cuda::mr::resource_ref, property_without_value >) == - (3 * sizeof(void*))); + (3 * sizeof(void*)), ""); static_assert(sizeof(cuda::mr::resource_ref, property_with_value >) == - (3 * sizeof(void*))); + (3 * sizeof(void*)), ""); static_assert(sizeof(cuda::mr::resource_ref, property_without_value >) == - (2 * sizeof(void*))); + (2 * sizeof(void*)), ""); _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::property_with_value) // + (requires (!cuda::property_with_value)) // int InvokeIfWithValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::property_with_value) // + (requires cuda::property_with_value) // typename Property::value_type InvokeIfWithValue(const Ref& ref) { return get_property(ref, Property{}); } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires cuda::property_with_value) // + (requires cuda::property_with_value) // int InvokeIfWithoutValue(const Ref& ref) { return -1; } _LIBCUDACXX_TEMPLATE(class Property, class Ref) -(requires !cuda::property_with_value) // + (requires (!cuda::property_with_value)) // int InvokeIfWithoutValue(const Ref& ref) { get_property(ref, Property{}); return 1; @@ -143,11 +140,11 @@ void test_property_forwarding() { using ref = cuda::mr::resource_ref >; static_assert(cuda::mr::resource_with, - property_with_value >); + property_with_value >, ""); static_assert(!cuda::mr::resource_with, - property_with_value >); + property_with_value >, ""); - static_assert(cuda::mr::resource_with >); + static_assert(cuda::mr::resource_with >, ""); } void test_resource_ref() { diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp index 7b77968255..06a945778d 100644 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp +++ b/.upstream-tests/test/cuda/stream_ref/stream_ref.constructor.pass.cpp @@ -13,20 +13,28 @@ #include #include -static_assert(cuda::std::is_default_constructible::value); -static_assert(!cuda::std::is_constructible::value); -static_assert(!cuda::std::is_constructible::value); +static_assert(cuda::std::is_default_constructible::value, ""); +static_assert(!cuda::std::is_constructible::value, ""); +static_assert(!cuda::std::is_constructible::value, ""); template using void_t = void; +#if TEST_STD_VER < 14 +template +struct has_value_type : cuda::std::false_type {}; +template +struct has_value_type> : cuda::std::true_type {}; +static_assert(has_value_type::value, ""); +#else template constexpr bool has_value_type = false; template -constexpr bool has_value_type > = true; - +constexpr bool has_value_type_v > = true; static_assert(has_value_type, ""); +#endif + int main(int argc, char** argv) { #ifndef __CUDA_ARCH__ diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp deleted file mode 100644 index a298ac2ce5..0000000000 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.ready.fail.cpp +++ /dev/null @@ -1,36 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 -// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -// -//===----------------------------------------------------------------------===// - -// UNSUPPORTED: nvrtc - -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE - -#include -#include - -int main(int argc, char** argv) { -#ifndef __CUDA_ARCH__ - cudaStream_t stream = reinterpret_cast(42); - cuda::stream_ref ref{stream}; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { - assert(ref.ready()); - } catch (const cuda::cuda_error& e) { - assert(cudaGetErrorString(e) == ""); - } catch (...) { - assert(false && "Should have thrown"); - } - assert(false); -#else - assert(ref.ready()); -#endif // _LIBCUDACXX_NO_EXCEPTIONS -#endif - return 0; -} diff --git a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp b/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp deleted file mode 100644 index d3c9654260..0000000000 --- a/.upstream-tests/test/cuda/stream_ref/stream_ref.wait.fail.cpp +++ /dev/null @@ -1,36 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// 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 -// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved. -// -//===----------------------------------------------------------------------===// - -// UNSUPPORTED: nvrtc - -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE - -#include -#include - -int main(int argc, char** argv) { -#ifndef __CUDA_ARCH__ - cudaStream_t stream = reinterpret_cast(42); - cuda::stream_ref ref{stream}; -#ifndef _LIBCUDACXX_NO_EXCEPTIONS - try { - ref.wait(); - } catch (const cuda::cuda_error& e) { - assert(cudaGetErrorString(e) == "Failed to synchronize stream."); - } catch (...) { - assert(false && "Should have thrown"); - } - assert(false); -#else - ref.wait(); -#endif // _LIBCUDACXX_NO_EXCEPTIONS -#endif - return 0; -} diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index 181770d0d3..a74cee1e1d 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -437,7 +437,7 @@ using _Filtered_vtable = typename _Filtered<_Properties...>::_Filtered_vtable::_ template struct _Alloc_base { - static_assert(cuda::std::is_base_of_v<_Alloc_vtable, _Vtable>); + static_assert(cuda::std::is_base_of_v<_Alloc_vtable, _Vtable>, ""); _Alloc_base(void* __object_, const _Vtable* __static_vtabl_) noexcept : __object(__object_) @@ -462,7 +462,7 @@ protected: template struct _Async_alloc_base : public _Alloc_base<_Vtable> { - static_assert(cuda::std::is_base_of_v<_Async_alloc_vtable, _Vtable>); + static_assert(cuda::std::is_base_of_v<_Async_alloc_vtable, _Vtable>, ""); _Async_alloc_base(void* __object_, const _Vtable* __static_vtabl_) noexcept : _Alloc_base<_Vtable>(__object_, __static_vtabl_) @@ -588,11 +588,11 @@ public: } _LIBCUDACXX_TEMPLATE(class _Property) - (requires !property_with_value<_Property> && _CUDA_VSTD::_One_of<_Property, _Properties...>) // + (requires (!property_with_value<_Property>) _LIBCUDACXX_AND _CUDA_VSTD::_One_of<_Property, _Properties...>) // friend void get_property(const basic_resource_ref &, _Property) noexcept {} _LIBCUDACXX_TEMPLATE(class _Property) - (requires property_with_value<_Property> &&_CUDA_VSTD::_One_of<_Property, _Properties...>) // + (requires property_with_value<_Property> _LIBCUDACXX_AND _CUDA_VSTD::_One_of<_Property, _Properties...>) // friend __property_value_t<_Property> get_property( const basic_resource_ref &__res, _Property) noexcept { return __res._Property_vtable<_Property>::__property_fn(__res.__object);