From 530c0c28c05af58562a09a10cee0599e2382fbfa Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 14 Sep 2022 10:07:39 +0200 Subject: [PATCH] Implement resource wrappers for This includes * `cuda::mr::resource_ref` * `cuda::mr::async_resource_ref` --- .../async_resource_ref.construction.pass.cpp | 124 +++++++ .../async_resource_ref.conversion.cpp | 97 ++++++ .../async_resource_ref.equality.fail.cpp | 98 ++++++ .../async_resource_ref.equality.pass.cpp | 117 +++++++ .../async_resource_ref.properties.pass.cpp | 246 ++++++++++++++ .../resource_ref.construction.pass.cpp | 111 +++++++ .../resource_ref.conversion.cpp | 119 +++++++ .../resource_ref.equality.fail.cpp | 88 +++++ .../resource_ref.equality.pass.cpp | 106 ++++++ .../resource_ref.properties.pass.cpp | 184 +++++++++++ include/cuda/memory_resource | 311 ++++++++++++++++++ 11 files changed, 1601 insertions(+) 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.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.properties.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.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.properties.pass.cpp 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..c4e0f37880 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -0,0 +1,124 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return nullptr; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) {} + + inline __host__ __device__ void* allocate_async(std::size_t, std::size_t, + cuda::stream_ref) { + return &_val; + } + + inline __host__ __device__ 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); + } + + inline __host__ __device__ bool operator==(const async_resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.cpp new file mode 100644 index 0000000000..e6039ad45c --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.cpp @@ -0,0 +1,97 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ void* allocate_async(std::size_t, std::size_t, + cuda::stream_ref) { + return &_val; + } + + inline __host__ __device__ 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); + } + + inline __host__ __device__ bool operator==(const resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +template +__host__ __device__ void test_conversion_from_async_resource_ref() { + resource input{42}; + cuda::mr::async_resource_ref ref_input{input}; + cuda::mr::async_resource_ref ref{ref_input}; + + // 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**) { + 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 >(); + 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..fa77f7e359 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -0,0 +1,98 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return nullptr; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) {} + + inline __host__ __device__ void* allocate_async(std::size_t, std::size_t, + cuda::stream_ref) { + return &_val; + } + + inline __host__ __device__ 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); + } + + inline __host__ __device__ bool operator==(const async_resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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 >; + +__host__ __device__ 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**) { + test_equality(); + 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..d4a0991e34 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -0,0 +1,117 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return nullptr; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) {} + + inline __host__ __device__ void* allocate_async(std::size_t, std::size_t, + cuda::stream_ref) { + return &_val; + } + + inline __host__ __device__ 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); + } + + inline __host__ __device__ bool operator==(const async_resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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 >; + +__host__ __device__ 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**) { + test_equality(); + 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..ea0aedeb35 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -0,0 +1,246 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ void* allocate_async(std::size_t, std::size_t, + cuda::stream_ref) { + return &_val; + } + + inline __host__ __device__ 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); + } + + inline __host__ __device__ bool + operator==(const async_resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const async_resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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) // + __host__ __device__ int InvokeIfWithValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires cuda::mr::property_with_value) // + __host__ __device__ 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) // + __host__ __device__ int InvokeIfWithoutValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires !cuda::mr::property_with_value) // + __host__ __device__ int InvokeIfWithoutValue(const Ref& ref) { + get_property(ref, Property{}); + return 1; +} + +template +__host__ __device__ void test_async_resource_ref() { + 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, {})); + + // Check all the potentially stateful properties + const int properties_with_value[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value[] = { + ((cuda::mr::property_with_value) ? 42 : -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]); + } + + // 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); + + // 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_deallocate + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value2[i] == expected_with_value2[i]); + } +} + +// Ensure that a async resource properly works as a "plain" resource +template +__host__ __device__ void test_async_resource_ref_plain() { + 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)); + + // Check all the potentially stateful properties + const int properties_with_value[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value[] = { + ((cuda::mr::property_with_value) ? 42 : -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]); + } + + // 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); + + // 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_deallocate + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value2[i] == expected_with_value2[i]); + } +} + +__host__ __device__ 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 >); +} + +__host__ __device__ 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(); + + // Basic tests for plain alloc interface + test_async_resource_ref_plain, + property_with_value >(); +} +} // namespace resource_test + +int main(int, char**) { + resource_test::test_async_resource_ref(); + 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..1ce67d40d3 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.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::resource_ref construction +#include + +#include + +template +struct property_with_value { + using value_type = T; +}; + +template +struct property_without_value {}; + +template +struct resource { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ bool operator==(const resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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.cpp b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.cpp new file mode 100644 index 0000000000..6abf6bdc33 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.cpp @@ -0,0 +1,119 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ void* allocate_async(std::size_t, std::size_t, + cuda::stream_ref) { + return &_val; + } + + inline __host__ __device__ 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); + } + + inline __host__ __device__ bool operator==(const resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + friend typename Property::value_type + get_property(const resource& res, Property) noexcept { + return res._val; + } +}; + +template +__host__ __device__ void test_conversion_from_resource_ref() { + resource input{42}; + cuda::mr::resource_ref ref_input{input}; + cuda::mr::resource_ref ref{ref_input}; + + // 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 +__host__ __device__ void test_conversion_from_async_resource_ref() { + resource input{42}; + cuda::mr::async_resource_ref ref_input{input}; + cuda::mr::resource_ref ref{ref_input}; + + // 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**) { + 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 >(); + 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..dc84ff82ea --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -0,0 +1,88 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ bool operator==(const resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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 >; + +__host__ __device__ 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**) { + test_equality(); + 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..b2b6b04592 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -0,0 +1,106 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ bool operator==(const resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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 >; + +__host__ __device__ 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**) { + test_equality(); + 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..22c8fa9af6 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -0,0 +1,184 @@ +//===----------------------------------------------------------------------===// +// +// 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 { + inline __host__ __device__ void* allocate(std::size_t, std::size_t) { + return &_val; + } + + inline __host__ __device__ void deallocate(void* ptr, std::size_t, + std::size_t) { + // ensure that we did get the right inputs forwarded + _val = *static_cast(ptr); + } + + inline __host__ __device__ bool operator==(const resource& other) const { + return _val == other._val; + } + inline __host__ __device__ 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) // + inline __host__ __device__ + friend void get_property(const resource&, Property) noexcept {} + + _LIBCUDACXX_TEMPLATE(class Property) + (requires cuda::mr::property_with_value&& + _CUDA_VSTD::_One_of) // + inline __host__ __device__ // + 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) // + __host__ __device__ int InvokeIfWithValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires cuda::mr::property_with_value) // + __host__ __device__ 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) // + __host__ __device__ int InvokeIfWithoutValue(const Ref& ref) { + return -1; +} + +_LIBCUDACXX_TEMPLATE(class Property, class Ref) +(requires !cuda::mr::property_with_value) // + __host__ __device__ int InvokeIfWithoutValue(const Ref& ref) { + get_property(ref, Property{}); + return 1; +} + +template +__host__ __device__ void test_resource_ref() { + 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)); + + // Check all the potentially stateful properties + const int properties_with_value[] = {InvokeIfWithValue(ref)...}; + const int expected_with_value[] = { + ((cuda::mr::property_with_value) ? 42 : -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]); + } + + // 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); + + // 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_deallocate + : -1)...}; + for (std::size_t i = 0; i < sizeof...(Properties); ++i) { + assert(properties_with_value2[i] == expected_with_value2[i]); + } +} + +__host__ __device__ 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 >); +} + +__host__ __device__ 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**) { + resource_test::test_resource_ref(); + return 0; +} diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index e2afb524f7..fa5f6fa856 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -199,6 +199,317 @@ _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; + + _LIBCUDACXX_INLINE_VISIBILITY 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; + + _LIBCUDACXX_INLINE_VISIBILITY 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 + _LIBCUDACXX_INLINE_VISIBILITY + static typename _Property::value_type _Get_property(void* __res) noexcept { + return get_property(*static_cast(__res), _Property{}); + } + + template + _LIBCUDACXX_INLINE_VISIBILITY + static void* _Alloc(void* __object, size_t __bytes, size_t __alignment) { + return static_cast<_Resource *>(__object)->allocate(__bytes, __alignment); + } + + template + _LIBCUDACXX_INLINE_VISIBILITY + static void _Dealloc(void* __object, void* __ptr, size_t __bytes, size_t __alignment) { + return static_cast<_Resource *>(__object)->deallocate(__ptr, __bytes, __alignment); + } + + template + _LIBCUDACXX_INLINE_VISIBILITY + 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 + _LIBCUDACXX_INLINE_VISIBILITY + 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 + _LIBCUDACXX_INLINE_VISIBILITY + 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)) // + _LIBCUDACXX_INLINE_VISIBILITY static constexpr _Async_alloc_vtable _Create() noexcept + { + return {&_Resource_vtable_builder::_Alloc<_Resource>, + &_Resource_vtable_builder::_Dealloc<_Resource>, + &_Resource_vtable_builder::_Equal<_Resource>, + nullptr, + nullptr}; + } + + _LIBCUDACXX_TEMPLATE(class _Resource, _AllocType _Alloc_type) + (requires(_Alloc_type == _AllocType::_Async)) // + _LIBCUDACXX_INLINE_VISIBILITY 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 <_AllocType _Alloc_type, class _Resource> +__device__ _LIBCUDACXX_INLINE_VAR constexpr _Async_alloc_vtable + __alloc_vtable = _Resource_vtable_builder::template _Create<_Resource, _Alloc_type>(); + +template +struct _Property_vtable +{ + using _PropertyFn = typename _Property::value_type (*)(void*); + _PropertyFn __property_fn = nullptr; + + _Property_vtable() = default; + + _LIBCUDACXX_INLINE_VISIBILITY + constexpr _Property_vtable(_PropertyFn __property_fn_) noexcept + : __property_fn(__property_fn_) + {} +}; + +template +struct _Resource_vtable : public _Property_vtable<_Properties>... +{ + _Resource_vtable() = default; + + template + _LIBCUDACXX_INLINE_VISIBILITY constexpr _Resource_vtable(_PropertyFns... __property_fn_) noexcept + : _Property_vtable<_Properties>(__property_fn_)... + {} + + template + _LIBCUDACXX_INLINE_VISIBILITY 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 <_AllocType _Alloc_type, class... _Properties> // +class basic_resource_ref : protected _Filtered_vtable<_Properties...> +{ +private: + template <_AllocType, class...> + friend class basic_resource_ref; + + void* __object = nullptr; + const _Async_alloc_vtable* __static_vtable = nullptr; + +public: + // clang-format off + _LIBCUDACXX_TEMPLATE(class _Resource) + (requires ((_Alloc_type == _AllocType::_Default) && resource_with<_Resource, _Properties...>) // + || ((_Alloc_type == _AllocType::_Async) && async_resource_with<_Resource, _Properties...>)) // + _LIBCUDACXX_INLINE_VISIBILITY basic_resource_ref(_Resource& __res) noexcept + : _Filtered_vtable<_Properties...>(_Filtered_vtable<_Properties...>::template _Create<_Resource>()) + , __object(&__res) + , __static_vtable(&__alloc_vtable<_Alloc_type, _Resource>) + {} + + #if _LIBCUDACXX_STD_VER < 17 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #endif + _LIBCUDACXX_INLINE_VISIBILITY basic_resource_ref(basic_resource_ref<_Alloc_type, _OtherProperties...>& __ref) noexcept + : _Filtered_vtable<_Properties...>(__ref._Property_vtable<_Properties>::__property_fn...) + , __object(__ref.object) + , __static_vtable(__ref.__static_vtable) + {} + + #if _LIBCUDACXX_STD_VER < 17 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (_Alloc_type == _AllocType::_Default) + && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (_Alloc_type == _AllocType::_Default) + && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #endif + _LIBCUDACXX_INLINE_VISIBILITY basic_resource_ref(basic_resource_ref<_AllocType::_Async, _OtherProperties...>& __ref) noexcept + : _Filtered_vtable<_Properties...>(__ref._Property_vtable<_Properties>::__property_fn...) + , __object(__ref.object) + , __static_vtable(__ref.__static_vtable) + {} + + _LIBCUDACXX_INLINE_VISIBILITY void* allocate(size_t __bytes, size_t __alignment) { + return __static_vtable->__alloc_fn(__object, __bytes, __alignment); + } + + _LIBCUDACXX_INLINE_VISIBILITY void deallocate(void* _Ptr, size_t __bytes, size_t __alignment) { + __static_vtable->__dealloc_fn(__object, _Ptr, __bytes, __alignment); + } + + _LIBCUDACXX_TEMPLATE(_AllocType _Alloc_type2 = _Alloc_type) + (requires (_Alloc_type2 == _AllocType::_Async)) // + _LIBCUDACXX_INLINE_VISIBILITY void* allocate_async(size_t __bytes, size_t __alignment, cuda::stream_ref __stream) { + return __static_vtable->__async_alloc_fn(__object, __bytes, __alignment, __stream); + } + + _LIBCUDACXX_TEMPLATE(_AllocType _Alloc_type2 = _Alloc_type) + (requires (_Alloc_type2 == _AllocType::_Async)) // + _LIBCUDACXX_INLINE_VISIBILITY void deallocate_async( + void* _Ptr, size_t __bytes, size_t __alignment, cuda::stream_ref __stream) { + __static_vtable->__async_dealloc_fn(__object, _Ptr, __bytes, __alignment, __stream); + } + + #if _LIBCUDACXX_STD_VER < 17 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires(sizeof...(_Properties) == sizeof...(_OtherProperties)) + && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #endif + _LIBCUDACXX_INLINE_VISIBILITY bool operator==( + const basic_resource_ref<_Alloc_type, _OtherProperties...> &__right) const { + return (__static_vtable->__equal_fn == __right.__static_vtable->__equal_fn) // + && __static_vtable->__equal_fn(__object, __right.__object); + } + + #if _LIBCUDACXX_STD_VER < 17 + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires (sizeof...(_Properties) == sizeof...(_OtherProperties)) + && _CUDA_VSTD::conjunction_v<_CUDA_VSTD::bool_constant< + _CUDA_VSTD::_One_of<_Properties, _OtherProperties...>>...>) + #else + _LIBCUDACXX_TEMPLATE(class... _OtherProperties) + (requires(sizeof...(_Properties) == sizeof...(_OtherProperties)) + && (_CUDA_VSTD::_One_of<_Properties, _OtherProperties...> && ...)) + #endif + _LIBCUDACXX_INLINE_VISIBILITY 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...>) // + _LIBCUDACXX_INLINE_VISIBILITY 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...>) // + _LIBCUDACXX_INLINE_VISIBILITY 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 // +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