From 27253d72db83ce3e4e1d7f8cd6fa0c41ea38bab6 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 31 Jul 2024 15:01:32 +0200 Subject: [PATCH] [PoC]: Implement `cuda::experimental::uninitialized_buffer` (#1831) * Drop `cuda::get_property` CPO It serves no purpose as it only ever forwards via ADL and also breaks older nvcc * Ensure that we test memory resources * Implement `cuda::uninitialized_buffer` `cuda::uninitialized_buffer` provides an allocation of `N` elements of type `T` utilitzing a `cuda::mr::resource` to allocate the storage. `cuda::uninitialized_buffer` takes care of alignment and deallocation of the storage. The user is required to ensure that the lifetime of the memory resource exceeds the lifetime of the buffer. --- cudax/cmake/cudaxHeaderTesting.cmake | 3 +- .../__container/uninitialized_buffer.h | 203 ++++++++++++++++++ cudax/include/cuda/experimental/buffer | 25 +++ cudax/test/CMakeLists.txt | 11 +- cudax/test/containers/uninitialized_buffer.cu | 129 +++++++++++ docs/cudax/container.rst | 25 +++ docs/cudax/index.rst | 30 +++ docs/index.rst | 3 + docs/repo.toml | 59 ++++- .../cuda/__memory_resource/get_property.h | 20 -- libcudacxx/include/cuda/std/__cccl/compiler.h | 10 +- .../test/internal_headers/CMakeLists.txt | 3 +- libcudacxx/test/libcudacxx/CMakeLists.txt | 2 +- .../allocate.pass.cpp | 1 - .../equality.pass.cpp | 1 - .../traits.pass.cpp | 1 - .../cuda_memory_resource/allocate.pass.cpp | 1 - .../cuda_memory_resource/equality.pass.cpp | 1 - .../cuda_memory_resource/traits.pass.cpp | 1 - .../allocate.pass.cpp | 1 - .../equality.pass.cpp | 1 - .../traits.pass.cpp | 1 - .../get_property/forward_property.pass.cpp | 1 - .../get_property/get_property.pass.cpp | 24 +-- .../get_property/has_property.pass.cpp | 1 - .../async_resource_ref.allocate.pass.cpp | 1 - .../async_resource_ref.construction.pass.cpp | 1 - .../async_resource_ref.conversion.pass.cpp | 1 - .../async_resource_ref.equality.fail.cpp | 1 - .../async_resource_ref.equality.pass.cpp | 1 - .../async_resource_ref.inheritance.pass.cpp | 1 - .../async_resource_ref.properties.pass.cpp | 1 - .../async_resource.pass.cpp | 1 - .../async_resource_with.pass.cpp | 1 - .../resource.pass.cpp | 1 - .../resource_with.pass.cpp | 1 - .../resource_ref.allocate.pass.cpp | 1 - .../resource_ref.construction.pass.cpp | 1 - .../resource_ref.conversion.pass.cpp | 1 - .../resource_ref.equality.fail.cpp | 1 - .../resource_ref.equality.pass.cpp | 1 - .../resource_ref.inheritance.pass.cpp | 1 - .../resource_ref.properties.pass.cpp | 1 - libcudacxx/test/public_headers/CMakeLists.txt | 5 +- 44 files changed, 505 insertions(+), 76 deletions(-) create mode 100644 cudax/include/cuda/experimental/__container/uninitialized_buffer.h create mode 100644 cudax/include/cuda/experimental/buffer create mode 100644 cudax/test/containers/uninitialized_buffer.cu create mode 100644 docs/cudax/container.rst create mode 100644 docs/cudax/index.rst diff --git a/cudax/cmake/cudaxHeaderTesting.cmake b/cudax/cmake/cudaxHeaderTesting.cmake index 00be5587e6..29a3bd58ca 100644 --- a/cudax/cmake/cudaxHeaderTesting.cmake +++ b/cudax/cmake/cudaxHeaderTesting.cmake @@ -28,7 +28,8 @@ function(cudax_add_header_test label definitions) set(headertest_target ${config_prefix}.headers.${label}) add_library(${headertest_target} OBJECT ${headertest_srcs}) target_link_libraries(${headertest_target} PUBLIC ${cn_target}) - target_compile_definitions(${headertest_target} PRIVATE ${definitions}) + target_compile_definitions(${headertest_target} PRIVATE ${definitions} + "-DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE") cudax_clone_target_properties(${headertest_target} ${cn_target}) add_dependencies(cudax.all.headers ${headertest_target}) diff --git a/cudax/include/cuda/experimental/__container/uninitialized_buffer.h b/cudax/include/cuda/experimental/__container/uninitialized_buffer.h new file mode 100644 index 0000000000..7415db6c85 --- /dev/null +++ b/cudax/include/cuda/experimental/__container/uninitialized_buffer.h @@ -0,0 +1,203 @@ +//===----------------------------------------------------------------------===// +// +// Part of the CUDA Toolkit, 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) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef __CUDAX__CONTAINERS_UNINITIALIZED_BUFFER_H +#define __CUDAX__CONTAINERS_UNINITIALIZED_BUFFER_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include + +#if _CCCL_STD_VER >= 2014 && !defined(_CCCL_COMPILER_MSVC_2017) \ + && defined(LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE) + +//! @file The \c uninitialized_buffer class provides a typed buffer allocated from a given memory resource. +namespace cuda::experimental +{ + +//! @rst +//! .. _cudax-containers-uninitialized-buffer: +//! +//! Uninitialized type safe memory storage +//! --------------------------------------- +//! +//! ``uninitialized_buffer`` provides a typed buffer allocated from a given :ref:`memory resource +//! `. It handles alignment and release of the allocation. +//! The memory is uninitialized, so that a user needs to ensure elements are properly constructed. +//! +//! In addition to being type safe, ``uninitialized_buffer`` also takes a set of :ref:`properties +//! ` to ensure that e.g. execution space constraints are checked +//! at compile time. However, we can only forward stateless properties. If a user wants to use a stateful one, then they +//! need to implement :ref:`get_property(const device_buffer&, Property) +//! `. +//! +//! .. warning:: +//! +//! ``uninitialized_buffer`` stores a reference to the provided memory :ref:`memory resource +//! `. It is the user's resposibility to ensure the lifetime of +//! the resource exceeds the lifetime of the buffer. +//! +//! @endrst +//! @tparam _T the type to be stored in the buffer +//! @tparam _Properties... The properties the allocated memory satisfies +template +class uninitialized_buffer +{ +private: + _CUDA_VMR::resource_ref<_Properties...> __mr_; + size_t __count_ = 0; + void* __buf_ = nullptr; + + //! @brief Determines the allocation size given the alignment and size of `T` + _CCCL_NODISCARD _CCCL_HOST_DEVICE static constexpr size_t __get_allocation_size(const size_t __count) noexcept + { + constexpr size_t __alignment = alignof(_Tp); + return (__count * sizeof(_Tp) + (__alignment - 1)) & ~(__alignment - 1); + } + + //! @brief Determines the properly aligned start of the buffer given the alignment and size of `T` + _CCCL_NODISCARD _CCCL_HOST_DEVICE _Tp* __get_data() const noexcept + { + constexpr size_t __alignment = alignof(_Tp); + size_t __space = __get_allocation_size(__count_); + void* __ptr = __buf_; + return _CUDA_VSTD::launder( + reinterpret_cast<_Tp*>(_CUDA_VSTD::align(__alignment, __count_ * sizeof(_Tp), __ptr, __space))); + } + +public: + using value_type = _Tp; + using reference = _Tp&; + using pointer = _Tp*; + using size_type = size_t; + + //! @brief Constructs a \c uninitialized_buffer, allocating sufficient storage for \p __count elements through \p __mr + //! @param __mr The memory resource to allocate the buffer with. + //! @param __count The desired size of the buffer. + //! @note Depending on the alignment requirements of `T` the size of the underlying allocation might be larger + //! than `count * sizeof(T)`. + //! @note Only allocates memory when \p __count > 0 + uninitialized_buffer(_CUDA_VMR::resource_ref<_Properties...> __mr, const size_t __count) + : __mr_(__mr) + , __count_(__count) + , __buf_(__count_ == 0 ? nullptr : __mr_.allocate(__get_allocation_size(__count_))) + {} + + uninitialized_buffer(const uninitialized_buffer&) = delete; + uninitialized_buffer& operator=(const uninitialized_buffer&) = delete; + + //! @brief Move construction + //! @param __other Another \c uninitialized_buffer + uninitialized_buffer(uninitialized_buffer&& __other) noexcept + : __mr_(__other.__mr_) + , __count_(__other.__count_) + , __buf_(__other.__buf_) + { + __other.__count_ = 0; + __other.__buf_ = nullptr; + } + + //! @brief Move assignment + //! @param __other Another \c uninitialized_buffer + uninitialized_buffer& operator=(uninitialized_buffer&& __other) noexcept + { + if (__buf_) + { + __mr_.deallocate(__buf_, __get_allocation_size(__count_)); + } + __mr_ = __other.__mr_; + __count_ = __other.__count_; + __buf_ = __other.__buf_; + __other.__count_ = 0; + __other.__buf_ = nullptr; + return *this; + } + + //! @brief Destroys an \c uninitialized_buffer deallocating the buffer + //! @warning The destructor does not destroy any objects that may or may not reside within the buffer. It is the + //! user's responsibility to ensure that all objects within the buffer have been properly destroyed. + ~uninitialized_buffer() + { + if (__buf_) + { + __mr_.deallocate(__buf_, __get_allocation_size(__count_)); + } + } + + //! @brief Returns an aligned pointer to the buffer + _CCCL_NODISCARD _CCCL_HOST_DEVICE pointer begin() const noexcept + { + return __get_data(); + } + + //! @brief Returns an aligned pointer to end of the buffer + _CCCL_NODISCARD _CCCL_HOST_DEVICE pointer end() const noexcept + { + return __get_data() + __count_; + } + + //! @brief Returns an aligned pointer to the buffer + _CCCL_NODISCARD _CCCL_HOST_DEVICE pointer data() const noexcept + { + return __get_data(); + } + + //! @brief Returns the size of the buffer + _CCCL_NODISCARD _CCCL_HOST_DEVICE constexpr size_type size() const noexcept + { + return __count_; + } + + //! @rst + //! Returns the :ref:`resource_ref ` used to allocate + //! the buffer + //! @endrst + _CCCL_NODISCARD _CCCL_HOST_DEVICE _CUDA_VMR::resource_ref<_Properties...> resource() const noexcept + { + return __mr_; + } + + //! @brief Swaps the contents with those of another \c uninitialized_buffer + //! @param __other The other \c uninitialized_buffer. + _CCCL_HOST_DEVICE constexpr void swap(uninitialized_buffer& __other) noexcept + { + _CUDA_VSTD::swap(__mr_, __other.__mr_); + _CUDA_VSTD::swap(__count_, __other.__count_); + _CUDA_VSTD::swap(__buf_, __other.__buf_); + } + +# ifndef DOXYGEN_SHOULD_SKIP_THIS // friend functions are currently brocken + //! @brief Forwards the passed Properties + _LIBCUDACXX_TEMPLATE(class _Property) + _LIBCUDACXX_REQUIRES((!property_with_value<_Property>) _LIBCUDACXX_AND _CUDA_VSTD::_One_of<_Property, _Properties...>) + friend constexpr void get_property(const uninitialized_buffer&, _Property) noexcept {} +# endif // DOXYGEN_SHOULD_SKIP_THIS +}; + +template +using uninitialized_device_buffer = uninitialized_buffer<_Tp, _CUDA_VMR::device_accessible>; + +} // namespace cuda::experimental + +#endif // _CCCL_STD_VER >= 2014 && !_CCCL_COMPILER_MSVC_2017 && LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE + +#endif //__CUDAX__CONTAINERS_UNINITIALIZED_BUFFER_H diff --git a/cudax/include/cuda/experimental/buffer b/cudax/include/cuda/experimental/buffer new file mode 100644 index 0000000000..2d9005f446 --- /dev/null +++ b/cudax/include/cuda/experimental/buffer @@ -0,0 +1,25 @@ +//===----------------------------------------------------------------------===// +// +// Part of the CUDA Toolkit, 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) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef __CUDAX_BUFFER +#define __CUDAX_BUFFER + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include + +#endif //_CUDA_BUFFER diff --git a/cudax/test/CMakeLists.txt b/cudax/test/CMakeLists.txt index e0c73a7b76..cb76d492e8 100644 --- a/cudax/test/CMakeLists.txt +++ b/cudax/test/CMakeLists.txt @@ -4,6 +4,9 @@ CPMAddPackage("gh:catchorg/Catch2@2.13.9") add_library(catch2_main STATIC catch2_helpers/catch2_main.cpp) target_link_libraries(catch2_main PUBLIC Catch2::Catch2) +find_package(Thrust REQUIRED CONFIG) +thrust_create_target(cudax::Thrust) + ## cudax_add_test # # Add a catch2 test executable and register it with ctest. @@ -24,6 +27,9 @@ function(cudax_add_catch2_test target_name_var test_name cn_target) # ARGN=test add_executable(${test_target} ${test_sources}) target_link_libraries(${test_target} PRIVATE ${cn_target} Catch2::Catch2 catch2_main) + target_link_libraries(${test_target} PRIVATE ${cn_target} cudax::Thrust) + target_compile_options(${test_target} PRIVATE "-DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE") + target_compile_options(${test_target} PRIVATE $<$:--extended-lambda>) cudax_clone_target_properties(${test_target} ${cn_target}) set_target_properties(${test_target} PROPERTIES CUDA_ARCHITECTURES "${CMAKE_CUDA_ARCHITECTURES}" @@ -56,7 +62,6 @@ foreach(cn_target IN LISTS cudax_TARGETS) launch/launch_smoke.cu launch/configuration.cu ) - target_compile_options(${test_target} PRIVATE $<$:--extended-lambda>) Cudax_add_catch2_test(test_target device_tests ${cn_target} device/device_smoke.cu @@ -75,4 +80,8 @@ foreach(cn_target IN LISTS cudax_TARGETS) cudax_add_catch2_test(test_target misc_tests ${cn_target} utility/driver_api.cu ) + + cudax_add_catch2_test(test_target containers ${cn_target} + containers/uninitialized_buffer.cu + ) endforeach() diff --git a/cudax/test/containers/uninitialized_buffer.cu b/cudax/test/containers/uninitialized_buffer.cu new file mode 100644 index 0000000000..b50527ed17 --- /dev/null +++ b/cudax/test/containers/uninitialized_buffer.cu @@ -0,0 +1,129 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// 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) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include +#include +#include + +#include +#include +#include +#include +#include + +#include + +#include + +struct do_not_construct +{ + do_not_construct() + { + CHECK(false); + } +}; + +struct non_trivial +{ + int val_ = 0; + non_trivial() = default; + __host__ __device__ constexpr non_trivial(const int val) noexcept + : val_(val) + {} + + __host__ __device__ constexpr friend bool operator==(const non_trivial& lhs, const non_trivial& rhs) + { + return lhs.val_ == rhs.val_; + } +}; + +struct my_property +{ + using value_type = int; +}; +constexpr int get_property(const cuda::experimental::uninitialized_buffer&, my_property) +{ + return 42; +} + +TEMPLATE_TEST_CASE( + "uninitialized_buffer", "[container]", char, short, int, long, long long, float, double, do_not_construct) +{ + using uninitialized_buffer = cuda::experimental::uninitialized_buffer; + static_assert(!cuda::std::is_default_constructible::value, ""); + static_assert(!cuda::std::is_copy_constructible::value, ""); + static_assert(!cuda::std::is_copy_assignable::value, ""); + + cuda::mr::cuda_memory_resource resource{}; + + SECTION("construction") + { + { + uninitialized_buffer from_count{resource, 42}; + CHECK(from_count.data() != nullptr); + CHECK(from_count.size() == 42); + } + { + uninitialized_buffer input{resource, 42}; + const TestType* ptr = input.data(); + + uninitialized_buffer from_rvalue{cuda::std::move(input)}; + CHECK(from_rvalue.data() == ptr); + CHECK(from_rvalue.size() == 42); + + // Ensure that we properly reset the input buffer + CHECK(input.data() == nullptr); + CHECK(input.size() == 0); + } + } + + SECTION("access") + { + uninitialized_buffer buf{resource, 42}; + CHECK(buf.data() != nullptr); + CHECK(buf.size() == 42); + CHECK(buf.begin() == buf.data()); + CHECK(buf.end() == buf.begin() + buf.size()); + CHECK(buf.resource() == resource); + + CHECK(cuda::std::as_const(buf).data() != nullptr); + CHECK(cuda::std::as_const(buf).size() == 42); + CHECK(cuda::std::as_const(buf).begin() == buf.data()); + CHECK(cuda::std::as_const(buf).end() == buf.begin() + buf.size()); + CHECK(cuda::std::as_const(buf).resource() == resource); + } + + SECTION("properties") + { + static_assert(cuda::has_property, + cuda::mr::device_accessible>, + ""); + static_assert(cuda::has_property, my_property>, ""); + } + + SECTION("convertion to span") + { + uninitialized_buffer buf{resource, 42}; + const cuda::std::span as_span{buf}; + CHECK(as_span.data() == buf.data()); + CHECK(as_span.size() == 42); + } + + SECTION("Actually use memory") + { + if constexpr (!cuda::std::is_same_v) + { + uninitialized_buffer buf{resource, 42}; + thrust::fill(thrust::device, buf.begin(), buf.end(), TestType{2}); + const auto res = thrust::reduce(thrust::device, buf.begin(), buf.end(), TestType{0}, thrust::plus()); + CHECK(res == TestType{84}); + } + } +} diff --git a/docs/cudax/container.rst b/docs/cudax/container.rst new file mode 100644 index 0000000000..3a0ab05f45 --- /dev/null +++ b/docs/cudax/container.rst @@ -0,0 +1,25 @@ +.. _cudax-containers: + +Containers library +=================== + +.. toctree:: + :glob: + :maxdepth: 1 + + ${repo_docs_api_path}/class*uninitialized__buffer* + +The headers of the container library provide facilities to store elements on the heap. They are heavily inspired by the +C++ `containers library `__ but deviate from the standard provided ones due to different requirements from +heterogeneous systems. + +They build upon :ref:`memory_resources ` to ensure that e.g. execution space +annotations are checked by the type system. + +.. list-table:: + :widths: 25 45 30 + :header-rows: 0 + + * - :ref:` ` + - Facilities providing uninitialized *heterogeneous* storage satisfying a set of properties + - cudax 2.7.0 / CCCL 2.7.0 diff --git a/docs/cudax/index.rst b/docs/cudax/index.rst new file mode 100644 index 0000000000..9b066b61c0 --- /dev/null +++ b/docs/cudax/index.rst @@ -0,0 +1,30 @@ +.. _cudax-module: + +CUDA Experimental +================= + +.. toctree:: + :hidden: + :maxdepth: 3 + + container + ${repo_docs_api_path}/cudax_api + +``CUDA Experimental`` (``cudax``) provides experimental new features that are still in development and subject to change. +However, any feature within this library has important use cases and we encourage users to experiment with them. + +Specifically, ``cudax`` provides: + - :ref:`uninitialized storage ` + - dimensions description functionality + +Stability Guarantees +--------------------- + +There are no stability guarantees whatsoever. We reserve the right to change both the ABI and the API of any feature +within ``cudax`` at any time without notice. + +Availability +------------- + +Due to its experimental nature and the lack of stability guarantees, ``cudax`` is not shipped with the CUDA toolkit but +is solely available through github. diff --git a/docs/index.rst b/docs/index.rst index 9fb4b82366..09355eafe1 100644 --- a/docs/index.rst +++ b/docs/index.rst @@ -8,6 +8,7 @@ CUDA C++ Core Libraries libcu++ CUB Thrust + Cuda Experimental Welcome to the CUDA C++ Core Libraries (CCCL) where our mission is to make CUDA C++ more delightful. @@ -20,5 +21,7 @@ Naturally, there was a lot of overlap among the three projects, and it became cl - `Thrust `__ is the C++ parallel algorithms library which inspired the introduction of parallel algorithms to the C++ Standard Library. Thrust's high-level interface greatly enhances programmer productivity while enabling performance portability between GPUs and multicore CPUs via configurable backends that allow using multiple parallel programming frameworks (such as CUDA, TBB, and OpenMP). +- `Cuda Experimental `__ is a library of exerimental features that are still in the design process. + The main goal of CCCL is to fill a similar role that the Standard C++ Library fills for Standard C++: provide general-purpose, speed-of-light tools to CUDA C++ developers, allowing them to focus on solving the problems that matter. Unifying these projects is the first step towards realizing that goal. diff --git a/docs/repo.toml b/docs/repo.toml index 81fddea229..74ebb0be7d 100644 --- a/docs/repo.toml +++ b/docs/repo.toml @@ -25,7 +25,7 @@ sphinx_exclude_patterns = [ "VERSION.md", ] -project_build_order = [ "libcudacxx", "cub", "thrust", "cccl", "pycuda" ] +project_build_order = [ "libcudacxx", "cudax", "cub", "thrust", "cccl", "pycuda" ] # deps can be used to link to other projects' documentation deps = [ @@ -303,3 +303,60 @@ enhanced_search_enabled = true python_paths = [ "${root}/../python/cuda" ] + +[repo_docs.projects.cudax] +name = "Cudax: Experimental new features" +docs_root = "cudax" +logo = "../img/logo.png" + +repo_url = "https://github.com/NVIDIA/cccl/cudax" +social_media_set = "" +social_media = [ + [ "github", "https://github.com/NVIDIA/cccl" ], +] + +enhanced_search_enabled = true +api_output_directory = "api" +use_fast_doxygen_conversion = true +sphinx_generate_doxygen_groups = true +sphinx_generate_doxygen_pages = true +sphinx_exclude_patterns = [] + +# deps can be used to link to other projects' documentation +deps = [ + [ "libcudacxx", "_build/docs/libcudacxx/latest" ], +] + +doxygen_input = [ + "../../cudax/include/cuda/experimental/__container/*.h", + "../../cudax/include/cuda/experimental/__memory_resource/*.h", +] + +doxygen_predefined = [ + "_CCCL_DEVICE=", + "_CCCL_EXEC_CHECK_DISABLE=", + "_CCCL_FORCEINLINE=", + "_CCCL_HOST=", + "_CCCL_HOST_DEVICE=", + "_CCCL_NODISCARD=[[nodiscard]]", + "_CCCL_NODISCARD_FRIEND=", + "_CCCL_STD_VER=2020", + "_CUDA_VMR=cuda::mr", + "_CUDA_VSTD=cuda::std", + "LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE=", + "_LIBCUDACXX_EAT_REST(x)=", + "_LIBCUDACXX_TRAILING_REQUIRES(x)=-> x _LIBCUDACXX_EAT_REST", +] + +# make sure to use ./fetch_imgs.sh +doxygen_conf_extra = """ + IMAGE_PATH = ${config_root}/img + DISTRIBUTE_GROUP_DOC = Yes + DOXYFILE_ENCODING = UTF-8 + INPUT_ENCODING = UTF-8 + EXTENSION_MAPPING = cuh=c++ cu=c++ + EXAMPLE_RECURSIVE = NO + EXAMPLE_PATTERNS = *.cu + EXCLUDE_SYMBOLS = "__*" + AUTOLINK_SUPPORT = YES +""" diff --git a/libcudacxx/include/cuda/__memory_resource/get_property.h b/libcudacxx/include/cuda/__memory_resource/get_property.h index 4208d07232..d39c6ac0d3 100644 --- a/libcudacxx/include/cuda/__memory_resource/get_property.h +++ b/libcudacxx/include/cuda/__memory_resource/get_property.h @@ -159,26 +159,6 @@ _LIBCUDACXX_END_NAMESPACE_CPO template using forward_property = __forward_property::__fn<_Derived, _Upstream>; -_LIBCUDACXX_BEGIN_NAMESPACE_CPO(__get_property) -void get_property(); - -struct __fn -{ - _CCCL_EXEC_CHECK_DISABLE - template - _LIBCUDACXX_INLINE_VISIBILITY constexpr auto operator()(const _Upstream& __res, _Property __prop) const - noexcept(noexcept(get_property(__res, __prop))) -> decltype(get_property(__res, __prop)) - { - return get_property(__res, __prop); - } -}; -_LIBCUDACXX_END_NAMESPACE_CPO - -inline namespace __cpo -{ -_CCCL_GLOBAL_CONSTANT auto get_property = __get_property::__fn{}; -} // namespace __cpo - _LIBCUDACXX_END_NAMESPACE_CUDA # endif // _CCCL_STD_VER >= 2014 diff --git a/libcudacxx/include/cuda/std/__cccl/compiler.h b/libcudacxx/include/cuda/std/__cccl/compiler.h index 80c20a9bc1..d4d68a86f2 100644 --- a/libcudacxx/include/cuda/std/__cccl/compiler.h +++ b/libcudacxx/include/cuda/std/__cccl/compiler.h @@ -72,19 +72,19 @@ #endif // __CUDACC__ || _CCCL_CUDA_COMPILER_NVHPC // Some convenience macros to filter CUDACC versions -#if defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1102000 +#if !defined(_CCCL_CUDA_COMPILER) || (defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1102000) # define _CCCL_CUDACC_BELOW_11_2 #endif // defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1102000 -#if defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1103000 +#if !defined(_CCCL_CUDA_COMPILER) || (defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1103000) # define _CCCL_CUDACC_BELOW_11_3 #endif // defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1103000 -#if defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1108000 +#if !defined(_CCCL_CUDA_COMPILER) || (defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1108000) # define _CCCL_CUDACC_BELOW_11_8 #endif // defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1108000 -#if defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1202000 +#if !defined(_CCCL_CUDA_COMPILER) || (defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1202000) # define _CCCL_CUDACC_BELOW_12_2 #endif // defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1203000 -#if defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1203000 +#if !defined(_CCCL_CUDA_COMPILER) || (defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1203000) # define _CCCL_CUDACC_BELOW_12_3 #endif // defined(_CCCL_CUDACC) && _CCCL_CUDACC_VER < 1203000 diff --git a/libcudacxx/test/internal_headers/CMakeLists.txt b/libcudacxx/test/internal_headers/CMakeLists.txt index 08cdb49350..c9c060c2dd 100644 --- a/libcudacxx/test/internal_headers/CMakeLists.txt +++ b/libcudacxx/test/internal_headers/CMakeLists.txt @@ -39,7 +39,8 @@ function(libcudacxx_create_internal_header_test header_name, headertest_src, fal target_compile_options(headertest_${header_name} PRIVATE $<$:${headertest_warning_levels_device}> - $<$:${headertest_warning_levels_host}>) + $<$:${headertest_warning_levels_host}> + -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE) target_link_libraries(headertest_${header_name} CUDA::cudart) if(fallback) diff --git a/libcudacxx/test/libcudacxx/CMakeLists.txt b/libcudacxx/test/libcudacxx/CMakeLists.txt index dad2ce9d91..f4c3b038f8 100644 --- a/libcudacxx/test/libcudacxx/CMakeLists.txt +++ b/libcudacxx/test/libcudacxx/CMakeLists.txt @@ -36,7 +36,7 @@ else() # NOT LIBCUDACXX_TEST_WITH_NVRTC set(LIBCUDACXX_FORCE_INCLUDE "-include ${libcudacxx_SOURCE_DIR}/test/libcudacxx/force_include.h") set(LIBCUDACXX_CUDA_COMPILER "${CMAKE_CUDA_COMPILER}") set(LIBCUDACXX_CUDA_TEST_WITH_NVRTC "False") - set(LIBCUDACXX_TEST_COMPILER_FLAGS "") + set(LIBCUDACXX_TEST_COMPILER_FLAGS "-DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE") endif() # enable exceptions in tests diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/allocate.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/allocate.pass.cpp index 85e365f3dc..6e9fd76f8e 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/allocate.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/allocate.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/equality.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/equality.pass.cpp index 837d82c98c..80fb2ab57e 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/equality.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/equality.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/traits.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/traits.pass.cpp index 07582b59fd..299247ff2e 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/traits.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_managed_memory_resource/traits.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/allocate.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/allocate.pass.cpp index 19a251b315..2c88483e6c 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/allocate.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/allocate.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/equality.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/equality.pass.cpp index d28b22e87b..94d659f90f 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/equality.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/equality.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/traits.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/traits.pass.cpp index c7fbc997f7..8229f4eaea 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/traits.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_memory_resource/traits.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/allocate.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/allocate.pass.cpp index 3d384068f6..7b9e374805 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/allocate.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/allocate.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/equality.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/equality.pass.cpp index 477cc6a466..7cab309a33 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/equality.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/equality.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/traits.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/traits.pass.cpp index 77e0522325..3909ac7238 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/traits.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/cuda_pinned_memory_resource/traits.pass.cpp @@ -11,7 +11,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/forward_property.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/forward_property.pass.cpp index 4c124b08bd..358913c7b9 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/forward_property.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/forward_property.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::forward_property #include diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/get_property.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/get_property.pass.cpp index a6aa889641..1e6ef7fa51 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/get_property.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/get_property.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::get_property @@ -28,44 +27,37 @@ struct prop struct upstream_with_valueless_property { - friend constexpr void get_property(const upstream_with_valueless_property&, prop) {} + __host__ __device__ friend constexpr void get_property(const upstream_with_valueless_property&, prop) {} }; -static_assert(cuda::std::invocable, ""); -static_assert(!cuda::std::invocable, - ""); struct upstream_with_stateful_property { - friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) + __host__ __device__ friend constexpr int get_property(const upstream_with_stateful_property&, prop_with_value) { return 42; } }; -static_assert(!cuda::std::invocable, ""); -static_assert(cuda::std::invocable, ""); 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) + __host__ __device__ friend constexpr void get_property(const upstream_with_both_properties&, prop) {} + __host__ __device__ friend constexpr int get_property(const upstream_with_both_properties&, prop_with_value) { return 42; } }; -static_assert(cuda::std::invocable, ""); -static_assert(cuda::std::invocable, ""); __host__ __device__ constexpr bool test() { upstream_with_valueless_property with_valueless{}; - cuda::get_property(with_valueless, prop{}); + get_property(with_valueless, prop{}); upstream_with_stateful_property with_value{}; - assert(cuda::get_property(with_value, prop_with_value{}) == 42); + assert(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); + get_property(with_both, prop{}); + assert(get_property(with_both, prop_with_value{}) == 42); return true; } diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/has_property.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/has_property.pass.cpp index de32624357..7d769e69bc 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/has_property.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/get_property/has_property.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::has_property, cuda::has_property_with diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp index 9bd69a381b..266ed7965a 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.allocate.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp index 92190272e9..c6dee19660 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.construction.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource_ref construction diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp index 3a2cc16a0d..1b081bde05 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.conversion.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp index b33206ff0e..08582c00fb 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.fail.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource_ref equality diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp index 9140a0eacd..3c5b992c2e 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.equality.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource_ref equality diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp index 2fed448857..2f5b125c78 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.inheritance.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp index d9680365c2..4db4074e26 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.async_resource_ref/async_resource_ref.properties.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp index c97ec6dc48..dcfd09b9c2 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp index 1ff360bb6b..78734d0cac 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/async_resource_with.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::async_resource_with diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp index 876ffc5b58..bd789d9bc4 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp index 84c544cfbf..386de3da32 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.concepts/resource_with.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_with diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp index e12761e762..5dbcac073e 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.allocate.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp index 9e86f30eb8..2d76172d35 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.construction.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref construction diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp index f2dfea34a4..7e5e23d8a2 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.conversion.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp index eca3afe0c2..8d51bb2177 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.fail.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref equality diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp index 1fa107995c..e835aed39a 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.equality.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref equality diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp index c99fd2f14d..4fcc1d5df6 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.inheritance.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref properties diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp index b099915d22..b4119a1227 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/memory_resource.resource_ref/resource_ref.properties.pass.cpp @@ -10,7 +10,6 @@ // UNSUPPORTED: c++03, c++11 // UNSUPPORTED: msvc-19.16 // UNSUPPORTED: nvrtc -#define LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE // cuda::mr::resource_ref properties diff --git a/libcudacxx/test/public_headers/CMakeLists.txt b/libcudacxx/test/public_headers/CMakeLists.txt index 231a8a5cb8..df84e5c95c 100644 --- a/libcudacxx/test/public_headers/CMakeLists.txt +++ b/libcudacxx/test/public_headers/CMakeLists.txt @@ -42,7 +42,10 @@ function(libcudacxx_add_public_header_test header) set(headertest_${header_name} verify_${header_name}) add_library(headertest_${header_name} SHARED "${headertest_src}.cu") target_include_directories(headertest_${header_name} PRIVATE "${libcudacxx_SOURCE_DIR}/include") - target_compile_options(headertest_${header_name} PRIVATE ${headertest_warning_levels_device}) + target_compile_options(headertest_${header_name} + PRIVATE + ${headertest_warning_levels_device} + -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE) # Ensure that if this is an atomic header, we only include the right architectures string(REGEX MATCH "atomic|barrier|latch|semaphore|annotated_ptr|pipeline" match "${header}")