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..56da661cf5 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -0,0 +1,116 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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_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.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..4667c52c24 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -0,0 +1,90 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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_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..42c2c25bf7 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -0,0 +1,109 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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_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..43392dfbe3 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -0,0 +1,213 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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 vtable_tests { +template +constexpr bool + valid_size = sizeof(cuda::mr::_Filtered_vtable) == + (N * sizeof(void*)); + +// Ensure, that we have the proper size of the vtable +static_assert(valid_size<3, >, ""); +static_assert(valid_size<4, property_with_value >, ""); +static_assert( + valid_size<5, property_with_value, property_with_value >, ""); + +// Ensure, that we do not store duplicated entries multiple times +static_assert( + valid_size<4, property_with_value, property_with_value >, ""); +static_assert(valid_size<5, property_with_value, + property_with_value, property_with_value >, + ""); + +// Ensure that we do not try to store entries of valueless properties +static_assert(valid_size<3, property_without_value >, ""); +static_assert( + valid_size<4, property_with_value, property_without_value >, + ""); +static_assert( + valid_size<4, property_with_value, property_without_value, + property_with_value >, + ""); +} // namespace vtable_tests + +namespace resource_test { + +template +struct async_resource { + 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; + } +}; + +_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_any_async_resource_ref() { + async_resource input{42}; + cuda::mr::async_resource_ref ref{input}; + + // Ensure that we propery 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]); + } +} +__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_any_async_resource_ref() { + // Test some basic combinations of properties w/o state + test_any_async_resource_ref, + property_with_value >(); + test_any_async_resource_ref, + property_without_value >(); + test_any_async_resource_ref, + property_without_value >(); + + // Test duplicated properties + test_any_async_resource_ref, + property_with_value, + property_with_value >(); + + test_any_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**) { + resource_test::test_any_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..7ca50eedfd --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp @@ -0,0 +1,110 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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.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..54768a5804 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -0,0 +1,87 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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..89fdd3ccb0 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -0,0 +1,105 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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..1a88f9ead0 --- /dev/null +++ b/.upstream-tests/test/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -0,0 +1,206 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// 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 vtable_tests { +template +constexpr bool valid_size = + sizeof(cuda::mr::_Filtered_vtable) == (N * sizeof(void*)); + +// Ensure, that we have the proper size of the vtable +static_assert(valid_size<3, >, ""); +static_assert(valid_size<4, property_with_value >, ""); +static_assert( + valid_size<5, property_with_value, property_with_value >, ""); + +// Ensure, that we do not store duplicated entries multiple times +static_assert( + valid_size<4, property_with_value, property_with_value >, ""); +static_assert(valid_size<5, property_with_value, + property_with_value, property_with_value >, + ""); + +// Ensure that we do not try to store entries of valueless properties +static_assert(valid_size<3, property_without_value >, ""); +static_assert( + valid_size<4, property_with_value, property_without_value >, + ""); +static_assert( + valid_size<4, property_with_value, property_without_value, + property_with_value >, + ""); +} // namespace vtable_tests + +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; + } +}; + +_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_any_resource_ref() { + resource input{42}; + cuda::mr::resource_ref ref{input}; + + // Ensure that we propery 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_any_resource_ref() { + // Test some basic combinations of properties w/o state + test_any_resource_ref, + property_with_value >(); + test_any_resource_ref, + property_without_value >(); + test_any_resource_ref, + property_without_value >(); + + // Test duplicated properties + test_any_resource_ref, property_with_value, + property_with_value >(); + + test_any_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_any_resource_ref(); + return 0; +} diff --git a/include/cuda/memory_resource b/include/cuda/memory_resource index e211ac507a..8fb2ebbbd9 100644 --- a/include/cuda/memory_resource +++ b/include/cuda/memory_resource @@ -198,6 +198,278 @@ _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, +}; + +// 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); + } +}; +// clang-format on + +template +struct _Property_vtable +{ + using _PropertyFn = typename _Property::value_type (*)(void *); + _PropertyFn __property_fn = nullptr; + + _LIBCUDACXX_INLINE_VISIBILITY + constexpr _Property_vtable(_PropertyFn __property_fn_) noexcept + : __property_fn(__property_fn_) + {} +}; + +template <_AllocType _Alloc_type> +struct _Alloc_vtable +{ + using _AllocFn = _CUDA_VSTD::_If<_Alloc_type == _AllocType::_Default, + void *(*)(void *, size_t, size_t), + void *(*)(void *, size_t, size_t, cuda::stream_ref)>; + using _DeallocFn = _CUDA_VSTD::_If<_Alloc_type == _AllocType::_Default, + void (*)(void *, void *, size_t, size_t), + void (*)(void *, void *, size_t, size_t, cuda::stream_ref)>; + using _EqualFn = bool (*)(void *, void *); + + _AllocFn __alloc_fn = nullptr; + _DeallocFn __dealloc_fn = nullptr; + _EqualFn __equal_fn = nullptr; + + _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_) + {} +}; + +template <_AllocType _Alloc_type, class... _Properties> +struct _Resource_vtable + : public _Property_vtable<_Properties>... + , public _Alloc_vtable<_Alloc_type> +{ + template + _LIBCUDACXX_INLINE_VISIBILITY constexpr _Resource_vtable(_AllocFn __alloc_fn_, + _DeallocFn __dealloc_fn_, + _EqualFn __equal_fn_, + _PropertyFns... __property_fn_) noexcept + : _Alloc_vtable<_Alloc_type>(__alloc_fn_, __dealloc_fn_, __equal_fn_) + , _Property_vtable<_Properties>(__property_fn_)... + {} + + _LIBCUDACXX_TEMPLATE(class _Resource, _AllocType _Alloc_type2 = _Alloc_type) + (requires(_Alloc_type2 == _AllocType::_Default)) // + _LIBCUDACXX_INLINE_VISIBILITY static constexpr _Resource_vtable _Create() noexcept + { + return {&_Resource_vtable_builder::_Alloc<_Resource>, + &_Resource_vtable_builder::_Dealloc<_Resource>, + &_Resource_vtable_builder::_Equal<_Resource>, + &_Resource_vtable_builder::_Get_property<_Resource, _Properties>...}; + } + + _LIBCUDACXX_TEMPLATE(class _Resource, _AllocType _Alloc_type2 = _Alloc_type) + (requires(_Alloc_type2 == _AllocType::_Async)) // + _LIBCUDACXX_INLINE_VISIBILITY static constexpr _Resource_vtable _Create() noexcept + { + return {&_Resource_vtable_builder::_Alloc_async<_Resource>, + &_Resource_vtable_builder::_Dealloc_async<_Resource>, + &_Resource_vtable_builder::_Equal<_Resource>, + &_Resource_vtable_builder::_Get_property<_Resource, _Properties>...}; + } + + template + _LIBCUDACXX_INLINE_VISIBILITY constexpr typename _Property::value_type _Get_property(void *__object, + _Property) const noexcept + { + return _Property_vtable<_Property>::__property_fn(__object); + } +}; + +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...>; + + template <_AllocType _Alloc_type> + using _Vtable = _Resource_vtable<_Alloc_type, _Property, _Properties...>; +}; + +template <> +struct _Filtered<> +{ + using _Filtered_vtable = _Filtered<>; + + template + using _Append_property = _Filtered<_OtherPropery>; + + template <_AllocType _Alloc_type> + using _Vtable = _Resource_vtable<_Alloc_type>; +}; + +template <_AllocType _Alloc_type, class... _Properties> +using _Filtered_vtable = typename _Filtered<_Properties...>::_Filtered_vtable::_Vtable<_Alloc_type>; + +template <_AllocType _Alloc_type, class _Resource, class... _Properties> +__device__ _LIBCUDACXX_INLINE_VAR constexpr _Filtered_vtable<_Alloc_type, _Properties...> + __resource_vtable = _Filtered_vtable<_Alloc_type, _Properties...>::template _Create<_Resource>(); + +template <_AllocType _Alloc_type, class... _Properties> // +class basic_resource_ref +{ +private: + template <_AllocType, class...> + friend class basic_resource_ref; + + void *__object = nullptr; + const _Filtered_vtable<_Alloc_type, _Properties...> *__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 + : __object(&__res), __vtable(&__resource_vtable<_Alloc_type, _Resource, _Properties...>) + {} + + _LIBCUDACXX_TEMPLATE(_AllocType _Alloc_type2 = _Alloc_type) + (requires (_Alloc_type2 == _AllocType::_Default)) + _LIBCUDACXX_INLINE_VISIBILITY void* allocate(size_t __bytes, size_t __alignment) { + return __vtable->__alloc_fn(__object, __bytes, __alignment); + } + + _LIBCUDACXX_TEMPLATE(_AllocType _Alloc_type2 = _Alloc_type) + (requires (_Alloc_type2 == _AllocType::_Default)) // + _LIBCUDACXX_INLINE_VISIBILITY void deallocate(void* _Ptr, size_t __bytes, size_t __alignment) { + __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 __vtable->__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) { + __vtable->__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 (__vtable->__equal_fn == __right.__vtable->__equal_fn) // + && __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 __prop) noexcept { + return __res.__vtable->_Get_property(__res.__object, __prop); + } + // 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