Use forward_property for property forwarding in resource_ref wrappers (#2328)

Replace the hand-rolled friend `get_property` templates in `cccl_resource_ref` and `cccl_async_resource_ref` with inheritance from `cuda::forward_property`. This delegates property forwarding to CCCL's own machinery, which correctly handles `dynamic_accessibility_property` ([NVIDIA/cccl#7727](NVIDIA/cccl#7727)) and any future properties without ambiguity.

Each wrapper now exposes `upstream_resource()` returning the inner `ResourceType`, as required by `forward_property` for stateful properties.

Tests add minimal `forward_property` adaptors using RMM resource refs as upstream, exercising the exact scenario that causes the ambiguity.

Note: this is a temporary solution for the `main` branch -- resolving #2323 / #2325 will remove this code on the `staging` branch while I continue working on CCCL MR migrations.

Closes #2322.

Authors:
  - Bradley Dice (https://github.com/bdice)

Approvers:
  - Rong Ou (https://github.com/rongou)

URL: #2328

Author: bdice

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 <typename ResourceType>
-class cccl_resource_ref {
+class cccl_resource_ref
+  : public cuda::forward_property<cccl_resource_ref<ResourceType>, 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 <typename Property>
-  friend auto constexpr get_property(cccl_resource_ref const& ref, Property prop) noexcept
-    -> decltype(get_property(std::declval<ResourceType const&>(), 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 <typename ResourceType>
-class cccl_async_resource_ref {
+class cccl_async_resource_ref
+  : public cuda::forward_property<cccl_async_resource_ref<ResourceType>, 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 <typename Property>
-  friend auto constexpr get_property(cccl_async_resource_ref const& ref, Property prop) noexcept
-    -> decltype(get_property(std::declval<ResourceType const&>(), 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.

cpp/tests/mr/resource_ref_conversion_tests.cpp

@@ -7,9 +7,11 @@
 
 #include <rmm/cuda_stream.hpp>
 #include <rmm/detail/error.hpp>
+#include <rmm/mr/cuda_memory_resource.hpp>
 #include <rmm/mr/pinned_host_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <cuda/memory_resource>
 #include <thrust/host_vector.h>
 
 #include <gtest/gtest.h>
@@ -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<forwarding_adaptor, rmm::device_async_resource_ref> {
+  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<forwarding_sync_adaptor, rmm::device_resource_ref> {
+  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, cuda::mr::device_accessible>,
+              "forwarding_adaptor must have device_accessible property via forward_property");
+
+static_assert(cuda::has_property<forwarding_sync_adaptor, cuda::mr::device_accessible>,
+              "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<cuda::mr::device_accessible> 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<cuda::mr::device_accessible> 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{});
+}