Implement resource wrappers for <memory_resource>

This includes
* `cuda::mr::resource_ref`
* `cuda::mr::async_resource_ref`

.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp

@@ -0,0 +1,124 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11
+
+// cuda::mr::async_resource_ref construction
+#include <cuda/memory_resource>
+
+#include <cuda/std/cstdint>
+
+template <class T>
+struct property_with_value {
+  using value_type = T;
+};
+
+template <class T>
+struct property_without_value {};
+
+template <class... Properties>
+struct async_resource {
+  inline __host__ __device__ void* allocate(std::size_t, std::size_t) {
+    return nullptr;
+  }
+
+  inline __host__ __device__ void deallocate(void* ptr, std::size_t,
+                                             std::size_t) {}
+
+  inline __host__ __device__ void* allocate_async(std::size_t, std::size_t,
+                                                  cuda::stream_ref) {
+    return &_val;
+  }
+
+  inline __host__ __device__ void
+  deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) {
+    // ensure that we did get the right inputs forwarded
+    _val = *static_cast<int*>(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<Property> &&
+   _CUDA_VSTD::_One_of<Property, Properties...>) //
+      inline __host__ __device__
+      friend void get_property(const async_resource&, Property) noexcept {}
+
+  _LIBCUDACXX_TEMPLATE(class Property)
+  (requires cuda::mr::property_with_value<Property>&&
+       _CUDA_VSTD::_One_of<Property, Properties...>) //
+      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<int>,
+                                     property_with_value<double>,
+                                     property_without_value<std::size_t> >;
+
+using matching_properties = async_resource<property_with_value<double>,
+                                           property_without_value<std::size_t>,
+                                           property_with_value<int> >;
+
+using missing_stateful_property =
+    async_resource<property_with_value<int>,
+                   property_without_value<std::size_t> >;
+using missing_stateless_property =
+    async_resource<property_with_value<int>, property_with_value<double> >;
+
+using cuda::std::is_constructible;
+static_assert(is_constructible<ref, matching_properties&>::value, "");
+static_assert(!is_constructible<ref, missing_stateful_property&>::value, "");
+static_assert(!is_constructible<ref, missing_stateless_property&>::value, "");
+
+static_assert(is_constructible<ref, ref&>::value, "");
+
+// Ensure we require a mutable valid reference and do not bind against rvalues
+static_assert(!is_constructible<ref, matching_properties>::value, "");
+static_assert(!is_constructible<ref, const matching_properties&>::value, "");
+
+static_assert(cuda::std::is_copy_constructible<ref>::value, "");
+static_assert(cuda::std::is_move_constructible<ref>::value, "");
+} // namespace constructible
+
+namespace assignable {
+using ref =
+    cuda::mr::async_resource_ref<property_with_value<int>,
+                                     property_with_value<double>,
+                                     property_without_value<std::size_t> >;
+
+using res =
+    async_resource<property_with_value<int>, property_with_value<double>,
+                   property_without_value<std::size_t> >;
+
+using other_res =
+    async_resource<property_without_value<int>, property_with_value<int>,
+                   property_with_value<double>,
+                   property_without_value<std::size_t> >;
+
+using cuda::std::is_assignable;
+static_assert(cuda::std::is_assignable<ref, res&>::value, "");
+static_assert(cuda::std::is_assignable<ref, other_res&>::value, "");
+
+static_assert(cuda::std::is_copy_assignable<ref>::value, "");
+static_assert(cuda::std::is_move_assignable<ref>::value, "");
+} // namespace assignable
+
+int main(int, char**) { return 0; }

.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.cpp

@@ -0,0 +1,97 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11
+
+// cuda::mr::resource_ref properties
+#include <cuda/memory_resource>
+
+#include <cuda/std/cassert>
+#include <cuda/std/cstdint>
+
+template <class T>
+struct property_with_value {
+  using value_type = T;
+};
+
+template <class T>
+struct property_without_value {};
+
+template <class... Properties>
+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<int*>(ptr);
+  }
+
+  inline __host__ __device__ void* allocate_async(std::size_t, std::size_t,
+                                                  cuda::stream_ref) {
+    return &_val;
+  }
+
+  inline __host__ __device__ void
+  deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) {
+    // ensure that we did get the right inputs forwarded
+    _val = *static_cast<int*>(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<Property> &&
+   _CUDA_VSTD::_One_of<Property, Properties...>) //
+      inline __host__ __device__
+      friend void get_property(const resource&, Property) noexcept {}
+
+  _LIBCUDACXX_TEMPLATE(class Property)
+  (requires cuda::mr::property_with_value<Property>&&
+       _CUDA_VSTD::_One_of<Property, Properties...>) //
+      inline __host__ __device__                     //
+      friend typename Property::value_type
+      get_property(const resource& res, Property) noexcept {
+    return res._val;
+  }
+};
+
+template <class PropA, class PropB>
+__host__ __device__ void test_conversion_from_async_resource_ref() {
+  resource<PropA, PropB> input{42};
+  cuda::mr::async_resource_ref<PropA, PropB> ref_input{input};
+  cuda::mr::async_resource_ref<PropB> ref{ref_input};
+
+  // Ensure that we properly pass on the allocate function
+  assert(input.allocate_async(0, 0) == ref.allocate_async(0, 0));
+
+  // Ensure we are deallocating properly
+  int expected_after_deallocate = 1337;
+  ref.deallocate_async(static_cast<void*>(&expected_after_deallocate), 0, 0);
+  assert(input._val == expected_after_deallocate);
+}
+
+int main(int, char**) {                        
+  test_conversion_from_async_resource_ref<property_with_value<short>,
+                        property_with_value<int> >();
+  test_conversion_from_async_resource_ref<property_with_value<short>,
+                        property_without_value<int> >();
+  test_conversion_from_async_resource_ref<property_without_value<short>,
+                        property_without_value<int> >();
+  return 0;
+}

.upstream-tests/test/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp

@@ -0,0 +1,98 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11
+
+// cuda::mr::async_resource_ref equality
+#include <cuda/memory_resource>
+#include <cuda/stream_ref>
+
+#include <cuda/std/cassert>
+#include <cuda/std/cstdint>
+
+template <class T>
+struct property_with_value {
+  using value_type = T;
+};
+
+template <class T>
+struct property_without_value {};
+
+template <class... Properties>
+struct async_resource {
+  inline __host__ __device__ void* allocate(std::size_t, std::size_t) {
+    return nullptr;
+  }
+
+  inline __host__ __device__ void deallocate(void* ptr, std::size_t,
+                                             std::size_t) {}
+
+  inline __host__ __device__ void* allocate_async(std::size_t, std::size_t,
+                                                  cuda::stream_ref) {
+    return &_val;
+  }
+
+  inline __host__ __device__ void
+  deallocate_async(void* ptr, std::size_t, std::size_t, cuda::stream_ref) {
+    // ensure that we did get the right inputs forwarded
+    _val = *static_cast<int*>(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<Property> &&
+   _CUDA_VSTD::_One_of<Property, Properties...>) //
+      inline __host__ __device__
+      friend void get_property(const async_resource&, Property) noexcept {}
+
+  _LIBCUDACXX_TEMPLATE(class Property)
+  (requires cuda::mr::property_with_value<Property>&&
+       _CUDA_VSTD::_One_of<Property, Properties...>) //
+      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<int>,
+                                     property_with_value<double>,
+                                     property_without_value<std::size_t> >;
+using different_properties =
+    cuda::mr::async_resource_ref<property_with_value<short>,
+                                     property_with_value<int>,
+                                     property_without_value<std::size_t> >;
+
+using res =
+    async_resource<property_with_value<int>, property_with_value<double>,
+                   property_without_value<std::size_t> >;
+
+__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;
+}