diff --git a/cpp/include/rmm/detail/cccl_adaptors.hpp b/cpp/include/rmm/detail/cccl_adaptors.hpp index d8f6999e0..c7f930e8c 100644 --- a/cpp/include/rmm/detail/cccl_adaptors.hpp +++ b/cpp/include/rmm/detail/cccl_adaptors.hpp @@ -47,10 +47,15 @@ class cccl_async_resource_ref; * interface of the underlying type. It enables constructing resource refs from * device_memory_resource pointers by wrapping them in a device_memory_resource_view. * + * Inherits from cuda::forward_property to delegate property queries to the wrapped + * resource_ref, avoiding ambiguity with CCCL's default get_property overloads + * (e.g. for dynamic_accessibility_property, NVIDIA/cccl#7727). + * * @tparam ResourceType The underlying CCCL synchronous_resource_ref type */ template -class cccl_resource_ref { +class cccl_resource_ref + : public cuda::forward_property, ResourceType> { public: using wrapped_type = ResourceType; @@ -244,14 +249,11 @@ class cccl_resource_ref { } /** - * @brief Forwards a property query to the wrapped resource_ref. + * @brief Returns a const reference to the wrapped resource_ref. + * + * Required by cuda::forward_property to forward stateful property queries. */ - template - friend auto constexpr get_property(cccl_resource_ref const& ref, Property prop) noexcept - -> decltype(get_property(std::declval(), prop)) - { - return get_property(ref.ref_, prop); - } + [[nodiscard]] ResourceType const& upstream_resource() const noexcept { return ref_; } /** * @brief Attempts to get a property from the wrapped resource_ref. @@ -276,6 +278,10 @@ class cccl_resource_ref { * to avoid recursive constraint satisfaction issues with CCCL 3.2's basic_any-based * resource_ref types. It provides both synchronous and asynchronous allocation methods. * + * Inherits from cuda::forward_property to delegate property queries to the wrapped + * resource_ref, avoiding ambiguity with CCCL's default get_property overloads + * (e.g. for dynamic_accessibility_property, NVIDIA/cccl#7727). + * * @tparam ResourceType The underlying CCCL resource_ref type (async) */ // Suppress spurious warning about calling a __host__ function from __host__ __device__ context @@ -286,7 +292,8 @@ class cccl_resource_ref { #pragma nv_diag_suppress 20011 #endif template -class cccl_async_resource_ref { +class cccl_async_resource_ref + : public cuda::forward_property, ResourceType> { public: using wrapped_type = ResourceType; @@ -528,14 +535,11 @@ class cccl_async_resource_ref { } /** - * @brief Forwards a property query to the wrapped resource_ref. + * @brief Returns a const reference to the wrapped resource_ref. + * + * Required by cuda::forward_property to forward stateful property queries. */ - template - friend auto constexpr get_property(cccl_async_resource_ref const& ref, Property prop) noexcept - -> decltype(get_property(std::declval(), prop)) - { - return get_property(ref.ref_, prop); - } + [[nodiscard]] ResourceType const& upstream_resource() const noexcept { return ref_; } /** * @brief Attempts to get a property from the wrapped resource_ref. diff --git a/cpp/tests/mr/resource_ref_conversion_tests.cpp b/cpp/tests/mr/resource_ref_conversion_tests.cpp index dac965d14..348192275 100644 --- a/cpp/tests/mr/resource_ref_conversion_tests.cpp +++ b/cpp/tests/mr/resource_ref_conversion_tests.cpp @@ -7,9 +7,11 @@ #include #include +#include #include #include +#include #include #include @@ -218,3 +220,126 @@ TEST(ResourceRefConversionAllocator, VectorMove) auto vec2 = std::move(vec1); ASSERT_EQ(vec2[0], 42); } + +// -------------------------------------------------------------------------- +// Tests for forward_property adaptor with RMM resource_ref as upstream. +// +// This exercises the fix for https://github.com/rapidsai/rmm/issues/2322: +// an unconstrained friend get_property in cccl_async_resource_ref caused +// ambiguity with CCCL's default get_property for dynamic_accessibility_property. +// -------------------------------------------------------------------------- + +// A minimal adaptor that uses cuda::forward_property with an RMM async resource ref as upstream. +struct forwarding_adaptor + : cuda::forward_property { + explicit forwarding_adaptor(rmm::device_async_resource_ref upstream) : upstream_{upstream} {} + + rmm::device_async_resource_ref upstream_resource() const { return upstream_; } + + void* allocate(cuda::stream_ref stream, std::size_t bytes, std::size_t alignment) + { + return upstream_.allocate(stream, bytes, alignment); + } + void deallocate(cuda::stream_ref stream, + void* ptr, + std::size_t bytes, + std::size_t alignment) noexcept + { + upstream_.deallocate(stream, ptr, bytes, alignment); + } + void* allocate_sync(std::size_t bytes, std::size_t alignment) + { + return upstream_.allocate_sync(bytes, alignment); + } + void deallocate_sync(void* ptr, std::size_t bytes, std::size_t alignment) noexcept + { + upstream_.deallocate_sync(ptr, bytes, alignment); + } + + friend bool operator==(forwarding_adaptor const& lhs, forwarding_adaptor const& rhs) + { + return lhs.upstream_ == rhs.upstream_; + } + friend bool operator!=(forwarding_adaptor const& lhs, forwarding_adaptor const& rhs) + { + return !(lhs == rhs); + } + + private: + rmm::device_async_resource_ref upstream_; +}; + +// A minimal adaptor using forward_property with an RMM sync resource ref as upstream. +struct forwarding_sync_adaptor + : cuda::forward_property { + explicit forwarding_sync_adaptor(rmm::device_resource_ref upstream) : upstream_{upstream} {} + + rmm::device_resource_ref upstream_resource() const { return upstream_; } + + void* allocate_sync(std::size_t bytes, std::size_t alignment) + { + return upstream_.allocate_sync(bytes, alignment); + } + void deallocate_sync(void* ptr, std::size_t bytes, std::size_t alignment) noexcept + { + upstream_.deallocate_sync(ptr, bytes, alignment); + } + + friend bool operator==(forwarding_sync_adaptor const& lhs, forwarding_sync_adaptor const& rhs) + { + return lhs.upstream_ == rhs.upstream_; + } + friend bool operator!=(forwarding_sync_adaptor const& lhs, forwarding_sync_adaptor const& rhs) + { + return !(lhs == rhs); + } + + private: + rmm::device_resource_ref upstream_; +}; + +// Compile-time checks: verify that the forwarding adaptor satisfies resource concepts. +static_assert(cuda::has_property, + "forwarding_adaptor must have device_accessible property via forward_property"); + +static_assert(cuda::has_property, + "forwarding_sync_adaptor must have device_accessible property via forward_property"); + +// Type-erasing a forward_property adaptor into resource_ref triggers the ambiguity +// from issue #2322. If the constraint on cccl_async_resource_ref::get_property is missing, +// this will fail to compile when CCCL has dynamic_accessibility_property. +TEST(ForwardPropertyAdaptor, TypeEraseAsyncAdaptor) +{ + rmm::mr::cuda_memory_resource mr{}; + rmm::device_async_resource_ref upstream{mr}; + forwarding_adaptor adaptor{upstream}; + cuda::mr::resource_ref erased{adaptor}; + + rmm::cuda_stream stream{}; + void* ptr = erased.allocate(stream, 1024, 256); + ASSERT_NE(ptr, nullptr); + erased.deallocate(stream, ptr, 1024, 256); +} + +TEST(ForwardPropertyAdaptor, TypeEraseSyncAdaptor) +{ + rmm::mr::cuda_memory_resource mr{}; + rmm::device_resource_ref upstream{mr}; + forwarding_sync_adaptor adaptor{upstream}; + cuda::mr::synchronous_resource_ref erased{adaptor}; + + void* ptr = erased.allocate_sync(1024); + ASSERT_NE(ptr, nullptr); + erased.deallocate_sync(ptr, 1024); +} + +// Verify that get_property still works correctly through the forwarding adaptor. +TEST(ForwardPropertyAdaptor, GetPropertyDeviceAccessible) +{ + rmm::mr::cuda_memory_resource mr{}; + rmm::device_async_resource_ref upstream{mr}; + forwarding_adaptor adaptor{upstream}; + + // Should compile and not throw - device_accessible is a stateless property + get_property(adaptor, cuda::mr::device_accessible{}); +}