Implement {async_}resource_ref

.upstream-tests/test/cuda/memory_resource/get_property/forward_property.pass.cpp

@@ -0,0 +1,129 @@
+//===----------------------------------------------------------------------===//
+//
+// 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::forward_property
+
+#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE
+
+#include <cuda/std/cassert>
+#include <cuda/memory_resource>
+
+namespace has_upstream_resource {
+  struct Upstream{};
+
+  __device__ Upstream upstream{};
+
+  struct with_reference {
+      Upstream& upstream_resource() const { return upstream; }
+  };
+  static_assert(cuda::__has_upstream_resource<with_reference, Upstream>, "");
+
+  struct with_const_reference {
+      const Upstream& upstream_resource() const { return upstream; }
+  };
+  static_assert(cuda::__has_upstream_resource<with_const_reference, Upstream>, "");
+
+  struct with_value {
+      Upstream upstream_resource() const { return Upstream{}; }
+  };
+  static_assert(cuda::__has_upstream_resource<with_value, Upstream>, "");
+
+  struct with_const_value {
+      const Upstream upstream_resource() const { return Upstream{}; }
+  };
+  static_assert(cuda::__has_upstream_resource<with_const_value, Upstream>, "");
+
+  struct Convertible {
+      operator Upstream() { return Upstream{}; }
+  };
+
+  struct with_conversion {
+      Convertible upstream_resource() const { return Convertible{}; }
+  };
+  static_assert(!cuda::__has_upstream_resource<with_conversion, Upstream>, "");
+} // namespace has_upstream_resource
+
+namespace forward_property {
+  struct prop_with_value { using value_type = int; };
+  struct prop {};
+
+  template<class Upstream>
+  struct derived_plain : public cuda::forward_property<derived_plain<Upstream>, Upstream>
+  {
+    constexpr Upstream upstream_resource() const noexcept { return Upstream{}; }
+  };
+
+  struct upstream_with_valueless_property {
+    friend constexpr void get_property(const upstream_with_valueless_property&, prop) {}
+  };
+  static_assert( cuda::has_property<derived_plain<upstream_with_valueless_property>, prop>, "");
+  static_assert(!cuda::has_property<derived_plain<upstream_with_valueless_property>, prop_with_value>, "");
+
+  struct upstream_with_stateful_property {
+    friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) {
+      return 42;
+    }
+  };
+  static_assert(!cuda::has_property<derived_plain<upstream_with_stateful_property>, prop>, "");
+  static_assert( cuda::has_property<derived_plain<upstream_with_stateful_property>, prop_with_value>, "");
+
+  struct upstream_with_both_properties {
+    friend constexpr void get_property(const upstream_with_both_properties&, prop) {}
+    friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) {
+      return 42;
+    }
+  };
+  static_assert( cuda::has_property<derived_plain<upstream_with_both_properties>, prop>, "");
+  static_assert( cuda::has_property<derived_plain<upstream_with_both_properties>, prop_with_value>, "");
+
+  struct derived_override : public cuda::forward_property<derived_override, upstream_with_both_properties>
+  {
+    constexpr upstream_with_both_properties upstream_resource() const noexcept {
+      return upstream_with_both_properties{};
+    }
+    // Get called directly so needs to be annotated
+    __host__ __device__ friend constexpr int get_property(const derived_override&, prop_with_value) {
+      return 1337;
+    }
+  };
+
+  struct convertible_to_upstream {
+    operator upstream_with_both_properties() const noexcept {
+      return upstream_with_both_properties{};
+    }
+  };
+
+  struct derived_with_converstin_upstream_resource : public cuda::forward_property<derived_with_converstin_upstream_resource, upstream_with_both_properties>
+  {
+    constexpr convertible_to_upstream upstream_resource() const noexcept {
+      return convertible_to_upstream{};
+    }
+  };
+  static_assert(!cuda::has_property<derived_with_converstin_upstream_resource, prop_with_value>, "");
+
+  __host__ __device__ constexpr bool test_stateful() {
+    using derived_no_override = derived_plain<upstream_with_stateful_property>;
+    const derived_no_override without_override{};
+    assert(get_property(without_override, prop_with_value{}) == 42);
+
+    const derived_override with_override{};
+    assert(get_property(with_override, prop_with_value{}) == 1337);
+
+    return true;
+  }
+} // namespace forward_property
+
+
+int main(int, char**) {
+  forward_property::test_stateful();
+  static_assert(forward_property::test_stateful(), "");
+  return 0;
+}

.upstream-tests/test/cuda/memory_resource/get_property/get_property.pass.cpp

@@ -0,0 +1,64 @@
+//===----------------------------------------------------------------------===//
+//
+// 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::get_property
+
+#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE
+
+#include <cuda/std/cassert>
+#include <cuda/memory_resource>
+
+struct prop_with_value {
+  using value_type = int;
+};
+struct prop {};
+
+struct upstream_with_valueless_property {
+  friend constexpr void get_property(const upstream_with_valueless_property&, prop) {}
+};
+static_assert( cuda::std::invocable<decltype(cuda::get_property), upstream_with_valueless_property, prop>, "");
+static_assert(!cuda::std::invocable<decltype(cuda::get_property), upstream_with_valueless_property, prop_with_value>, "");
+
+struct upstream_with_stateful_property {
+  friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) {
+    return 42;
+  }
+};
+static_assert(!cuda::std::invocable<decltype(cuda::get_property), upstream_with_stateful_property, prop>, "");
+static_assert( cuda::std::invocable<decltype(cuda::get_property), upstream_with_stateful_property, prop_with_value>, "");
+
+struct upstream_with_both_properties {
+  friend constexpr void get_property(const upstream_with_both_properties&, prop) {}
+  friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) {
+    return 42;
+  }
+};
+static_assert( cuda::std::invocable<decltype(cuda::get_property), upstream_with_both_properties, prop>, "");
+static_assert( cuda::std::invocable<decltype(cuda::get_property), upstream_with_both_properties, prop_with_value>, "");
+
+__host__ __device__ constexpr bool test() {
+  upstream_with_valueless_property with_valueless{};
+  cuda::get_property(with_valueless, prop{});
+
+  upstream_with_stateful_property with_value{};
+  assert(cuda::get_property(with_value, prop_with_value{}) == 42);
+
+  upstream_with_both_properties with_both{};
+  cuda::get_property(with_both, prop{});
+  assert(cuda::get_property(with_both, prop_with_value{}) == 42);
+  return true;
+}
+
+int main(int, char**) {
+  test();
+  static_assert(test(), "");
+  return 0;
+}

.upstream-tests/test/cuda/memory_resource/get_property/has_property.pass.cpp

@@ -0,0 +1,63 @@
+//===----------------------------------------------------------------------===//
+//
+// 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::has_property, cuda::has_property_with
+
+#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE
+
+#include <cuda/memory_resource>
+
+struct prop_with_value {
+  using value_type = int;
+};
+struct prop {};
+
+static_assert(cuda::property_with_value<prop_with_value>, "");
+static_assert(!cuda::property_with_value<prop>, "");
+
+struct valid_property {
+  friend void get_property(const valid_property&, prop) {}
+};
+static_assert(!cuda::has_property<valid_property, prop_with_value>, "");
+static_assert(cuda::has_property<valid_property, prop>, "");
+static_assert(!cuda::has_property_with<valid_property, prop, int>, "");
+
+struct valid_property_with_value {
+  friend int get_property(const valid_property_with_value&, prop_with_value) {
+    return 42;
+  }
+};
+static_assert(
+    cuda::has_property<valid_property_with_value, prop_with_value>, "");
+static_assert(!cuda::has_property<valid_property_with_value, prop>, "");
+static_assert(cuda::has_property_with<valid_property_with_value,
+                                          prop_with_value, int>,
+              "");
+static_assert(!cuda::has_property_with<valid_property_with_value,
+                                           prop_with_value, double>,
+              "");
+
+struct derived_from_property : public valid_property {
+  friend int get_property(const derived_from_property&, prop_with_value) {
+    return 42;
+  }
+};
+static_assert(cuda::has_property<derived_from_property, prop_with_value>,
+              "");
+static_assert(cuda::has_property<derived_from_property, prop>, "");
+static_assert(
+    cuda::has_property_with<derived_from_property, prop_with_value, int>,
+    "");
+static_assert(!cuda::has_property_with<derived_from_property,
+                                           prop_with_value, double>,
+              "");
+
+int main(int, char**) { return 0; }

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

@@ -0,0 +1,107 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 properties
+
+#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE
+
+#include <cuda/memory_resource>
+
+#include <cuda/std/cassert>
+#include <cuda/std/cstdint>
+
+struct async_resource {
+  void* allocate(std::size_t, std::size_t) { return &_val; }
+
+  void deallocate(void* ptr, std::size_t, std::size_t) {
+    // ensure that we did get the right inputs forwarded
+    _val = *static_cast<int*>(ptr);
+  }
+
+  void* allocate_async(std::size_t, std::size_t, cuda::stream_ref) {
+    return &_val;
+  }
+
+  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);
+  }
+
+  bool operator==(const async_resource& other) const {
+    return _val == other._val;
+  }
+  bool operator!=(const async_resource& other) const {
+    return _val != other._val;
+  }
+
+  int _val = 0;
+};
+
+void test_allocate() {
+  { // allocate(size)
+    async_resource input{42};
+    cuda::mr::async_resource_ref<> ref{input};
+
+    // Ensure that we properly pass on the allocate function
+    assert(input.allocate(0, 0) == ref.allocate(0));
+
+    int expected_after_deallocate = 1337;
+    ref.deallocate(static_cast<void*>(&expected_after_deallocate), 0);
+    assert(input._val == expected_after_deallocate);
+  }
+
+  { // allocate(size, alignment)
+    async_resource input{42};
+    cuda::mr::async_resource_ref<> ref{input};
+
+    // Ensure that we properly pass on the allocate function
+    assert(input.allocate(0, 0) == ref.allocate(0, 0));
+
+    int expected_after_deallocate = 1337;
+    ref.deallocate(static_cast<void*>(&expected_after_deallocate), 0, 0);
+    assert(input._val == expected_after_deallocate);
+  }
+}
+
+void test_allocate_async() {
+  { // allocate(size)
+    async_resource input{42};
+    cuda::mr::async_resource_ref<> ref{input};
+
+    // Ensure that we properly pass on the allocate function
+    assert(input.allocate_async(0, 0, {}) == ref.allocate_async(0, {}));
+
+    int expected_after_deallocate = 1337;
+    ref.deallocate_async(static_cast<void*>(&expected_after_deallocate), 0, {});
+    assert(input._val == expected_after_deallocate);
+  }
+
+  { // allocate(size, alignment)
+    async_resource input{42};
+    cuda::mr::async_resource_ref<> ref{input};
+
+    // Ensure that we properly pass on the allocate function
+    assert(input.allocate_async(0, 0, {}) == ref.allocate_async(0, 0, {}));
+
+    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**) {
+#ifndef __CUDA_ARCH__
+  test_allocate();
+  test_allocate_async();
+#endif
+  return 0;
+}