diff --git a/cpp/benchmarks/utilities/simulated_memory_resource.hpp b/cpp/benchmarks/utilities/simulated_memory_resource.hpp index ff5254b83..0a66ffb43 100644 --- a/cpp/benchmarks/utilities/simulated_memory_resource.hpp +++ b/cpp/benchmarks/utilities/simulated_memory_resource.hpp @@ -6,7 +6,6 @@ #include #include -#include #include #include @@ -112,8 +111,8 @@ class simulated_memory_resource final { bool operator==(simulated_memory_resource const&) const noexcept { return true; } bool operator!=(simulated_memory_resource const&) const noexcept { return false; } - RMM_CONSTEXPR_FRIEND void get_property(simulated_memory_resource const&, - cuda::mr::device_accessible) noexcept + constexpr friend void get_property(simulated_memory_resource const&, + cuda::mr::device_accessible) noexcept { } diff --git a/cpp/include/rmm/detail/cccl_adaptors.hpp b/cpp/include/rmm/detail/cccl_adaptors.hpp index d2b33ec4d..b8ab9ae37 100644 --- a/cpp/include/rmm/detail/cccl_adaptors.hpp +++ b/cpp/include/rmm/detail/cccl_adaptors.hpp @@ -6,11 +6,8 @@ #include #include -#include -#include #include -#include #include #include @@ -83,11 +80,11 @@ inline constexpr bool is_cccl_async_resource_ref_v> = /** * @brief A wrapper around CCCL synchronous_resource_ref that adds compatibility with - * device_memory_resource pointers. + * shared_resource-derived types. * * This class uses composition to wrap a CCCL resource_ref type and provides the full * interface of the underlying type. It enables constructing resource refs from - * device_memory_resource pointers by wrapping them in a device_memory_resource_view. + * shared_resource-derived types by casting to the shared_resource base. * * @tparam ResourceType The underlying CCCL synchronous_resource_ref type */ @@ -100,77 +97,28 @@ class cccl_resource_ref { template friend class cccl_resource_ref; - /** - * @brief Constructs a resource reference from a raw `device_memory_resource` pointer. - * - * This constructor enables compatibility with CCCL 3.2 by wrapping the pointer in a - * `device_memory_resource_view`, which is copyable unlike the virtual base class. - * - * @param ptr Non-null pointer to a `device_memory_resource` - */ - cccl_resource_ref(rmm::mr::device_memory_resource* ptr) : view_{ptr}, ref_{*view_} {} - - /** - * @brief Constructs a resource reference from a `device_memory_resource` reference. - * - * This constructor enables compatibility with CCCL 3.2 by wrapping the address in a - * `device_memory_resource_view`, which is copyable unlike the virtual base class. - * - * @param res Reference to a `device_memory_resource` - */ - cccl_resource_ref(rmm::mr::device_memory_resource& res) : view_{&res}, ref_{*view_} {} - /** * @brief Constructs a resource reference from a CCCL resource_ref directly. * - * This constructor enables interoperability with CCCL 3.2 resource_ref types, - * allowing RMM resource_ref types to be constructed from CCCL resource_ref types. - * * @param ref A CCCL resource_ref of the appropriate type */ template - cccl_resource_ref(cuda::mr::synchronous_resource_ref const& ref) - : view_{cuda::std::nullopt}, ref_{ref} + cccl_resource_ref(cuda::mr::synchronous_resource_ref const& ref) : ref_{ref} { } /** * @brief Constructs a resource reference from a CCCL resource_ref directly (move). * - * This constructor enables interoperability with CCCL 3.2 resource_ref types, - * allowing RMM resource_ref types to be constructed from CCCL resource_ref types - * using move semantics. - * * @param ref A CCCL resource_ref of the appropriate type */ template - cccl_resource_ref(cuda::mr::synchronous_resource_ref&& ref) - : view_{cuda::std::nullopt}, ref_{std::move(ref)} + cccl_resource_ref(cuda::mr::synchronous_resource_ref&& ref) : ref_{std::move(ref)} { } - /** - * @brief Copy constructor that properly reconstructs the ref to point to the new view. - * - * If the view is present (e.g., when constructed from device_memory_resource*), we reconstruct - * the ref from our local view. Otherwise, we copy the ref directly. - */ - cccl_resource_ref(cccl_resource_ref const& other) - : view_{other.view_}, ref_{view_.has_value() ? ResourceType{*view_} : other.ref_} - { - } - - /** - * @brief Move constructor that properly reconstructs the ref to point to the new view. - * - * If the view is present (e.g., when constructed from device_memory_resource*), we reconstruct - * the ref from our local view. Otherwise, we move the ref directly. - */ - cccl_resource_ref(cccl_resource_ref&& other) noexcept - : view_{std::move(other.view_)}, - ref_{view_.has_value() ? ResourceType{*view_} : std::move(other.ref_)} - { - } + cccl_resource_ref(cccl_resource_ref const&) = default; + cccl_resource_ref(cccl_resource_ref&&) noexcept = default; /** * @brief Conversion constructor from a cccl_resource_ref with a convertible ResourceType. @@ -183,8 +131,7 @@ class cccl_resource_ref { * @param other The source resource_ref to convert from */ template - cccl_resource_ref(cccl_resource_ref const& other) - : view_{other.view_}, ref_{view_.has_value() ? ResourceType{*view_} : ResourceType{other.ref_}} + cccl_resource_ref(cccl_resource_ref const& other) : ref_{other.ref_} { } @@ -200,7 +147,7 @@ class cccl_resource_ref { template ::value>* = nullptr> cccl_resource_ref(OtherResourceType& other) - : view_{}, ref_{ResourceType{shared_resource_cast::cast(other)}} + : ref_{ResourceType{shared_resource_cast::cast(other)}} { } @@ -208,11 +155,11 @@ class cccl_resource_ref { * @brief Construct a ref from a resource. * * This constructor accepts CCCL resource types but NOT CCCL resource_ref types, - * our own wrapper types, device_memory_resource derived types, or - * shared_resource-derived types (handled by dedicated constructor above). + * our own wrapper types, or shared_resource-derived types (handled by dedicated + * constructor above). * The exclusions are checked FIRST to prevent recursive constraint satisfaction. * - * @tparam OtherResourceType A CCCL resource type (not a resource_ref, wrapper, DMR, + * @tparam OtherResourceType A CCCL resource type (not a resource_ref, wrapper, * or shared_resource) * @param other The resource to construct a ref from */ @@ -225,42 +172,13 @@ class cccl_resource_ref { not is_cccl_resource_ref_v> and not is_cccl_async_resource_ref_v> and not shared_resource_cast::value and - not std::is_base_of_v> and cuda::mr::synchronous_resource>* = nullptr> - cccl_resource_ref(OtherResourceType& other) : view_{}, ref_{ResourceType{other}} + cccl_resource_ref(OtherResourceType& other) : ref_{ResourceType{other}} { } - /** - * @brief Copy assignment operator. - * - * If the view is present, we reconstruct the ref from our local view. - * Otherwise, we copy the ref directly. - */ - cccl_resource_ref& operator=(cccl_resource_ref const& other) - { - if (this != std::addressof(other)) { - view_ = other.view_; - ref_ = view_.has_value() ? ResourceType{*view_} : other.ref_; - } - return *this; - } - - /** - * @brief Move assignment operator. - * - * If the view is present, we reconstruct the ref from our local view. - * Otherwise, we move the ref directly. - */ - cccl_resource_ref& operator=(cccl_resource_ref&& other) noexcept - { - if (this != std::addressof(other)) { - view_ = std::move(other.view_); - ref_ = view_.has_value() ? ResourceType{*view_} : std::move(other.ref_); - } - return *this; - } + cccl_resource_ref& operator=(cccl_resource_ref const&) = default; + cccl_resource_ref& operator=(cccl_resource_ref&&) noexcept = default; void* allocate_sync(std::size_t bytes) { return ref_.allocate_sync(bytes); } @@ -324,13 +242,12 @@ class cccl_resource_ref { } protected: - cuda::std::optional view_; ResourceType ref_; }; /** * @brief A wrapper around CCCL resource_ref (async) that adds compatibility with - * device_memory_resource pointers. + * shared_resource-derived types. * * This class is a standalone implementation (not inheriting from cccl_resource_ref) * to avoid recursive constraint satisfaction issues with CCCL 3.2's basic_any-based @@ -354,52 +271,23 @@ class cccl_async_resource_ref { template friend class cccl_async_resource_ref; - /** - * @brief Constructs a resource reference from a raw `device_memory_resource` pointer. - * - * This constructor enables compatibility with CCCL 3.2 by wrapping the pointer in a - * `device_memory_resource_view`, which is copyable unlike the virtual base class. - * - * @param ptr Non-null pointer to a `device_memory_resource` - */ - cccl_async_resource_ref(rmm::mr::device_memory_resource* ptr) : view_{ptr}, ref_{*view_} {} - - /** - * @brief Constructs a resource reference from a `device_memory_resource` reference. - * - * This constructor enables compatibility with CCCL 3.2 by wrapping the address in a - * `device_memory_resource_view`, which is copyable unlike the virtual base class. - * - * @param res Reference to a `device_memory_resource` - */ - cccl_async_resource_ref(rmm::mr::device_memory_resource& res) : view_{&res}, ref_{*view_} {} - /** * @brief Constructs a resource reference from a CCCL resource_ref directly. * - * This constructor enables interoperability with CCCL 3.2 resource_ref types, - * allowing RMM resource_ref types to be constructed from CCCL resource_ref types. - * * @param ref A CCCL resource_ref of the appropriate type */ template - cccl_async_resource_ref(cuda::mr::resource_ref const& ref) - : view_{cuda::std::nullopt}, ref_{ref} + cccl_async_resource_ref(cuda::mr::resource_ref const& ref) : ref_{ref} { } /** * @brief Constructs a resource reference from a CCCL resource_ref directly (move). * - * This constructor enables interoperability with CCCL 3.2 resource_ref types, - * allowing RMM resource_ref types to be constructed from CCCL resource_ref types - * using move semantics. - * * @param ref A CCCL resource_ref of the appropriate type */ template - cccl_async_resource_ref(cuda::mr::resource_ref&& ref) - : view_{cuda::std::nullopt}, ref_{std::move(ref)} + cccl_async_resource_ref(cuda::mr::resource_ref&& ref) : ref_{std::move(ref)} { } @@ -412,33 +300,12 @@ class cccl_async_resource_ref { * @param res A CCCL any_resource to reference */ template - cccl_async_resource_ref(cuda::mr::any_resource& res) - : view_{cuda::std::nullopt}, ref_{res} + cccl_async_resource_ref(cuda::mr::any_resource& res) : ref_{res} { } - /** - * @brief Copy constructor that properly reconstructs the ref to point to the new view. - * - * If the view is present (e.g., when constructed from device_memory_resource*), we reconstruct - * the ref from our local view. Otherwise, we copy the ref directly. - */ - cccl_async_resource_ref(cccl_async_resource_ref const& other) - : view_{other.view_}, ref_{view_.has_value() ? ResourceType{*view_} : other.ref_} - { - } - - /** - * @brief Move constructor that properly reconstructs the ref to point to the new view. - * - * If the view is present (e.g., when constructed from device_memory_resource*), we reconstruct - * the ref from our local view. Otherwise, we move the ref directly. - */ - cccl_async_resource_ref(cccl_async_resource_ref&& other) noexcept - : view_{std::move(other.view_)}, - ref_{view_.has_value() ? ResourceType{*view_} : std::move(other.ref_)} - { - } + cccl_async_resource_ref(cccl_async_resource_ref const&) = default; + cccl_async_resource_ref(cccl_async_resource_ref&&) noexcept = default; /** * @brief Conversion constructor from a cccl_async_resource_ref with a convertible ResourceType. @@ -447,16 +314,12 @@ class cccl_async_resource_ref { * where the source type has a superset of properties compared to the target type. * The underlying CCCL resource_ref types handle the actual property compatibility check. * - * IMPORTANT: This constructor must copy the view_ from the source to preserve the - * device_memory_resource pointer. Without this, the converted resource_ref will have - * an empty view_, causing corrupt pointer dereferences during deallocation. - * * @tparam OtherResourceType A CCCL async resource_ref type that is convertible to ResourceType * @param other The source async resource_ref to convert from */ template cccl_async_resource_ref(cccl_async_resource_ref const& other) - : view_{other.view_}, ref_{view_.has_value() ? ResourceType{*view_} : ResourceType{other.ref_}} + : ref_{other.ref_} { } @@ -472,7 +335,7 @@ class cccl_async_resource_ref { template ::value>* = nullptr> cccl_async_resource_ref(OtherResourceType& other) - : view_{}, ref_{ResourceType{shared_resource_cast::cast(other)}} + : ref_{ResourceType{shared_resource_cast::cast(other)}} { } @@ -480,12 +343,12 @@ class cccl_async_resource_ref { * @brief Construct a ref from a resource. * * This constructor accepts CCCL resource types but NOT CCCL resource_ref types, - * our own wrapper types, any_resource types, device_memory_resource derived types, - * or shared_resource-derived types (handled by dedicated constructor above). + * our own wrapper types, any_resource types, or shared_resource-derived types + * (handled by dedicated constructor above). * The exclusions are checked FIRST to prevent recursive constraint satisfaction. * * @tparam OtherResourceType A CCCL resource type (not a resource_ref, wrapper, any_resource, - * DMR, or shared_resource) + * or shared_resource) * @param other The resource to construct a ref from */ template < @@ -498,42 +361,13 @@ class cccl_async_resource_ref { not is_cccl_resource_ref_v> and not is_cccl_async_resource_ref_v> and not shared_resource_cast::value and - not std::is_base_of_v> and cuda::mr::resource>* = nullptr> - cccl_async_resource_ref(OtherResourceType& other) : view_{}, ref_{ResourceType{other}} + cccl_async_resource_ref(OtherResourceType& other) : ref_{ResourceType{other}} { } - /** - * @brief Copy assignment operator. - * - * If the view is present, we reconstruct the ref from our local view. - * Otherwise, we copy the ref directly. - */ - cccl_async_resource_ref& operator=(cccl_async_resource_ref const& other) - { - if (this != std::addressof(other)) { - view_ = other.view_; - ref_ = view_.has_value() ? ResourceType{*view_} : other.ref_; - } - return *this; - } - - /** - * @brief Move assignment operator. - * - * If the view is present, we reconstruct the ref from our local view. - * Otherwise, we move the ref directly. - */ - cccl_async_resource_ref& operator=(cccl_async_resource_ref&& other) noexcept - { - if (this != std::addressof(other)) { - view_ = std::move(other.view_); - ref_ = view_.has_value() ? ResourceType{*view_} : std::move(other.ref_); - } - return *this; - } + cccl_async_resource_ref& operator=(cccl_async_resource_ref const&) = default; + cccl_async_resource_ref& operator=(cccl_async_resource_ref&&) noexcept = default; // Synchronous allocation methods (delegated to the underlying ref) void* allocate_sync(std::size_t bytes) { return ref_.allocate_sync(bytes); } @@ -632,12 +466,10 @@ class cccl_async_resource_ref { template operator cuda::mr::any_resource() const { - if (view_.has_value()) { return cuda::mr::any_resource{*view_}; } return cuda::mr::any_resource{ref_}; } protected: - cuda::std::optional view_; ResourceType ref_; }; #ifdef __CUDACC__ diff --git a/cpp/include/rmm/mr/arena_memory_resource.hpp b/cpp/include/rmm/mr/arena_memory_resource.hpp index d91da19df..e0611d3c1 100644 --- a/cpp/include/rmm/mr/arena_memory_resource.hpp +++ b/cpp/include/rmm/mr/arena_memory_resource.hpp @@ -24,7 +24,7 @@ namespace mr { /** * @brief A suballocator that emphasizes fragmentation avoidance and scalable concurrency support. * - * Allocation (do_allocate()) and deallocation (do_deallocate()) are thread-safe. Also, + * Allocation and deallocation are thread-safe. Also, * this class is compatible with CUDA per-thread default stream. * * GPU memory is divided into a global arena, per-thread arenas for default streams, and per-stream diff --git a/cpp/include/rmm/mr/callback_memory_resource.hpp b/cpp/include/rmm/mr/callback_memory_resource.hpp index e8b40b1c9..f174ff676 100644 --- a/cpp/include/rmm/mr/callback_memory_resource.hpp +++ b/cpp/include/rmm/mr/callback_memory_resource.hpp @@ -28,12 +28,10 @@ namespace mr { * * * Returns a pointer to an allocation of at least `bytes` usable immediately on * `stream`. The stream-ordered behavior requirements are identical to - * `device_memory_resource::allocate`. + * `allocate`. * - * * This signature is compatible with `do_allocate` but adds the extra function - * parameter `arg`. The `arg` is provided to the constructor of the - * `callback_memory_resource` and will be forwarded along to every invocation - * of the callback function. + * * The `arg` is provided to the constructor of the `callback_memory_resource` + * and will be forwarded along to every invocation of the callback function. */ using allocate_callback_t = std::function; @@ -46,12 +44,10 @@ using allocate_callback_t = std::function; diff --git a/cpp/include/rmm/mr/cuda_async_view_memory_resource.hpp b/cpp/include/rmm/mr/cuda_async_view_memory_resource.hpp index 8c505b442..e2aa54b45 100644 --- a/cpp/include/rmm/mr/cuda_async_view_memory_resource.hpp +++ b/cpp/include/rmm/mr/cuda_async_view_memory_resource.hpp @@ -68,8 +68,6 @@ class cuda_async_view_memory_resource final { cuda_async_view_memory_resource& operator=(cuda_async_view_memory_resource&&) = default; ///< @default_move_assignment{cuda_async_view_memory_resource} - // -- CCCL memory resource interface (hides device_memory_resource versions) -- - /** * @brief Allocates memory of size at least \p bytes. * diff --git a/cpp/include/rmm/mr/cuda_memory_resource.hpp b/cpp/include/rmm/mr/cuda_memory_resource.hpp index 0819c1ecf..41a5c98bd 100644 --- a/cpp/include/rmm/mr/cuda_memory_resource.hpp +++ b/cpp/include/rmm/mr/cuda_memory_resource.hpp @@ -34,8 +34,6 @@ class cuda_memory_resource final { cuda_memory_resource& operator=(cuda_memory_resource&&) = default; ///< @default_move_assignment{cuda_memory_resource} - // -- CCCL memory resource interface (hides device_memory_resource versions) -- - /** * @brief Allocates memory of size at least \p bytes. * diff --git a/cpp/include/rmm/mr/detail/arena.hpp b/cpp/include/rmm/mr/detail/arena.hpp index ec0dfabf1..3f6648efc 100644 --- a/cpp/include/rmm/mr/detail/arena.hpp +++ b/cpp/include/rmm/mr/detail/arena.hpp @@ -480,8 +480,6 @@ inline auto max_free_size(std::set const& superblocks) * * The global arena is a shared memory pool from which other arenas allocate superblocks. * - * @tparam Upstream Memory resource to use for allocating the arena. Implements - * rmm::mr::device_memory_resource interface. */ class global_arena final { public: @@ -778,8 +776,6 @@ class global_arena final { * An arena is a per-thread or per-non-default-stream memory pool. It allocates * superblocks from the global arena, and returns them when the superblocks become empty. * - * @tparam Upstream Memory resource to use for allocating the global arena. Implements - * rmm::mr::device_memory_resource interface. */ class arena { public: @@ -957,8 +953,6 @@ class arena { * * This is useful when a thread is about to terminate, and it contains a per-thread arena. * - * @tparam Upstream Memory resource to use for allocating the global arena. Implements - * rmm::mr::device_memory_resource interface. */ class arena_cleaner { public: diff --git a/cpp/include/rmm/mr/detail/device_memory_resource_view.hpp b/cpp/include/rmm/mr/detail/device_memory_resource_view.hpp deleted file mode 100644 index e24834d20..000000000 --- a/cpp/include/rmm/mr/detail/device_memory_resource_view.hpp +++ /dev/null @@ -1,180 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include - -#include - -#include -#include - -namespace RMM_NAMESPACE { -namespace mr::detail { - -/** - * @brief A copyable view wrapping a `device_memory_resource*` pointer. - * - * This class serves as a temporary bridge to enable compatibility with CCCL 3.2's memory resource - * design, which requires resource types to be copyable but no longer accepts raw pointers directly. - * Since `device_memory_resource` is a virtual base class that cannot be copied, this view provides - * a copyable wrapper around a `device_memory_resource*` pointer. - * - * This is an internal implementation detail and should not be used directly by users. It will be - * removed once RMM fully migrates away from the `device_memory_resource` virtual base class. - * - * @note This class does NOT manage the lifetime of the wrapped pointer. The caller is responsible - * for ensuring the pointed-to resource remains valid for the lifetime of this view. - */ -class device_memory_resource_view { - public: - /** - * @brief Constructs a view wrapping the given `device_memory_resource` pointer. - * - * @throws rmm::logic_error if `ptr` is null - * - * @param ptr Non-null pointer to a `device_memory_resource` - */ - device_memory_resource_view(device_memory_resource* ptr) : resource_ptr_{ptr} - { - RMM_EXPECTS(ptr != nullptr, "device_memory_resource_view cannot wrap a null pointer"); - } - - /** - * @brief Synchronously allocates memory of size at least `bytes`. - * - * @param bytes The size of the allocation - * @param alignment The alignment of the allocation - * @return void* Pointer to the newly allocated memory - */ - [[nodiscard]] void* allocate_sync(std::size_t bytes, - std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) - { - return resource_ptr_->allocate_sync(bytes, alignment); - } - - /** - * @brief Synchronously deallocates memory pointed to by `ptr`. - * - * @param ptr Pointer to be deallocated - * @param bytes The size in bytes of the allocation - * @param alignment The alignment of the allocation - */ - void deallocate_sync(void* ptr, - std::size_t bytes, - std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept - { - resource_ptr_->deallocate_sync(ptr, bytes, alignment); - } - - /** - * @brief Asynchronously allocates memory of size at least `bytes` on the specified stream. - * - * @param stream The stream on which to perform the allocation - * @param bytes The size of the allocation - * @param alignment The alignment of the allocation - * @return void* Pointer to the newly allocated memory - */ - [[nodiscard]] void* allocate(cuda_stream_view stream, - std::size_t bytes, - std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) - { - return resource_ptr_->allocate(stream, bytes, alignment); - } - - /** - * @brief Asynchronously deallocates memory pointed to by `ptr` on the specified stream. - * - * @param stream The stream on which to perform the deallocation - * @param ptr Pointer to be deallocated - * @param bytes The size in bytes of the allocation - * @param alignment The alignment of the allocation - */ - void deallocate(cuda_stream_view stream, - void* ptr, - std::size_t bytes, - std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept - { - resource_ptr_->deallocate(stream, ptr, bytes, alignment); - } - - /** - * @brief Returns the wrapped `device_memory_resource` pointer. - * - * @return device_memory_resource* The wrapped pointer - */ - [[nodiscard]] device_memory_resource* get() const noexcept { return resource_ptr_; } - - /** - * @brief Compares this view to another for equality. - * - * Two views are equal if they wrap pointers to resources that compare equal. - * - * @param other The other view to compare to - * @return true If the wrapped resources are equivalent - */ - [[nodiscard]] bool operator==(device_memory_resource_view const& other) const noexcept - { - // If both pointers are null, they're equal - if (resource_ptr_ == nullptr && other.resource_ptr_ == nullptr) { return true; } - // If only one is null, they're not equal - if (resource_ptr_ == nullptr || other.resource_ptr_ == nullptr) { return false; } - // Otherwise, compare the resources they point to - return resource_ptr_->is_equal(*other.resource_ptr_); - } - - /** - * @brief Compares this view to another for inequality. - * - * @param other The other view to compare to - * @return true If the wrapped resources are not equivalent - */ - [[nodiscard]] bool operator!=(device_memory_resource_view const& other) const noexcept - { - return !(*this == other); - } - - /** - * @brief Enables the `cuda::mr::device_accessible` property - * - * This property declares that the wrapped `device_memory_resource` provides device accessible - * memory. - */ - friend void get_property(device_memory_resource_view const&, cuda::mr::device_accessible) noexcept - { - } - - /** - * @brief Enables the `cuda::mr::host_accessible` property - * - * This property declares that the wrapped `device_memory_resource` may provide host accessible - * memory. This is needed for resources like pinned_host_memory_resource that are both host and - * device accessible. - */ - friend void get_property(device_memory_resource_view const&, cuda::mr::host_accessible) noexcept - { - } - - private: - device_memory_resource* resource_ptr_{nullptr}; -}; - -// Static assertions to verify that device_memory_resource_view satisfies CCCL resource concepts -static_assert(cuda::mr::resource_with, - "device_memory_resource_view must satisfy async resource concept"); -static_assert( - cuda::mr::synchronous_resource_with, - "device_memory_resource_view must satisfy synchronous resource concept"); - -// Verify copyability - required for resource_ref construction -static_assert(cuda::std::copyable, - "device_memory_resource_view must satisfy copyable concept"); -static_assert(cuda::std::copy_constructible, - "device_memory_resource_view must be copy constructible"); - -} // namespace mr::detail -} // namespace RMM_NAMESPACE diff --git a/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp b/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp index 99f3bcd5b..e8de8b65e 100644 --- a/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp +++ b/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp @@ -76,39 +76,39 @@ class stream_ordered_memory_resource : public crtp { stream_ordered_memory_resource& operator=(stream_ordered_memory_resource&&) = delete; /** - * @brief Allocates memory of size at least `size` bytes. + * @brief Allocates memory of size at least `bytes` bytes. * * The returned pointer has at least 256B alignment. * * @throws `std::bad_alloc` if the requested allocation could not be fulfilled * * @param stream The stream in which to order this allocation - * @param size The size in bytes of the allocation + * @param bytes The size in bytes of the allocation * @param alignment Unused; alignment is always at least `CUDA_ALLOCATION_ALIGNMENT` * @return void* Pointer to the newly allocated memory */ - void* allocate(cuda::stream_ref stream, std::size_t size, std::size_t /*alignment*/) + void* allocate(cuda::stream_ref stream, std::size_t bytes, std::size_t /*alignment*/) { auto const strm = cuda_stream_view{stream}; - RMM_LOG_TRACE("[A][stream %s][%zuB]", rmm::detail::format_stream(strm), size); + RMM_LOG_TRACE("[A][stream %s][%zuB]", rmm::detail::format_stream(strm), bytes); - if (size == 0) { return nullptr; } + if (bytes == 0) { return nullptr; } lock_guard lock(mtx_); auto stream_event = get_event(strm); - size = rmm::align_up(size, rmm::CUDA_ALLOCATION_ALIGNMENT); - RMM_EXPECTS(size <= this->underlying().get_maximum_allocation_size(), + bytes = rmm::align_up(bytes, rmm::CUDA_ALLOCATION_ALIGNMENT); + RMM_EXPECTS(bytes <= this->underlying().get_maximum_allocation_size(), std::string("Maximum allocation size exceeded (failed to allocate ") + - rmm::detail::format_bytes(size) + ")", + rmm::detail::format_bytes(bytes) + ")", rmm::out_of_memory); - auto const block = this->underlying().get_block(size, stream_event); + auto const block = this->underlying().get_block(bytes, stream_event); RMM_LOG_TRACE("[A][stream %s][%zuB][%p]", rmm::detail::format_stream(stream_event.stream), - size, + bytes, block.pointer()); log_summary_trace(); @@ -121,29 +121,29 @@ class stream_ordered_memory_resource : public crtp { * * @param stream The stream in which to order this deallocation * @param ptr Pointer to be deallocated - * @param size The size in bytes of the allocation to deallocate + * @param bytes The size in bytes of the allocation to deallocate * @param alignment Unused */ void deallocate(cuda::stream_ref stream, void* ptr, - std::size_t size, + std::size_t bytes, std::size_t /*alignment*/) noexcept { auto const strm = cuda_stream_view{stream}; - RMM_LOG_TRACE("[D][stream %s][%zuB][%p]", rmm::detail::format_stream(strm), size, ptr); + RMM_LOG_TRACE("[D][stream %s][%zuB][%p]", rmm::detail::format_stream(strm), bytes, ptr); - if (size == 0 || ptr == nullptr) { return; } + if (bytes == 0 || ptr == nullptr) { return; } lock_guard lock(mtx_); auto stream_event = get_event(strm); - size = rmm::align_up(size, rmm::CUDA_ALLOCATION_ALIGNMENT); - auto const block = this->underlying().free_block(ptr, size); + bytes = rmm::align_up(bytes, rmm::CUDA_ALLOCATION_ALIGNMENT); + auto const block = this->underlying().free_block(ptr, bytes); // TODO: cudaEventRecord has significant overhead on deallocations. For the non-PTDS case - // we may be able to delay recording the event in some situations. But using events rather than - // streams allows stealing from deleted streams. + // we may be able to delay recording the event in some situations. But using events rather + // than streams allows stealing from deleted streams. RMM_ASSERT_CUDA_SUCCESS(cudaEventRecord(stream_event.event, strm.value())); stream_free_blocks_[stream_event].insert(block); diff --git a/cpp/include/rmm/mr/device_memory_resource.hpp b/cpp/include/rmm/mr/device_memory_resource.hpp deleted file mode 100644 index 36779eb18..000000000 --- a/cpp/include/rmm/mr/device_memory_resource.hpp +++ /dev/null @@ -1,288 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ -#pragma once - -#include -#include -#include -#include -#include - -#include - -#include -#include - -namespace RMM_NAMESPACE { -namespace mr { -/** - * @addtogroup memory_resources - * @{ - * @file - */ - -/** - * @brief Base class for all librmm device memory allocation. - * - * This class serves as the interface that all custom device memory - * implementations must satisfy. - * - * There are two private, pure virtual functions that all derived classes must implement: - *`do_allocate` and `do_deallocate`. Optionally, derived classes may also override `is_equal`. By - * default, `is_equal` simply performs an identity comparison. - * - * The public, non-virtual functions `allocate`, `deallocate`, and `is_equal` simply call the - * private virtual functions. The reason for this is to allow implementing shared, default behavior - * in the base class. For example, the base class' `allocate` function may log every allocation, no - * matter what derived class implementation is used. - * - * The `allocate` and `deallocate` APIs and implementations provide stream-ordered memory - * allocation. This allows optimizations such as re-using memory deallocated on the same stream - * without the overhead of stream synchronization. - * - * A call to `allocate(bytes, stream_a)` (on any derived class) returns a pointer that is valid to - * use on `stream_a`. Using the memory on a different stream (say `stream_b`) is Undefined Behavior - * unless the two streams are first synchronized, for example by using - * `cudaStreamSynchronize(stream_a)` or by recording a CUDA event on `stream_a` and then - * calling `cudaStreamWaitEvent(stream_b, event)`. - * - * The stream specified to deallocate() should be a stream on which it is valid to use the - * deallocated memory immediately for another allocation. Typically this is the stream on which the - * allocation was *last* used before the call to deallocate(). The passed stream may be used - * internally by a device_memory_resource for managing available memory with minimal - * synchronization, and it may also be synchronized at a later time, for example using a call to - * `cudaStreamSynchronize()`. - * - * For this reason, it is Undefined Behavior to destroy a CUDA stream that is passed to - * deallocate(). If the stream on which the allocation was last used has been destroyed before - * calling deallocate() or it is known that it will be destroyed, it is likely better to synchronize - * the stream (before destroying it) and then pass a different stream to deallocate() (e.g. the - * default stream). - * - * A device_memory_resource should only be used when the active CUDA device is the same device - * that was active when the device_memory_resource was created. Otherwise behavior is undefined. - * - * Creating a device_memory_resource for each device requires care to set the current device - * before creating each resource, and to maintain the lifetime of the resources as long as they - * are set as per-device resources. Here is an example loop that creates `unique_ptr`s to - * pool_memory_resource objects for each device and sets them as the per-device resource for that - * device. - * - * @code{.cpp} - * using pool_mr = rmm::mr::pool_memory_resource; - * std::vector> per_device_pools; - * for(int i = 0; i < N; ++i) { - * cudaSetDevice(i); - * // Note: for brevity, omitting creation of upstream and computing initial_size - * per_device_pools.push_back(std::make_unique(upstream, initial_size)); - * set_per_device_resource(cuda_device_id{i}, &per_device_pools.back()); - * } - * @endcode - */ -class device_memory_resource { - public: - device_memory_resource() = default; - virtual ~device_memory_resource() = default; - device_memory_resource(device_memory_resource const&) = default; ///< @default_copy_constructor - device_memory_resource(device_memory_resource&&) noexcept = - default; ///< @default_move_constructor - device_memory_resource& operator=(device_memory_resource const&) = - default; ///< @default_copy_assignment{device_memory_resource} - device_memory_resource& operator=(device_memory_resource&&) noexcept = - default; ///< @default_move_assignment{device_memory_resource} - - /** - * @brief Allocates memory of size at least \p bytes. - * - * The returned pointer will have 256 byte alignment regardless of the value - * of alignment. Higher alignments must use the aligned_resource_adaptor. - * - * The returned pointer is immediately valid on all streams. - * - * @throws rmm::bad_alloc When the requested `bytes` cannot be allocated. - * - * @param bytes The size of the allocation - * @param alignment The alignment of the allocation (see notes above) - * @return void* Pointer to the newly allocated memory - */ - void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) - { - RMM_EXPECTS( - alignment <= rmm::CUDA_ALLOCATION_ALIGNMENT && rmm::is_supported_alignment(alignment), - "Alignment must be less than or equal to 256 and a power of two", - rmm::bad_alloc); - auto const stream = cuda_stream_view{}; - void* ptr = do_allocate(bytes, stream); - stream.synchronize(); - return ptr; - } - - /** - * @brief Deallocate memory pointed to by \p ptr. - * - * @note All stream-ordered work on `ptr` must be complete before calling this function otherwise - * behavior is undefined. - * - * @param ptr Pointer to be deallocated - * @param bytes The size in bytes of the allocation. This must be equal to the - * value of `bytes` that was passed to the `allocate` call that returned `ptr`. - * @param alignment The alignment that was passed to the `allocate` call that returned `ptr` - */ - void deallocate_sync( - void* ptr, - std::size_t bytes, - [[maybe_unused]] std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept - { - do_deallocate(ptr, bytes, cuda_stream_view{}); - } - - /** - * @brief Allocates memory of size at least \p bytes on the specified stream. - * - * The returned pointer will have 256 byte alignment regardless of the value - * of alignment. Higher alignments must use the aligned_resource_adaptor. - * - * @throws rmm::bad_alloc When the requested `bytes` cannot be allocated. - * - * @param stream The stream on which to perform the allocation - * @param bytes The size of the allocation - * @param alignment The alignment of the allocation (see notes above) - * @return void* Pointer to the newly allocated memory - */ - void* allocate(cuda_stream_view stream, - std::size_t bytes, - std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) - { - RMM_EXPECTS( - alignment <= rmm::CUDA_ALLOCATION_ALIGNMENT && rmm::is_supported_alignment(alignment), - "Alignment must be less than or equal to 256 and a power of two", - rmm::bad_alloc); - return do_allocate(bytes, stream); - } - - /** - * @brief Deallocate memory pointed to by \p ptr on the specified stream. - * - * @param stream The stream on which to perform the deallocation - * @param ptr Pointer to be deallocated - * @param bytes The size in bytes of the allocation. This must be equal to the - * value of `bytes` that was passed to the `allocate` call that returned `ptr`. - * @param alignment The alignment that was passed to the `allocate` call that returned `ptr` - */ - void deallocate(cuda_stream_view stream, - void* ptr, - std::size_t bytes, - [[maybe_unused]] std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept - { - do_deallocate(ptr, bytes, stream); - } - - /** - * @brief Compare this resource to another. - * - * Two device_memory_resources compare equal if and only if memory allocated - * from one device_memory_resource can be deallocated from the other and vice - * versa. - * - * By default, simply checks if \p *this and \p other refer to the same - * object, i.e., does not check if they are two objects of the same class. - * - * @param other The other resource to compare to - * @returns If the two resources are equivalent - */ - [[nodiscard]] bool is_equal(device_memory_resource const& other) const noexcept - { - return do_is_equal(other); - } - - /** - * @brief Comparison operator with another device_memory_resource - * - * @param other The other resource to compare to - * @return true If the two resources are equivalent - * @return false If the two resources are not equivalent - */ - [[nodiscard]] bool operator==(device_memory_resource const& other) const noexcept - { - return do_is_equal(other); - } - - /** - * @brief Comparison operator with another device_memory_resource - * - * @param other The other resource to compare to - * @return false If the two resources are equivalent - * @return true If the two resources are not equivalent - */ - [[nodiscard]] bool operator!=(device_memory_resource const& other) const noexcept - { - return !do_is_equal(other); - } - - /** - * @brief Enables the `cuda::mr::device_accessible` property - * - * This property declares that a `device_memory_resource` provides device accessible memory - */ - friend void get_property(device_memory_resource const&, cuda::mr::device_accessible) noexcept {} - - private: - /** - * @brief Allocates memory of size at least \p bytes. - * - * The returned pointer will have at minimum 256 byte alignment. - * - * If supported, this operation may optionally be executed on a stream. - * Otherwise, the stream is ignored and the null stream is used. - * - * @param bytes The size of the allocation - * @param stream Stream on which to perform allocation - * @return void* Pointer to the newly allocated memory - */ - virtual void* do_allocate(std::size_t bytes, cuda_stream_view stream) = 0; - - /** - * @brief Deallocate memory pointed to by \p ptr. - * - * If supported, this operation may optionally be executed on a stream. - * Otherwise, the stream is ignored and the null stream is used. - * - * @param ptr Pointer to be deallocated - * @param bytes The size in bytes of the allocation. This must be equal to the - * value of `bytes` that was passed to the `allocate` call that returned `ptr`. - * @param stream Stream on which to perform deallocation - */ - virtual void do_deallocate(void* ptr, std::size_t bytes, cuda_stream_view stream) noexcept = 0; - - /** - * @brief Compare this resource to another. - * - * Two device_memory_resources compare equal if and only if memory allocated - * from one device_memory_resource can be deallocated from the other and vice - * versa. - * - * By default, simply checks if \p *this and \p other refer to the same - * object, i.e., does not check if they are two objects of the same class. - * - * @param other The other resource to compare to - * @return true If the two resources are equivalent - * @return false If the two resources are not equal - */ - [[nodiscard]] virtual bool do_is_equal(device_memory_resource const& other) const noexcept - { - return this == std::addressof(other); - } -}; - -// static property checks -static_assert(cuda::mr::synchronous_resource); -static_assert(cuda::mr::resource); -static_assert( - cuda::mr::synchronous_resource_with); -static_assert(cuda::mr::resource_with); - -/** @} */ // end of group -} // namespace mr -} // namespace RMM_NAMESPACE diff --git a/cpp/include/rmm/mr/fixed_size_memory_resource.hpp b/cpp/include/rmm/mr/fixed_size_memory_resource.hpp index 07d558d47..8ca19f2e4 100644 --- a/cpp/include/rmm/mr/fixed_size_memory_resource.hpp +++ b/cpp/include/rmm/mr/fixed_size_memory_resource.hpp @@ -21,7 +21,7 @@ namespace mr { */ /** - * @brief A `device_memory_resource` which allocates memory blocks of a single fixed size. + * @brief A memory resource which allocates memory blocks of a single fixed size. * * Supports only allocations of size smaller than the configured block_size. * diff --git a/cpp/include/rmm/mr/managed_memory_resource.hpp b/cpp/include/rmm/mr/managed_memory_resource.hpp index 9f972fa1e..203ef315c 100644 --- a/cpp/include/rmm/mr/managed_memory_resource.hpp +++ b/cpp/include/rmm/mr/managed_memory_resource.hpp @@ -34,8 +34,6 @@ class managed_memory_resource final { managed_memory_resource& operator=(managed_memory_resource&&) = default; ///< @default_move_assignment{managed_memory_resource} - // -- CCCL memory resource interface (hides device_memory_resource versions) -- - /** * @brief Allocates memory of size at least \p bytes. * diff --git a/cpp/include/rmm/mr/pinned_host_memory_resource.hpp b/cpp/include/rmm/mr/pinned_host_memory_resource.hpp index d9fa48fc1..c96662033 100644 --- a/cpp/include/rmm/mr/pinned_host_memory_resource.hpp +++ b/cpp/include/rmm/mr/pinned_host_memory_resource.hpp @@ -27,8 +27,8 @@ namespace mr { /** * @brief Memory resource class for allocating pinned host memory. * - * This class uses CUDA's `cudaHostAlloc` to allocate pinned host memory. It implements the - * `cuda::mr::memory_resource` and `cuda::mr::device_memory_resource` concepts, and + * This class uses CUDA's `cudaHostAlloc` to allocate pinned host memory. It satisfies the + * `cuda::mr::resource` and `cuda::mr::synchronous_resource` concepts, and * the `cuda::mr::host_accessible` and `cuda::mr::device_accessible` properties. */ class pinned_host_memory_resource final { @@ -44,8 +44,6 @@ class pinned_host_memory_resource final { pinned_host_memory_resource& operator=(pinned_host_memory_resource&&) = default; ///< @default_move_assignment{pinned_host_memory_resource} - // -- CCCL memory resource interface (hides device_memory_resource versions) -- - /** * @brief Allocates pinned host memory of size at least \p bytes bytes. * @@ -69,7 +67,6 @@ class pinned_host_memory_resource final { // don't allocate anything if the user requested zero bytes if (0 == bytes) { return nullptr; } - // TODO: Use the alignment parameter as an argument to do_allocate std::size_t constexpr alloc_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; return rmm::detail::aligned_host_allocate(bytes, alloc_alignment, [](std::size_t size) { void* ptr{nullptr}; @@ -94,7 +91,6 @@ class pinned_host_memory_resource final { std::size_t bytes, [[maybe_unused]] std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept { - // TODO: Use the alignment parameter as an argument to do_deallocate std::size_t constexpr alloc_alignment = rmm::CUDA_ALLOCATION_ALIGNMENT; rmm::detail::aligned_host_deallocate(ptr, bytes, alloc_alignment, [](void* ptr) { RMM_ASSERT_CUDA_SUCCESS_SAFE_SHUTDOWN(cudaFreeHost(ptr)); diff --git a/cpp/include/rmm/mr/polymorphic_allocator.hpp b/cpp/include/rmm/mr/polymorphic_allocator.hpp index 0a2aca298..1d7bcf858 100644 --- a/cpp/include/rmm/mr/polymorphic_allocator.hpp +++ b/cpp/include/rmm/mr/polymorphic_allocator.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -21,11 +21,11 @@ namespace mr { * @file */ /** - * @brief A stream ordered Allocator using a `rmm::mr::device_memory_resource` to satisfy + * @brief A stream ordered Allocator using a `device_async_resource_ref` to satisfy * (de)allocations. * * Similar to `std::pmr::polymorphic_allocator`, uses the runtime polymorphism of - * `device_memory_resource` to allow containers with `polymorphic_allocator` as their static + * type-erased resource refs to allow containers with `polymorphic_allocator` as their static * allocator type to be interoperable, but exhibit different behavior depending on resource used. * * Unlike STL allocators, `polymorphic_allocator`'s `allocate` and `deallocate` functions are stream diff --git a/cpp/include/rmm/mr/system_memory_resource.hpp b/cpp/include/rmm/mr/system_memory_resource.hpp index 6156c71e3..f47e889a5 100644 --- a/cpp/include/rmm/mr/system_memory_resource.hpp +++ b/cpp/include/rmm/mr/system_memory_resource.hpp @@ -80,8 +80,6 @@ class system_memory_resource final { system_memory_resource& operator=(system_memory_resource&&) = default; ///< @default_move_assignment{system_memory_resource} - // -- CCCL memory resource interface (hides device_memory_resource versions) -- - /** * @brief Allocates memory of size at least \p bytes. * diff --git a/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp b/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp index 80194a18e..ee02f6a28 100644 --- a/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp +++ b/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp @@ -133,7 +133,10 @@ class thrust_allocator : public thrust::device_malloc_allocator { * * This property declares that a `thrust_allocator` provides device accessible memory */ - friend void get_property(thrust_allocator const&, cuda::mr::device_accessible) noexcept {} + RMM_CONSTEXPR_FRIEND void get_property(thrust_allocator const&, + cuda::mr::device_accessible) noexcept + { + } private: cuda_stream_view _stream{}; diff --git a/cpp/include/rmm/resource_ref.hpp b/cpp/include/rmm/resource_ref.hpp index c664d0fcd..c1b1a7324 100644 --- a/cpp/include/rmm/resource_ref.hpp +++ b/cpp/include/rmm/resource_ref.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -7,7 +7,6 @@ #include #include #include -#include namespace RMM_NAMESPACE { @@ -75,22 +74,6 @@ device_async_resource_ref to_device_async_resource_ref_checked(Resource* res) return device_async_resource_ref{*res}; } -// Verify that the device_memory_resource bridge constructors work correctly. -// These assertions validate that resource_ref types can be constructed from raw -// pointers/references. -static_assert( - std::is_constructible_v, - "device_resource_ref must be constructible from device_memory_resource* (via bridge)"); -static_assert( - std::is_constructible_v, - "device_async_resource_ref must be constructible from device_memory_resource* (via bridge)"); -static_assert( - std::is_constructible_v, - "device_resource_ref must be constructible from device_memory_resource& (via bridge)"); -static_assert( - std::is_constructible_v, - "device_async_resource_ref must be constructible from device_memory_resource& (via bridge)"); - // Verify that RMM resource_ref types can be constructed from corresponding CCCL resource_ref types. static_assert( std::is_constructible_v #include #include -#include #include #include #include diff --git a/cpp/tests/device_check_resource_adaptor.hpp b/cpp/tests/device_check_resource_adaptor.hpp index 89812bee0..8b2a1e568 100644 --- a/cpp/tests/device_check_resource_adaptor.hpp +++ b/cpp/tests/device_check_resource_adaptor.hpp @@ -2,19 +2,24 @@ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ +#pragma once +#include #include #include -#include +#include #include #include +#include +#include +#include + #include #include -#include -class device_check_resource_adaptor final : public rmm::mr::device_memory_resource { +class device_check_resource_adaptor final { public: device_check_resource_adaptor(rmm::device_async_resource_ref upstream) : device_id{rmm::get_current_cuda_device()}, upstream_(upstream) @@ -29,33 +34,61 @@ class device_check_resource_adaptor final : public rmm::mr::device_memory_resour return upstream_; } - private: - [[nodiscard]] bool check_device_id() const { return device_id == rmm::get_current_cuda_device(); } - - void* do_allocate(std::size_t bytes, rmm::cuda_stream_view stream) override + void* allocate(cuda::stream_ref stream, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) { bool const is_correct_device = check_device_id(); EXPECT_TRUE(is_correct_device); - if (is_correct_device) { return get_upstream_resource().allocate(stream, bytes); } + if (is_correct_device) { return get_upstream_resource().allocate(stream, bytes, alignment); } return nullptr; } - void do_deallocate(void* ptr, std::size_t bytes, rmm::cuda_stream_view stream) noexcept override + void deallocate(cuda::stream_ref stream, + void* ptr, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept { bool const is_correct_device = check_device_id(); EXPECT_TRUE(is_correct_device); - if (is_correct_device) { get_upstream_resource().deallocate(stream, ptr, bytes); } + if (is_correct_device) { get_upstream_resource().deallocate(stream, ptr, bytes, alignment); } + } + + void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) + { + rmm::cuda_stream_view stream{}; + auto* ptr = allocate(stream, bytes, alignment); + stream.synchronize(); + return ptr; + } + + void deallocate_sync(void* ptr, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept + { + deallocate(rmm::cuda_stream_view{}, ptr, bytes, alignment); } - [[nodiscard]] bool do_is_equal( - rmm::mr::device_memory_resource const& other) const noexcept override + bool operator==(device_check_resource_adaptor const& other) const noexcept { - if (this == std::addressof(other)) { return true; } - auto const* cast = dynamic_cast(&other); - if (cast == nullptr) { return false; } - return get_upstream_resource() == cast->get_upstream_resource(); + return get_upstream_resource() == other.get_upstream_resource(); } + bool operator!=(device_check_resource_adaptor const& other) const noexcept + { + return !(*this == other); + } + + constexpr friend void get_property(device_check_resource_adaptor const&, + cuda::mr::device_accessible) noexcept + { + } + + private: + [[nodiscard]] bool check_device_id() const { return device_id == rmm::get_current_cuda_device(); } + rmm::cuda_device_id device_id; rmm::device_async_resource_ref upstream_; }; + +static_assert(cuda::mr::resource_with); diff --git a/cpp/tests/device_scalar_tests.cpp b/cpp/tests/device_scalar_tests.cpp index a3a3d275c..0276914c6 100644 --- a/cpp/tests/device_scalar_tests.cpp +++ b/cpp/tests/device_scalar_tests.cpp @@ -5,7 +5,6 @@ #include #include -#include #include #include diff --git a/cpp/tests/logger_tests.cpp b/cpp/tests/logger_tests.cpp index 9a578f21d..c14718a09 100644 --- a/cpp/tests/logger_tests.cpp +++ b/cpp/tests/logger_tests.cpp @@ -85,7 +85,7 @@ class raii_temp_directory { * * Events in the log file are expected to occur in the same order as in `expected_events`. * - * @note This function accounts for the fact that `device_memory_resource` automatically pads + * @note This function accounts for the fact that memory resources automatically pad * allocations to a multiple of 8 bytes by rounding up the expected allocation sizes to a multiple * of 8. * diff --git a/cpp/tests/mock_resource.hpp b/cpp/tests/mock_resource.hpp index b65d1ea4d..023da63fe 100644 --- a/cpp/tests/mock_resource.hpp +++ b/cpp/tests/mock_resource.hpp @@ -4,23 +4,91 @@ */ #pragma once -#include +#include +#include + +#include +#include #include +#include + namespace rmm::test { -class mock_resource : public rmm::mr::device_memory_resource { +class mock_resource { public: - MOCK_METHOD(void*, do_allocate, (std::size_t, cuda_stream_view), (override)); - MOCK_METHOD(void, do_deallocate, (void*, std::size_t, cuda_stream_view), (noexcept, override)); + MOCK_METHOD(void*, allocate, (cuda::stream_ref, std::size_t, std::size_t)); + MOCK_METHOD(void, deallocate, (cuda::stream_ref, void*, std::size_t, std::size_t), (noexcept)); + + void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) + { + return allocate(rmm::cuda_stream_view{}, bytes, alignment); + } + + void deallocate_sync(void* ptr, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept + { + deallocate(rmm::cuda_stream_view{}, ptr, bytes, alignment); + } + bool operator==(mock_resource const&) const noexcept { return true; } bool operator!=(mock_resource const&) const { return false; } - friend void get_property(mock_resource const&, cuda::mr::device_accessible) noexcept {} + constexpr friend void get_property(mock_resource const&, cuda::mr::device_accessible) noexcept {} using size_pair = std::pair; }; // static property checks static_assert(cuda::mr::resource_with); +// Copyable wrapper around mock_resource that satisfies CCCL basic_any's requirements. +// GMock types are not copyable, so they cannot be type-erased by CCCL's resource_ref +// (which uses basic_any internally). This thin forwarding layer solves that. +class mock_resource_wrapper { + public: + explicit mock_resource_wrapper(mock_resource* mock) noexcept : mock_{mock} {} + + void* allocate(cuda::stream_ref stream, std::size_t bytes, std::size_t alignment) + { + return mock_->allocate(stream, bytes, alignment); + } + + void deallocate(cuda::stream_ref stream, + void* ptr, + std::size_t bytes, + std::size_t alignment) noexcept + { + mock_->deallocate(stream, ptr, bytes, alignment); + } + + void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) + { + return mock_->allocate_sync(bytes, alignment); + } + + void deallocate_sync(void* ptr, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept + { + mock_->deallocate_sync(ptr, bytes, alignment); + } + + bool operator==(mock_resource_wrapper const& other) const noexcept + { + return mock_ == other.mock_; + } + bool operator!=(mock_resource_wrapper const& other) const noexcept { return !(*this == other); } + + constexpr friend void get_property(mock_resource_wrapper const&, + cuda::mr::device_accessible) noexcept + { + } + + private: + mock_resource* mock_; +}; + +static_assert(cuda::mr::resource_with); + } // namespace rmm::test diff --git a/cpp/tests/mr/aligned_mr_tests.cpp b/cpp/tests/mr/aligned_mr_tests.cpp index e9a2b2cab..9c1da0173 100644 --- a/cpp/tests/mr/aligned_mr_tests.cpp +++ b/cpp/tests/mr/aligned_mr_tests.cpp @@ -20,6 +20,7 @@ namespace rmm::test { namespace { +using ::testing::_; using ::testing::Return; using aligned_adaptor = rmm::mr::aligned_resource_adaptor; @@ -78,32 +79,35 @@ TEST_P(allocation_size, MultiThreaded) TEST(AlignedTest, ThrowOnInvalidAllocationAlignment) { mock_resource mock; - auto construct_alignment = [](mock_resource& memres, std::size_t align) { - aligned_adaptor mr{memres, align}; + mock_resource_wrapper wrapper{&mock}; + auto construct_alignment = [](mock_resource_wrapper& w, std::size_t align) { + aligned_adaptor mr{device_async_resource_ref{w}, align}; }; - EXPECT_THROW(construct_alignment(mock, 255), rmm::logic_error); - EXPECT_NO_THROW(construct_alignment(mock, 256)); - EXPECT_THROW(construct_alignment(mock, 768), rmm::logic_error); + EXPECT_THROW(construct_alignment(wrapper, 255), rmm::logic_error); + EXPECT_NO_THROW(construct_alignment(wrapper, 256)); + EXPECT_THROW(construct_alignment(wrapper, 768), rmm::logic_error); } TEST(AlignedTest, SupportsGetMemInfo) { mock_resource mock; - aligned_adaptor mr{mock}; + mock_resource_wrapper wrapper{&mock}; + aligned_adaptor mr{device_async_resource_ref{wrapper}}; } TEST(AlignedTest, DefaultAllocationAlignmentPassthrough) { mock_resource mock; - aligned_adaptor mr{mock}; + mock_resource_wrapper wrapper{&mock}; + aligned_adaptor mr{device_async_resource_ref{wrapper}}; cuda_stream_view stream; void* const pointer = int_to_address(123); { auto const size{5}; - EXPECT_CALL(mock, do_allocate(size, stream)).WillOnce(Return(pointer)); - EXPECT_CALL(mock, do_deallocate(pointer, size, stream)).Times(1); + EXPECT_CALL(mock, allocate(_, size, _)).WillOnce(Return(pointer)); + EXPECT_CALL(mock, deallocate(_, pointer, size, _)).Times(1); } { @@ -116,16 +120,17 @@ TEST(AlignedTest, DefaultAllocationAlignmentPassthrough) TEST(AlignedTest, BelowAlignmentThresholdPassthrough) { mock_resource mock; + mock_resource_wrapper wrapper{&mock}; auto const alignment{4096}; auto const threshold{65536}; - aligned_adaptor mr{mock, alignment, threshold}; + aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold}; cuda_stream_view stream; void* const pointer = int_to_address(123); { auto const size{3}; - EXPECT_CALL(mock, do_allocate(size, stream)).WillOnce(Return(pointer)); - EXPECT_CALL(mock, do_deallocate(pointer, size, stream)).Times(1); + EXPECT_CALL(mock, allocate(_, size, _)).WillOnce(Return(pointer)); + EXPECT_CALL(mock, deallocate(_, pointer, size, _)).Times(1); } { @@ -137,8 +142,8 @@ TEST(AlignedTest, BelowAlignmentThresholdPassthrough) { auto const size{65528}; void* const pointer1 = int_to_address(456); - EXPECT_CALL(mock, do_allocate(size, stream)).WillOnce(Return(pointer1)); - EXPECT_CALL(mock, do_deallocate(pointer1, size, stream)).Times(1); + EXPECT_CALL(mock, allocate(_, size, _)).WillOnce(Return(pointer1)); + EXPECT_CALL(mock, deallocate(_, pointer1, size, _)).Times(1); EXPECT_EQ(mr.allocate(stream, size, rmm::CUDA_ALLOCATION_ALIGNMENT), pointer1); mr.deallocate(stream, pointer1, size, rmm::CUDA_ALLOCATION_ALIGNMENT); } @@ -147,17 +152,18 @@ TEST(AlignedTest, BelowAlignmentThresholdPassthrough) TEST(AlignedTest, UpstreamAddressAlreadyAligned) { mock_resource mock; + mock_resource_wrapper wrapper{&mock}; auto const alignment{4096}; auto const threshold{65536}; - aligned_adaptor mr{mock, alignment, threshold}; + aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold}; cuda_stream_view stream; void* const pointer = int_to_address(4096); { auto const size{69376}; - EXPECT_CALL(mock, do_allocate(size, stream)).WillOnce(Return(pointer)); - EXPECT_CALL(mock, do_deallocate(pointer, size, stream)).Times(1); + EXPECT_CALL(mock, allocate(_, size, _)).WillOnce(Return(pointer)); + EXPECT_CALL(mock, deallocate(_, pointer, size, _)).Times(1); } { @@ -170,16 +176,17 @@ TEST(AlignedTest, UpstreamAddressAlreadyAligned) TEST(AlignedTest, AlignUpstreamAddress) { mock_resource mock; + mock_resource_wrapper wrapper{&mock}; auto const alignment{4096}; auto const threshold{65536}; - aligned_adaptor mr{mock, alignment, threshold}; + aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold}; cuda_stream_view stream; { void* const pointer = int_to_address(256); auto const size{69376}; - EXPECT_CALL(mock, do_allocate(size, stream)).WillOnce(Return(pointer)); - EXPECT_CALL(mock, do_deallocate(pointer, size, stream)).Times(1); + EXPECT_CALL(mock, allocate(_, size, _)).WillOnce(Return(pointer)); + EXPECT_CALL(mock, deallocate(_, pointer, size, _)).Times(1); } { @@ -193,9 +200,10 @@ TEST(AlignedTest, AlignUpstreamAddress) TEST(AlignedTest, AlignMultiple) { mock_resource mock; + mock_resource_wrapper wrapper{&mock}; auto const alignment{4096}; auto const threshold{65536}; - aligned_adaptor mr{mock, alignment, threshold}; + aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold}; cuda_stream_view stream; @@ -206,12 +214,12 @@ TEST(AlignedTest, AlignMultiple) auto const size1{69376}; auto const size2{77568}; auto const size3{81664}; - EXPECT_CALL(mock, do_allocate(size1, stream)).WillOnce(Return(pointer1)); - EXPECT_CALL(mock, do_allocate(size2, stream)).WillOnce(Return(pointer2)); - EXPECT_CALL(mock, do_allocate(size3, stream)).WillOnce(Return(pointer3)); - EXPECT_CALL(mock, do_deallocate(pointer1, size1, stream)).Times(1); - EXPECT_CALL(mock, do_deallocate(pointer2, size2, stream)).Times(1); - EXPECT_CALL(mock, do_deallocate(pointer3, size3, stream)).Times(1); + EXPECT_CALL(mock, allocate(_, size1, _)).WillOnce(Return(pointer1)); + EXPECT_CALL(mock, allocate(_, size2, _)).WillOnce(Return(pointer2)); + EXPECT_CALL(mock, allocate(_, size3, _)).WillOnce(Return(pointer3)); + EXPECT_CALL(mock, deallocate(_, pointer1, size1, _)).Times(1); + EXPECT_CALL(mock, deallocate(_, pointer2, size2, _)).Times(1); + EXPECT_CALL(mock, deallocate(_, pointer3, size3, _)).Times(1); } { diff --git a/cpp/tests/mr/arena_mr_tests.cpp b/cpp/tests/mr/arena_mr_tests.cpp index 61b8d81ae..548bbe76c 100644 --- a/cpp/tests/mr/arena_mr_tests.cpp +++ b/cpp/tests/mr/arena_mr_tests.cpp @@ -10,9 +10,9 @@ #include #include #include -#include #include +#include #include #include @@ -30,23 +30,74 @@ namespace rmm::test { namespace { -class mock_memory_resource : public rmm::mr::device_memory_resource { +class mock_memory_resource { public: - MOCK_METHOD(void*, do_allocate, (std::size_t, cuda_stream_view)); - MOCK_METHOD(void, do_deallocate, (void*, std::size_t, cuda_stream_view), (noexcept)); MOCK_METHOD(void*, allocate_sync, (std::size_t, std::size_t)); MOCK_METHOD(void, deallocate_sync, (void*, std::size_t, std::size_t), (noexcept)); - MOCK_METHOD(void*, allocate, (cuda_stream_view, std::size_t, std::size_t)); - MOCK_METHOD(void, deallocate, (cuda_stream_view, void*, std::size_t, std::size_t), (noexcept)); + MOCK_METHOD(void*, allocate, (cuda::stream_ref, std::size_t, std::size_t)); + MOCK_METHOD(void, deallocate, (cuda::stream_ref, void*, std::size_t, std::size_t), (noexcept)); bool operator==(mock_memory_resource const&) const noexcept { return true; } bool operator!=(mock_memory_resource const&) const { return false; } - friend void get_property(mock_memory_resource const&, cuda::mr::device_accessible) noexcept {} + constexpr friend void get_property(mock_memory_resource const&, + cuda::mr::device_accessible) noexcept + { + } }; // static property checks static_assert(cuda::mr::resource_with); +// Copyable wrapper so the mock can be type-erased by CCCL basic_any. +class mock_memory_resource_wrapper { + public: + explicit mock_memory_resource_wrapper(mock_memory_resource* mock) noexcept : mock_{mock} {} + + void* allocate(cuda::stream_ref stream, std::size_t bytes, std::size_t alignment) + { + return mock_->allocate(stream, bytes, alignment); + } + + void deallocate(cuda::stream_ref stream, + void* ptr, + std::size_t bytes, + std::size_t alignment) noexcept + { + mock_->deallocate(stream, ptr, bytes, alignment); + } + + void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) + { + return mock_->allocate_sync(bytes, alignment); + } + + void deallocate_sync(void* ptr, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept + { + mock_->deallocate_sync(ptr, bytes, alignment); + } + + bool operator==(mock_memory_resource_wrapper const& other) const noexcept + { + return mock_ == other.mock_; + } + bool operator!=(mock_memory_resource_wrapper const& other) const noexcept + { + return !(*this == other); + } + + constexpr friend void get_property(mock_memory_resource_wrapper const&, + cuda::mr::device_accessible) noexcept + { + } + + private: + mock_memory_resource* mock_; +}; + +static_assert(cuda::mr::resource_with); + using rmm::mr::detail::arena::block; using rmm::mr::detail::arena::byte_span; using rmm::mr::detail::arena::superblock; @@ -65,14 +116,17 @@ auto const fake_address4 = reinterpret_cast(superblock::minimum_size * 2) struct ArenaTest : public ::testing::Test { void SetUp() override { - EXPECT_CALL(mock_mr, do_allocate(arena_size, ::testing::_)).WillOnce(Return(fake_address3)); - EXPECT_CALL(mock_mr, do_deallocate(fake_address3, arena_size, ::testing::_)); - global = std::make_unique(mock_mr, arena_size); + EXPECT_CALL(mock_mr, allocate_sync(arena_size, ::testing::_)).WillOnce(Return(fake_address3)); + EXPECT_CALL(mock_mr, deallocate_sync(fake_address3, arena_size, ::testing::_)); + mock_wrapper = std::make_unique(&mock_mr); + global = + std::make_unique(rmm::device_async_resource_ref{*mock_wrapper}, arena_size); per_thread = std::make_unique(*global); } std::size_t arena_size{superblock::minimum_size * 4}; mock_memory_resource mock_mr{}; + std::unique_ptr mock_wrapper{}; std::unique_ptr global{}; std::unique_ptr per_thread{}; }; diff --git a/cpp/tests/mr/callback_mr_tests.cpp b/cpp/tests/mr/callback_mr_tests.cpp index 75d98cf64..01f8d5d49 100644 --- a/cpp/tests/mr/callback_mr_tests.cpp +++ b/cpp/tests/mr/callback_mr_tests.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -25,10 +25,11 @@ using ::testing::_; TEST(CallbackTest, TestCallbacksAreInvoked) { - auto base_mr = mock_resource(); - auto base_ref = device_async_resource_ref{base_mr}; - EXPECT_CALL(base_mr, do_allocate(10_MiB, cuda_stream_view{})).Times(1); - EXPECT_CALL(base_mr, do_deallocate(_, 10_MiB, cuda_stream_view{})).Times(1); + auto base_mr = mock_resource(); + auto base_wrapper = mock_resource_wrapper{&base_mr}; + auto base_ref = device_async_resource_ref{base_wrapper}; + EXPECT_CALL(base_mr, allocate(_, 10_MiB, _)).Times(1); + EXPECT_CALL(base_mr, deallocate(_, _, 10_MiB, _)).Times(1); auto allocate_callback = [](std::size_t size, cuda_stream_view stream, void* arg) { auto base_mr = *static_cast(arg); diff --git a/cpp/tests/mr/delayed_memory_resource.hpp b/cpp/tests/mr/delayed_memory_resource.hpp index 00ffcf101..6288e0c51 100644 --- a/cpp/tests/mr/delayed_memory_resource.hpp +++ b/cpp/tests/mr/delayed_memory_resource.hpp @@ -48,7 +48,10 @@ class delayed_memory_resource { upstream_.deallocate(stream, ptr, bytes, alignment); std::this_thread::sleep_for(delay_); } - friend void get_property(delayed_memory_resource const&, cuda::mr::device_accessible) noexcept {} + constexpr friend void get_property(delayed_memory_resource const&, + cuda::mr::device_accessible) noexcept + { + } bool operator==(delayed_memory_resource const& other) const noexcept { return this == std::addressof(other); diff --git a/cpp/tests/mr/failure_callback_mr_tests.cpp b/cpp/tests/mr/failure_callback_mr_tests.cpp index 5af05bd89..bd835680c 100644 --- a/cpp/tests/mr/failure_callback_mr_tests.cpp +++ b/cpp/tests/mr/failure_callback_mr_tests.cpp @@ -5,12 +5,15 @@ #include "../byte_literals.hpp" +#include #include #include -#include #include #include +#include +#include + #include #include @@ -33,22 +36,43 @@ bool failure_handler(std::size_t /*bytes*/, void* arg) } template -class always_throw_memory_resource final : public mr::device_memory_resource { - private: - void* do_allocate(std::size_t /*bytes*/, cuda_stream_view /*stream*/) override +class always_throw_memory_resource final { + public: + void* allocate(cuda::stream_ref /*stream*/, + std::size_t /*bytes*/, + std::size_t /*alignment*/ = rmm::CUDA_ALLOCATION_ALIGNMENT) { throw ExceptionType{"foo"}; } - void do_deallocate(void* /*ptr*/, - std::size_t /*bytes*/, - cuda_stream_view /*stream*/) noexcept override {}; + void deallocate(cuda::stream_ref /*stream*/, + void* /*ptr*/, + std::size_t /*bytes*/, + std::size_t /*alignment*/ = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept {}; + + void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) + { + return allocate(rmm::cuda_stream_view{}, bytes, alignment); + } + void deallocate_sync(void* ptr, + std::size_t bytes, + std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept + { + deallocate(rmm::cuda_stream_view{}, ptr, bytes, alignment); + } + + bool operator==(always_throw_memory_resource const&) const noexcept { return true; } + bool operator!=(always_throw_memory_resource const&) const noexcept { return false; } + constexpr friend void get_property(always_throw_memory_resource const&, + cuda::mr::device_accessible) noexcept + { + } }; TEST(FailureCallbackTest, RetryAllocationOnce) { always_throw_memory_resource throwing_mr; bool retried{false}; - failure_callback_adaptor<> mr{&throwing_mr, failure_handler, &retried}; + failure_callback_adaptor<> mr{throwing_mr, failure_handler, &retried}; EXPECT_EQ(retried, false); EXPECT_THROW((void)mr.allocate_sync(1_MiB), rmm::bad_alloc); EXPECT_EQ(retried, true); @@ -67,7 +91,7 @@ TEST(FailureCallbackTest, DifferentExceptionTypes) { bool retried{false}; failure_callback_adaptor bad_alloc_callback_mr{ - &bad_alloc_mr, failure_handler, &retried}; + bad_alloc_mr, failure_handler, &retried}; EXPECT_EQ(retried, false); EXPECT_THROW((void)bad_alloc_callback_mr.allocate_sync(1_MiB), rmm::bad_alloc); @@ -79,8 +103,7 @@ TEST(FailureCallbackTest, DifferentExceptionTypes) { bool retried{false}; - failure_callback_adaptor oom_callback_mr{ - &oom_mr, failure_handler, &retried}; + failure_callback_adaptor oom_callback_mr{oom_mr, failure_handler, &retried}; EXPECT_EQ(retried, false); EXPECT_THROW((void)oom_callback_mr.allocate_sync(1_MiB), rmm::out_of_memory); EXPECT_EQ(retried, true); @@ -92,7 +115,7 @@ TEST(FailureCallbackTest, DifferentExceptionTypes) bool retried{false}; failure_callback_adaptor oom_callback_mr{ - &bad_alloc_mr, failure_handler, &retried}; + bad_alloc_mr, failure_handler, &retried}; EXPECT_EQ(retried, false); EXPECT_THROW((void)oom_callback_mr.allocate_sync(1_MiB), rmm::bad_alloc); // bad_alloc passes through diff --git a/cpp/tests/mr/resource_ref_conversion_tests.cpp b/cpp/tests/mr/resource_ref_conversion_tests.cpp index dac965d14..d3c17f0b3 100644 --- a/cpp/tests/mr/resource_ref_conversion_tests.cpp +++ b/cpp/tests/mr/resource_ref_conversion_tests.cpp @@ -54,7 +54,10 @@ class new_delete_memory_resource { bool operator!=(new_delete_memory_resource const& other) const { return !operator==(other); } // NOLINTBEGIN - friend void get_property(new_delete_memory_resource const&, cuda::mr::host_accessible) noexcept {} + constexpr friend void get_property(new_delete_memory_resource const&, + cuda::mr::host_accessible) noexcept + { + } // NOLINTEND };