diff --git a/CHANGELOG.md b/CHANGELOG.md index 3795a23463..aba78d4b6a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,8 +1,33 @@ -# Thrust 1.10.0 (NVIDIA HPC SDK 20.9) +# Thrust 1.11.0 (NVIDIA HPC SDK 20.11) + +Thrust 1.11.0 is the major release accompanying the NVIDIA HPC SDK 20.11 release. +It exposes common Thrust abstractions for dealing with memory that can be + accessed by both hosts and devices: `thrust::universal_vector`, + `thrust::universal_ptr`, and `thrust::universal_allocator`. +It also adds `thrust::async::inclusive_scan` and + `thrust::async::exclusive_scan`. +Finally, it brings substantial improvements to the performance of the sort + backend for CUDA. + +## New Features + +- `thrust::universal_allocator`: An allocator which allocates memory that is + accessible to both hosts and devices. +- `thrust::universal_ptr`: A tagged pointer to memory that is accessible to + both hosts and devices. +- `thrust::universal_vector`: A dynamically-allocated array of objects + accessible to both hosts and devices. + +## Other Enhancements + +- All backend `vector`, `pointer`, and `reference` types are now type aliases. + +# Thrust 1.10.0 (NVIDIA HPC SDK 20.9, CUDA Toolkit 11.2) ## Summary -Thrust 1.10.0 is the major release accompanying the NVIDIA HPC SDK 20.9 release. +Thrust 1.10.0 is the major release accompanying the NVIDIA HPC SDK 20.9 release + and the CUDA Toolkit 11.2 release. It drops support for C++03, GCC < 5, Clang < 6, and MSVC < 2017. It also overhauls CMake support. Finally, we now have a Code of Conduct for contributors: diff --git a/testing/cuda/managed_memory_pointer.cu b/testing/cuda/managed_memory_pointer.cu index 46a2191fab..7d6ad907c3 100644 --- a/testing/cuda/managed_memory_pointer.cu +++ b/testing/cuda/managed_memory_pointer.cu @@ -1,141 +1,19 @@ -#include +#include -#if THRUST_CPP_DIALECT >= 2011 +#include -# include - -# include -# include -# include -# include - -# include -# include - -namespace -{ - -template -using allocator = - thrust::mr::stateless_resource_allocator; - -// The managed_memory_pointer class should be identified as a -// contiguous_iterator -THRUST_STATIC_ASSERT( - thrust::is_contiguous_iterator::pointer>::value); +#include template -struct some_object { - some_object(T data) - : m_data(data) - {} - - void setter(T data) { m_data = data; } - T getter() const { return m_data; } - -private: - T m_data; -}; - -} // namespace - -template -void TestAllocateUnique() +void TestCudaManagedMemoryPointer() { - // Simple test to ensure that pointers created with universal_memory_resource - // can be dereferenced and used with STL code. This is necessary as some - // STL implementations break when using fancy references that overload - // `operator&`, so universal_memory_resource uses a special pointer type that - // returns regular C++ references that can be safely used host-side. - - // These operations fail to compile with fancy references: - auto pRaw = thrust::allocate_unique(allocator{}, 42); - auto pObj = - thrust::allocate_unique >(allocator >{}, 42); - - static_assert( - std::is_same >::value, - "Unexpected pointer returned from unique_ptr::get."); static_assert( - std::is_same > >::value, - "Unexpected pointer returned from unique_ptr::get."); - - ASSERT_EQUAL(*pRaw, T(42)); - ASSERT_EQUAL(*pRaw.get(), T(42)); - ASSERT_EQUAL(pObj->getter(), T(42)); - ASSERT_EQUAL((*pObj).getter(), T(42)); - ASSERT_EQUAL(pObj.get()->getter(), T(42)); - ASSERT_EQUAL((*pObj.get()).getter(), T(42)); -} -DECLARE_GENERIC_UNITTEST(TestAllocateUnique); - -template -void TestIterationRaw() -{ - auto array = thrust::allocate_unique_n(allocator{}, 6, 42); - - static_assert( - std::is_same >::value, - "Unexpected pointer returned from unique_ptr::get."); - - for (auto iter = array.get(), end = array.get() + 6; iter < end; ++iter) - { - ASSERT_EQUAL(*iter, T(42)); - ASSERT_EQUAL(*iter.get(), T(42)); - } -} -DECLARE_GENERIC_UNITTEST(TestIterationRaw); - -template -void TestIterationObj() -{ - auto array = - thrust::allocate_unique_n >(allocator >{}, - 6, - 42); - - static_assert( - std::is_same > >::value, - "Unexpected pointer returned from unique_ptr::get."); - - for (auto iter = array.get(), end = array.get() + 6; iter < end; ++iter) - { - ASSERT_EQUAL(iter->getter(), T(42)); - ASSERT_EQUAL((*iter).getter(), T(42)); - ASSERT_EQUAL(iter.get()->getter(), T(42)); - ASSERT_EQUAL((*iter.get()).getter(), T(42)); - } -} -DECLARE_GENERIC_UNITTEST(TestIterationObj); - -template -void TestStdVector() -{ - // Verify that a std::vector using the universal allocator will work with - // STL algorithms. - std::vector > v0; - - static_assert( - std::is_same::type::pointer, - thrust::system::cuda::detail::managed_memory_pointer< - T > >::value, - "Unexpected pointer returned from unique_ptr::get."); - - v0.resize(6); - std::iota(v0.begin(), v0.end(), 0); - ASSERT_EQUAL(v0[0], T(0)); - ASSERT_EQUAL(v0[1], T(1)); - ASSERT_EQUAL(v0[2], T(2)); - ASSERT_EQUAL(v0[3], T(3)); - ASSERT_EQUAL(v0[4], T(4)); - ASSERT_EQUAL(v0[5], T(5)); + std::is_same< + thrust::universal_ptr, + thrust::system::cuda::detail::managed_memory_pointer + >::value, + "thrust::universal_ptr is not thrust::system::cuda::detail::managed_memory_pointer." + ); } -DECLARE_GENERIC_UNITTEST(TestStdVector); +DECLARE_GENERIC_UNITTEST(TestCudaManagedMemoryPointer); -#endif // C++11 diff --git a/testing/unittest/assertions.h b/testing/unittest/assertions.h index 6803e8168d..3528e09b9c 100644 --- a/testing/unittest/assertions.h +++ b/testing/unittest/assertions.h @@ -3,6 +3,7 @@ #include #include #include +#include #include #include @@ -376,7 +377,7 @@ class almost_equal_to > double a_tol, r_tol; almost_equal_to(double _a_tol = DEFAULT_ABSOLUTE_TOL, double _r_tol = DEFAULT_RELATIVE_TOL) : a_tol(_a_tol), r_tol(_r_tol) {} bool operator()(const thrust::complex& a, const thrust::complex& b) const { - return almost_equal((double) a.real(), (double) b.real(), a_tol, r_tol) + return almost_equal((double) a.real(), (double) b.real(), a_tol, r_tol) && almost_equal((double) a.imag(), (double) b.imag(), a_tol, r_tol); } }; @@ -390,12 +391,12 @@ void assert_equal(ForwardIterator1 first1, ForwardIterator1 last1, ForwardIterat { typedef typename thrust::iterator_difference::type difference_type; typedef typename thrust::iterator_value::type InputType; - + bool failure = false; difference_type length1 = thrust::distance(first1, last1); difference_type length2 = thrust::distance(first2, last2); - + difference_type min_length = thrust::min(length1, length2); unittest::UnitTestFailure f; @@ -409,7 +410,7 @@ void assert_equal(ForwardIterator1 first1, ForwardIterator1 last1, ForwardIterat } // check values - + size_t mismatches = 0; for (difference_type i = 0; i < min_length; i++) @@ -472,7 +473,6 @@ void assert_almost_equal(ForwardIterator1 first1, ForwardIterator1 last1, Forwar assert_equal(first1, last1, first2, last2, almost_equal_to(a_tol, r_tol), filename, lineno); } - template void assert_equal(const thrust::host_vector& A, const thrust::host_vector& B, const std::string& filename = "unknown", int lineno = -1) @@ -480,14 +480,6 @@ void assert_equal(const thrust::host_vector& A, const thrust::host_vec assert_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno); } -template -void assert_almost_equal(const thrust::host_vector& A, const thrust::host_vector& B, - const std::string& filename = "unknown", int lineno = -1, - const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) -{ - assert_almost_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno, a_tol, r_tol); -} - template void assert_equal(const thrust::host_vector& A, const thrust::device_vector& B, const std::string& filename = "unknown", int lineno = -1) @@ -513,6 +505,58 @@ void assert_equal(const thrust::device_vector& A, const thrust::device assert_equal(A_host, B_host, filename, lineno); } +template +void assert_equal(const thrust::universal_vector& A, const thrust::universal_vector& B, + const std::string& filename = "unknown", int lineno = -1) +{ + assert_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno); +} + +template +void assert_equal(const thrust::host_vector& A, const thrust::universal_vector& B, + const std::string& filename = "unknown", int lineno = -1) +{ + assert_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno); +} + +template +void assert_equal(const thrust::universal_vector& A, const thrust::host_vector& B, + const std::string& filename = "unknown", int lineno = -1) +{ + assert_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno); +} + +template +void assert_equal(const thrust::device_vector& A, const thrust::universal_vector& B, + const std::string& filename = "unknown", int lineno = -1) +{ + thrust::host_vector A_host = A; + assert_equal(A_host, B, filename, lineno); +} + +template +void assert_equal(const thrust::universal_vector& A, const thrust::device_vector& B, + const std::string& filename = "unknown", int lineno = -1) +{ + thrust::host_vector B_host = B; + assert_equal(A, B_host, filename, lineno); +} + +template +void assert_equal(const std::vector& A, const std::vector& B, + const std::string& filename = "unknown", int lineno = -1) +{ + assert_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno); +} + +template +void assert_almost_equal(const thrust::host_vector& A, const thrust::host_vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + assert_almost_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno, a_tol, r_tol); +} + template void assert_almost_equal(const thrust::host_vector& A, const thrust::device_vector& B, const std::string& filename = "unknown", int lineno = -1, @@ -541,6 +585,56 @@ void assert_almost_equal(const thrust::device_vector& A, const thrust: assert_almost_equal(A_host, B_host, filename, lineno, a_tol, r_tol); } +template +void assert_almost_equal(const thrust::universal_vector& A, const thrust::universal_vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + assert_almost_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno, a_tol, r_tol); +} + +template +void assert_almost_equal(const thrust::host_vector& A, const thrust::universal_vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + assert_almost_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno, a_tol, r_tol); +} + +template +void assert_almost_equal(const thrust::universal_vector& A, const thrust::host_vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + assert_almost_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno, a_tol, r_tol); +} + +template +void assert_almost_equal(const thrust::device_vector& A, const thrust::universal_vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + thrust::host_vector A_host = A; + assert_almost_equal(A_host, B, filename, lineno, a_tol, r_tol); +} + +template +void assert_almost_equal(const thrust::universal_vector& A, const thrust::device_vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + thrust::host_vector B_host = B; + assert_almost_equal(A, B_host, filename, lineno, a_tol, r_tol); +} + +template +void assert_almost_equal(const std::vector& A, const std::vector& B, + const std::string& filename = "unknown", int lineno = -1, + const double a_tol = DEFAULT_ABSOLUTE_TOL, const double r_tol = DEFAULT_RELATIVE_TOL) +{ + assert_almost_equal(A.begin(), A.end(), B.begin(), B.end(), filename, lineno, a_tol, r_tol); +} + enum threw_status { did_not_throw diff --git a/testing/unittest/testframework.h b/testing/unittest/testframework.h index ec5c42bb65..396306fa78 100644 --- a/testing/unittest/testframework.h +++ b/testing/unittest/testframework.h @@ -359,7 +359,7 @@ class NAME##UnitTest : public UnitTest { \ public: \ NAME##UnitTest() : UnitTest(#NAME) {} \ void run(){ \ - TEST(); \ + TEST(); \ } \ }; \ NAME##UnitTest NAME##Instance @@ -388,15 +388,16 @@ void VTEST##Device(void) { \ VTEST< thrust::device_vector > >(); \ - VTEST< thrust::device_vector > >(); \ +} \ +void VTEST##Universal(void) { \ + VTEST< thrust::universal_vector >(); \ VTEST< thrust::device_vector > >();\ } \ DECLARE_UNITTEST(VTEST##Host); \ -DECLARE_UNITTEST(VTEST##Device); +DECLARE_UNITTEST(VTEST##Device); \ +DECLARE_UNITTEST(VTEST##Universal); // Same as above, but only for integral types #define DECLARE_INTEGRAL_VECTOR_UNITTEST(VTEST) \ @@ -410,8 +411,15 @@ void VTEST##Device(void) { \ VTEST< thrust::device_vector >(); \ VTEST< thrust::device_vector >(); \ } \ +void VTEST##Universal(void) { \ + VTEST< thrust::universal_vector >(); \ + VTEST< thrust::device_vector > >();\ +} \ DECLARE_UNITTEST(VTEST##Host); \ -DECLARE_UNITTEST(VTEST##Device); +DECLARE_UNITTEST(VTEST##Device); \ +DECLARE_UNITTEST(VTEST##Universal); // Macro to create instances of a test for several data types. #define DECLARE_GENERIC_UNITTEST(TEST) \ diff --git a/testing/universal_memory.cu b/testing/universal_memory.cu new file mode 100644 index 0000000000..18a30fbfe6 --- /dev/null +++ b/testing/universal_memory.cu @@ -0,0 +1,166 @@ +#include + +#include +#include +#include +#include + +#include +#include + +namespace +{ + +// The managed_memory_pointer class should be identified as a +// contiguous_iterator +THRUST_STATIC_ASSERT( + thrust::is_contiguous_iterator::pointer>::value); + +template +struct some_object { + some_object(T data) + : m_data(data) + {} + + void setter(T data) { m_data = data; } + T getter() const { return m_data; } + +private: + T m_data; +}; + +} // namespace + +template +void TestUniversalAllocateUnique() +{ + // Simple test to ensure that pointers created with universal_memory_resource + // can be dereferenced and used with STL code. This is necessary as some + // STL implementations break when using fancy references that overload + // operator&, so universal_memory_resource uses a special pointer type that + // returns regular C++ references that can be safely used host-side. + + // These operations fail to compile with fancy references: + auto raw = thrust::allocate_unique(thrust::universal_allocator{}, 42); + auto obj = thrust::allocate_unique>( + thrust::universal_allocator >{}, 42 + ); + + static_assert( + std::is_same >::value, + "Unexpected pointer type returned from std::unique_ptr::get."); + static_assert( + std::is_same > >::value, + "Unexpected pointer type returned from std::unique_ptr::get."); + + ASSERT_EQUAL(*raw, T(42)); + ASSERT_EQUAL(*raw.get(), T(42)); + ASSERT_EQUAL(obj->getter(), T(42)); + ASSERT_EQUAL((*obj).getter(), T(42)); + ASSERT_EQUAL(obj.get()->getter(), T(42)); + ASSERT_EQUAL((*obj.get()).getter(), T(42)); +} +DECLARE_GENERIC_UNITTEST(TestUniversalAllocateUnique); + +template +void TestUniversalIterationRaw() +{ + auto array = thrust::allocate_unique_n( + thrust::universal_allocator{}, 6, 42); + + static_assert( + std::is_same >::value, + "Unexpected pointer type returned from std::unique_ptr::get."); + + for (auto iter = array.get(), end = array.get() + 6; iter < end; ++iter) + { + ASSERT_EQUAL(*iter, T(42)); + ASSERT_EQUAL(*iter.get(), T(42)); + } +} +DECLARE_GENERIC_UNITTEST(TestUniversalIterationRaw); + +template +void TestUniversalIterationObj() +{ + auto array = thrust::allocate_unique_n>( + thrust::universal_allocator>{}, 6, 42); + + static_assert( + std::is_same>>::value, + "Unexpected pointer type returned from std::unique_ptr::get."); + + for (auto iter = array.get(), end = array.get() + 6; iter < end; ++iter) + { + ASSERT_EQUAL(iter->getter(), T(42)); + ASSERT_EQUAL((*iter).getter(), T(42)); + ASSERT_EQUAL(iter.get()->getter(), T(42)); + ASSERT_EQUAL((*iter.get()).getter(), T(42)); + } +} +DECLARE_GENERIC_UNITTEST(TestUniversalIterationObj); + +template +void TestUniversalRawPointerCast() +{ + auto obj = thrust::allocate_unique(thrust::universal_allocator{}, 42); + + static_assert( + std::is_same>::value, + "Unexpected pointer type returned from std::unique_ptr::get."); + + static_assert( + std::is_same::value, + "Unexpected pointer type returned from thrust::raw_pointer_cast."); + + *thrust::raw_pointer_cast(obj.get()) = T(17); + + ASSERT_EQUAL(*obj, T(17)); +} +DECLARE_GENERIC_UNITTEST(TestUniversalRawPointerCast); + +template +void TestUniversalThrustVector(std::size_t const n) +{ + thrust::host_vector host(n); + thrust::universal_vector universal(n); + + static_assert( + std::is_same::type::pointer, + thrust::universal_ptr>::value, + "Unexpected thrust::universal_vector pointer type."); + + thrust::sequence(host.begin(), host.end(), 0); + thrust::sequence(universal.begin(), universal.end(), 0); + + ASSERT_EQUAL(host.size(), n); + ASSERT_EQUAL(universal.size(), n); + ASSERT_EQUAL(host, universal); +} +DECLARE_VARIABLE_UNITTEST(TestUniversalThrustVector); + +// Verify that a std::vector using the universal allocator will work with +// Standard Library algorithms. +template +void TestUniversalStdVector(std::size_t const n) +{ + std::vector host(n); + std::vector> universal(n); + + static_assert( + std::is_same::type::pointer, + thrust::universal_ptr>::value, + "Unexpected std::vector pointer type."); + + std::iota(host.begin(), host.end(), 0); + std::iota(universal.begin(), universal.end(), 0); + + ASSERT_EQUAL(host.size(), n); + ASSERT_EQUAL(universal.size(), n); + ASSERT_EQUAL(host, universal); +} +DECLARE_VARIABLE_UNITTEST(TestUniversalStdVector); + diff --git a/thrust/detail/pointer.h b/thrust/detail/pointer.h index e9204978f5..f9b0344018 100644 --- a/thrust/detail/pointer.h +++ b/thrust/detail/pointer.h @@ -94,7 +94,7 @@ template thrust::detail::identity_ >, thrust::detail::identity_ > - >::type reference_arg; + >::type reference_type; typedef thrust::iterator_adaptor< derived_type, // pass along the type of our Derived class to iterator_adaptor @@ -102,7 +102,7 @@ template value_type, // the value type Tag, // system tag thrust::random_access_traversal_tag, // pointers have random access traversal - reference_arg, // pass along our Reference type + reference_type, // pass along our Reference type std::ptrdiff_t > type; }; // end pointer_base @@ -205,6 +205,9 @@ template __host__ __device__ Element *get() const; + __host__ __device__ + Element *operator->() const; + #if THRUST_CPP_DIALECT >= 2011 // NOTE: This is needed so that Thrust smart pointers can be used in // `std::unique_ptr`. diff --git a/thrust/detail/pointer.inl b/thrust/detail/pointer.inl index 464c3579ed..d1a814ed41 100644 --- a/thrust/detail/pointer.inl +++ b/thrust/detail/pointer.inl @@ -159,6 +159,15 @@ template } // end pointer::get +template + __host__ __device__ + Element *pointer + ::operator->() const +{ + return super_t::base(); +} // end pointer::operator-> + + #if THRUST_CPP_DIALECT >= 2011 template __host__ __device__ @@ -211,63 +220,37 @@ bool operator!=(pointer p, decltype(nullptr)) } #endif -namespace detail -{ +template +struct tagged_reference; -#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (_MSC_VER <= 1400) -// XXX WAR MSVC 2005 problem with correctly implementing -// pointer_raw_pointer for pointer by specializing it here -template - struct pointer_raw_pointer< thrust::pointer > +namespace detail { - typedef typename pointer::raw_pointer type; -}; // end pointer_raw_pointer -#endif - -#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) && (THRUST_GCC_VERSION < 40200) -// XXX WAR g++-4.1 problem with correctly implementing -// pointer_element for pointer by specializing it here -template - struct pointer_element< thrust::pointer > +template +struct rebind_pointer< + thrust::pointer>, + NewElement +> { - typedef Element type; -}; // end pointer_element - -template - struct pointer_element< thrust::pointer > - : pointer_element< thrust::pointer > -{}; // end pointer_element - -template - struct pointer_element< thrust::pointer > - : pointer_element< thrust::pointer > -{}; // end pointer_element - - + using type = thrust::pointer< + NewElement, Tag, thrust::tagged_reference + >; +}; -// XXX WAR g++-4.1 problem with correctly implementing -// rebind_pointer for pointer by specializing it here -template - struct rebind_pointer, NewElement> +template +struct rebind_pointer< + thrust::pointer< + Element, Tag, typename std::add_lvalue_reference::type + >, + NewElement +> { - // XXX note we don't attempt to rebind the pointer's Reference type (or Derived) - typedef thrust::pointer type; + using type = thrust::pointer< + NewElement, Tag, typename std::add_lvalue_reference::type + >; }; -template - struct rebind_pointer, NewElement> - : rebind_pointer, NewElement> -{}; - -template - struct rebind_pointer, NewElement> - : rebind_pointer, NewElement> -{}; -#endif - -} // end namespace detail - +} // namespace detail -} // end thrust +} // namespace thrust diff --git a/thrust/detail/reference.h b/thrust/detail/reference.h index 89bcf63ca7..ee218f6ba3 100644 --- a/thrust/detail/reference.h +++ b/thrust/detail/reference.h @@ -172,7 +172,101 @@ std::basic_ostream & operator<<(std::basic_ostream &os, const reference &y); -} // end thrust +template +class tagged_reference + : public thrust::reference< + Element, + thrust::pointer>, + tagged_reference + > +{ +private: + typedef thrust::reference< + Element, + thrust::pointer>, + tagged_reference + > super_t; + +public: + using value_type = typename super_t::value_type; + using pointer = typename super_t::pointer; + + /*! This constructor initializes this \p tagged_reference to refer to an + * object pointed to by the given \p pointer. After this + * \p tagged_reference is constructed, it shall refer to the object pointed + * to by \p ptr. + * + * \param ptr A \p pointer to copy from. + */ + __host__ __device__ explicit tagged_reference(const pointer &ptr) + : super_t(ptr) + { + } + + /*! This constructor accepts a const reference to another + * \p tagged_reference of related type. After this \p tagged_reference is + * constructed, it shall refer to the same object as \p other. + * + * \param other A \p tagged_reference to copy from. + * \tparam OtherT The element type of the other \p tagged_reference. + * + * \note This constructor is templated primarily to allow initialization + * of tagged_reference from + * tagged_reference. + */ + template + __host__ __device__ + tagged_reference(const tagged_reference &other, + typename thrust::detail::enable_if_convertible< + typename tagged_reference< + OtherElement, OtherTag + >::pointer, + pointer + >::type * = 0) + : super_t(other) + { + } + + /*! Copy assignment operator copy assigns from another \p tagged_reference + * of related type. + * + * \param other The other \p tagged_reference to assign from. + * \return *this + * \tparam OtherT The element type of the other \p tagged_reference. + */ + template + __host__ __device__ + tagged_reference & + operator=(const tagged_reference &other) + { + return super_t::operator=(other); + } + + /*! Assignment operator assigns from a \p value_type. + * + * \param x The \p value_type to assign from. + * \return *this + */ + __host__ __device__ + tagged_reference & + operator=(const value_type &x) + { + return super_t::operator=(x); + } +}; + +/*! Exchanges the values of two objects referred to by \p tagged_reference. + * \p x The first \p tagged_reference of interest. + * \p y The second \p tagged_reference of interest. + */ +template +__host__ __device__ +void swap(tagged_reference x, tagged_reference y) +{ + x.swap(y); +} + +} // namespace thrust #include diff --git a/thrust/device_allocator.h b/thrust/device_allocator.h index f5ff0d9654..6ae0cafab3 100644 --- a/thrust/device_allocator.h +++ b/thrust/device_allocator.h @@ -16,7 +16,8 @@ /*! \file device_allocator.h - * \brief An allocator which creates new elements in device memory + * \brief An allocator which creates new elements in memory accessible by + * devices. */ #pragma once @@ -83,13 +84,10 @@ class device_ptr_memory_resource THRUST_FINAL Upstream * m_upstream; }; -/*! \} - */ - -/*! \addtogroup memory_management Memory Management - * \addtogroup memory_management_classes Memory Management Classes - * \ingroup memory_management - * \{ +/*! \brief An allocator which creates new elements in memory accessible by + * devices. + * + * \see https://en.cppreference.com/w/cpp/named_req/Allocator */ template class device_allocator diff --git a/thrust/device_ptr.h b/thrust/device_ptr.h index fb3ad1ee02..f9149da148 100644 --- a/thrust/device_ptr.h +++ b/thrust/device_ptr.h @@ -16,7 +16,7 @@ /*! \file device_ptr.h - * \brief A pointer to a variable which resides in the "device" system's memory space + * \brief A pointer to a variable which resides memory accessible to devices. */ #pragma once @@ -89,7 +89,7 @@ template /*! \p device_ptr's copy constructor is templated to allow copying to a * device_ptr from a T *. - * + * * \param ptr A raw pointer to copy from, presumed to point to a location in * device memory. */ diff --git a/thrust/device_vector.h b/thrust/device_vector.h index fa52ec6623..892ecd0195 100644 --- a/thrust/device_vector.h +++ b/thrust/device_vector.h @@ -16,7 +16,8 @@ /*! \file device_vector.h - * \brief A dynamically-sizable array of elements which reside in the "device" memory space + * \brief A dynamically-sizable array of elements which resides in memory + * accessible to devices. */ #pragma once @@ -44,12 +45,13 @@ template class host_vector; * constant time removal of elements at the end, and linear time insertion * and removal of elements at the beginning or in the middle. The number of * elements in a \p device_vector may vary dynamically; memory management is - * automatic. The memory associated with a \p device_vector resides in the memory - * space of a parallel device. + * automatic. The memory associated with a \p device_vector resides in the + * memory accessible to devices. * - * \see http://www.sgi.com/tech/stl/Vector.html + * \see https://en.cppreference.com/w/cpp/container/vector * \see device_allocator * \see host_vector + * \see universal_vector */ template > class device_vector @@ -431,7 +433,7 @@ template > * \param x The exemplar element to copy & insert. * \return An iterator pointing to the newly inserted element. */ - iterator insert(iterator position, const T &x); + iterator insert(iterator position, const T &x); /*! This method inserts a copy of an exemplar value to a range at the * specified position in this vector. diff --git a/thrust/host_vector.h b/thrust/host_vector.h index ebe64216e2..c6407a96b3 100644 --- a/thrust/host_vector.h +++ b/thrust/host_vector.h @@ -16,7 +16,8 @@ /*! \file host_vector.h - * \brief A dynamically-sizable array of elements which reside in the "host" memory space + * \brief A dynamically-sizable array of elements which resides in memory + * accessible to hosts. */ #pragma once @@ -43,11 +44,12 @@ template class device_vector; * constant time removal of elements at the end, and linear time insertion * and removal of elements at the beginning or in the middle. The number of * elements in a \p host_vector may vary dynamically; memory management is - * automatic. The memory associated with a \p host_vector resides in the memory - * space of the host associated with a parallel device. + * automatic. The memory associated with a \p host_vector resides in memory + * accessible to hosts. * - * \see http://www.sgi.com/tech/stl/Vector.html + * \see https://en.cppreference.com/w/cpp/container/vector * \see device_vector + * \see universal_vector */ template > class host_vector @@ -450,7 +452,7 @@ template > * \param x The exemplar element to copy & insert. * \return An iterator pointing to the newly inserted element. */ - iterator insert(iterator position, const T &x); + iterator insert(iterator position, const T &x); /*! This method inserts a copy of an exemplar value to a range at the * specified position in this vector. diff --git a/thrust/system/cpp/detail/pointer.inl b/thrust/system/cpp/detail/pointer.inl deleted file mode 100644 index 7d9de3e55a..0000000000 --- a/thrust/system/cpp/detail/pointer.inl +++ /dev/null @@ -1,67 +0,0 @@ -/* - * Copyright 2008-2018 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -namespace thrust -{ - -// XXX WAR an issue with MSVC 2005 (cl v14.00) incorrectly implementing -// pointer_raw_pointer for pointer by specializing it here -#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (_MSC_VER <= 1400) -namespace detail -{ - -template - struct pointer_raw_pointer< thrust::cpp::pointer > -{ - typedef typename thrust::cpp::pointer::raw_pointer type; -}; // end pointer_raw_pointer - -} // end detail -#endif - -namespace system -{ -namespace cpp -{ - -template - template - reference & - reference - ::operator=(const reference &other) -{ - return super_t::operator=(other); -} // end reference::operator=() - -template - reference & - reference - ::operator=(const value_type &x) -{ - return super_t::operator=(x); -} // end reference::operator=() - -template -__host__ __device__ -void swap(reference a, reference b) -{ - a.swap(b); -} // end swap() - -} // end cpp -} // end system -} // end thrust - diff --git a/thrust/system/cpp/execution_policy.h b/thrust/system/cpp/execution_policy.h index 3bf521be34..d22b4ceeb7 100644 --- a/thrust/system/cpp/execution_policy.h +++ b/thrust/system/cpp/execution_policy.h @@ -14,12 +14,12 @@ * limitations under the License. */ -#pragma once - /*! \file thrust/system/cpp/execution_policy.h - * \brief Execution policies for Thrust's standard C++ system. + * \brief Execution policies for Thrust's Standard C++ system. */ +#pragma once + #include // get the execution policies definitions first @@ -104,7 +104,7 @@ struct execution_policy : thrust::execution_policy struct tag : thrust::system::cpp::execution_policy { unspecified }; -/*! +/*! * \p thrust::system::cpp::par is the parallel execution policy associated with Thrust's standard * C++ backend system. * diff --git a/thrust/system/cpp/memory.h b/thrust/system/cpp/memory.h index 18b31e758d..1f85f90fcf 100644 --- a/thrust/system/cpp/memory.h +++ b/thrust/system/cpp/memory.h @@ -15,7 +15,7 @@ */ /*! \file thrust/system/cpp/memory.h - * \brief Managing memory associated with Thrust's standard C++ system. + * \brief Managing memory associated with Thrust's Standard C++ system. */ #pragma once @@ -27,12 +27,9 @@ #include #include -namespace thrust -{ -namespace system -{ -namespace cpp +namespace thrust { namespace system { namespace cpp { + /*! Allocates an area of memory available to Thrust's cpp system. * \param n Number of bytes to allocate. * \return A cpp::pointer pointing to the beginning of the newly @@ -66,30 +63,37 @@ inline pointer malloc(std::size_t n); */ inline void free(pointer ptr); -/*! \p cpp::allocator is the default allocator used by the \p cpp system's containers such as - * cpp::vector if no user-specified allocator is provided. \p cpp::allocator allocates - * (deallocates) storage with \p cpp::malloc (\p cpp::free). +/*! \p cpp::allocator is the default allocator used by the \p cpp system's + * containers such as cpp::vector if no user-specified allocator is + * provided. \p cpp::allocator allocates (deallocates) storage with \p + * cpp::malloc (\p cpp::free). */ template -using allocator = thrust::mr::stateless_resource_allocator; +using allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::cpp::memory_resource +>; -} // end cpp +/*! \p cpp::universal_allocator allocates memory that can be used by the \p cpp + * system and host systems. + */ +template +using universal_allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::cpp::universal_memory_resource +>; -} // end system +} // namespace system::cpp /*! \namespace thrust::cpp * \brief \p thrust::cpp is a top-level alias for thrust::system::cpp. */ namespace cpp { - using thrust::system::cpp::malloc; using thrust::system::cpp::free; using thrust::system::cpp::allocator; +} // namespace cpp -} // end cpp - -} // end thrust +} // namespace thrust #include diff --git a/thrust/system/cpp/pointer.h b/thrust/system/cpp/pointer.h index 8efeb33c46..dac60a7e3c 100644 --- a/thrust/system/cpp/pointer.h +++ b/thrust/system/cpp/pointer.h @@ -1,5 +1,5 @@ /* - * Copyright 2008-2018 NVIDIA Corporation + * Copyright 2008-2020 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,116 +14,36 @@ * limitations under the License. */ +/*! \file thrust/system/cpp/memory.h + * \brief Managing memory associated with Thrust's TBB system. + */ + #pragma once #include +#include #include -#include #include #include -namespace thrust -{ -namespace system -{ -namespace cpp -{ - -template class pointer; - -} // end cpp -} // end system -} // end thrust - - -/*! \cond - */ - -// specialize thrust::iterator_traits to avoid problems with the name of -// pointer's constructor shadowing its nested pointer type -// do this before pointer is defined so the specialization is correctly -// used inside the definition -namespace thrust -{ - -template - struct iterator_traits > -{ - private: - typedef thrust::system::cpp::pointer ptr; - - public: - typedef typename ptr::iterator_category iterator_category; - typedef typename ptr::value_type value_type; - typedef typename ptr::difference_type difference_type; - typedef ptr pointer; - typedef typename ptr::reference reference; -}; // end iterator_traits - -} // end thrust - -/*! \endcond - */ - - -namespace thrust +namespace thrust { namespace system { namespace cpp { -namespace system -{ - -/*! \addtogroup system_backends Systems - * \ingroup system - * \{ - */ - -/*! \namespace thrust::system::cpp - * \brief \p thrust::system::cpp is the namespace containing functionality for allocating, manipulating, - * and deallocating memory available to Thrust's standard C++ backend system. - * The identifiers are provided in a separate namespace underneath thrust::system - * for import convenience but are also aliased in the top-level thrust::cpp - * namespace for easy access. - * - */ -namespace cpp -{ - -// forward declaration of reference for pointer -template class reference; - -/*! \cond - */ - -// XXX nvcc + msvc have trouble instantiating reference below -// this is a workaround -namespace detail -{ - -template - struct reference_msvc_workaround -{ - typedef thrust::system::cpp::reference type; -}; // end reference_msvc_workaround - -} // end detail - -/*! \endcond - */ - -/*! \p pointer stores a pointer to an object allocated in memory available to the cpp system. - * This type provides type safety when dispatching standard algorithms on ranges resident - * in cpp memory. +/*! \p cpp::pointer stores a pointer to an object allocated in memory accessible + * by the \p cpp system. This type provides type safety when dispatching + * algorithms on ranges resident in \p cpp memory. * - * \p pointer has pointer semantics: it may be dereferenced and manipulated with pointer arithmetic. + * \p cpp::pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. * - * \p pointer can be created with the function \p cpp::malloc, or by explicitly calling its constructor - * with a raw pointer. + * \p cpp::pointer can be created with the function \p cpp::malloc, or by + * explicitly calling its constructor with a raw pointer. * - * The raw pointer encapsulated by a \p pointer may be obtained by eiter its get member function - * or the \p raw_pointer_cast function. + * The raw pointer encapsulated by a \p cpp::pointer may be obtained by eiter its + * get member function or the \p raw_pointer_cast function. * - * \note \p pointer is not a "smart" pointer; it is the programmer's responsibility to deallocate memory - * pointed to by \p pointer. + * \note \p cpp::pointer is not a "smart" pointer; it is the programmer's + * responsibility to deallocate memory pointed to by \p cpp::pointer. * * \tparam T specifies the type of the pointee. * @@ -131,221 +51,66 @@ template * \see cpp::free * \see raw_pointer_cast */ -template - class pointer - : public thrust::pointer< - T, - thrust::system::cpp::tag, - thrust::system::cpp::reference, - thrust::system::cpp::pointer - > -{ - /*! \cond - */ - - private: - typedef thrust::pointer< - T, - thrust::system::cpp::tag, - //thrust::system::cpp::reference, - typename detail::reference_msvc_workaround::type, - thrust::system::cpp::pointer - > super_t; - - /*! \endcond - */ - - public: - // note that cpp::pointer's member functions need __host__ __device__ - // to interoperate with nvcc + iterators' dereference member function - - /*! \p pointer's no-argument constructor initializes its encapsulated pointer to \c 0. - */ - __host__ __device__ - pointer() : super_t() {} - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer(decltype(nullptr)) : super_t(nullptr) {} - #endif - - /*! This constructor allows construction of a pointer from a T*. - * - * \param ptr A raw pointer to copy from, presumed to point to a location in memory - * accessible by the \p cpp system. - * \tparam OtherT \p OtherT shall be convertible to \p T. - */ - template - __host__ __device__ - explicit pointer(OtherT *ptr) : super_t(ptr) {} - - /*! This constructor allows construction from another pointer-like object with related type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::cpp::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer - >::type * = 0) : super_t(other) {} - - /*! This constructor allows construction from another pointer-like object with \p void type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::cpp::tag and its element type shall be \p void. - */ - template - __host__ __device__ - explicit - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_void_pointer_is_system_convertible< - OtherPointer, - pointer - >::type * = 0) : super_t(other) {} - - /*! Assignment operator allows assigning from another pointer-like object with related type. - * - * \param other The other pointer-like object to assign from. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::cpp::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer, - pointer & - >::type - operator=(const OtherPointer &other) - { - return super_t::operator=(other); - } - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer& operator=(decltype(nullptr)) - { - super_t::operator=(nullptr); - return *this; - } - #endif -}; // end pointer - -/*! \p reference is a wrapped reference to an object stored in memory available to the \p cpp system. - * \p reference is the type of the result of dereferencing a \p cpp::pointer. +template +using pointer = thrust::pointer< + T, + thrust::system::cpp::tag, + thrust::tagged_reference +>; + +/*! \p cpp::universal_pointer stores a pointer to an object allocated in memory + * accessible by the \p cpp system and host systems. * - * \tparam T Specifies the type of the referenced object. + * \p cpp::universal_pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. + * + * \p cpp::universal_pointer can be created with \p cpp::universal_allocator + * or by explicitly calling its constructor with a raw pointer. + * + * The raw pointer encapsulated by a \p cpp::universal_pointer may be obtained + * by eiter its get member function or the \p raw_pointer_cast + * function. + * + * \note \p cpp::universal_pointer is not a "smart" pointer; it is the + * programmer's responsibility to deallocate memory pointed to by + * \p cpp::universal_pointer. + * + * \tparam T specifies the type of the pointee. + * + * \see cpp::universal_allocator + * \see raw_pointer_cast */ -template - class reference - : public thrust::reference< - T, - thrust::system::cpp::pointer, - thrust::system::cpp::reference - > -{ - /*! \cond - */ - - private: - typedef thrust::reference< - T, - thrust::system::cpp::pointer, - thrust::system::cpp::reference - > super_t; - - /*! \endcond - */ - - public: - /*! \cond - */ - - typedef typename super_t::value_type value_type; - typedef typename super_t::pointer pointer; - - /*! \endcond - */ - - /*! This constructor initializes this \p reference to refer to an object - * pointed to by the given \p pointer. After this \p reference is constructed, - * it shall refer to the object pointed to by \p ptr. - * - * \param ptr A \p pointer to copy from. - */ - __host__ __device__ - explicit reference(const pointer &ptr) - : super_t(ptr) - {} - - /*! This constructor accepts a const reference to another \p reference of related type. - * After this \p reference is constructed, it shall refer to the same object as \p other. - * - * \param other A \p reference to copy from. - * \tparam OtherT The element type of the other \p reference. - * - * \note This constructor is templated primarily to allow initialization of reference - * from reference. - */ - template - __host__ __device__ - reference(const reference &other, - typename thrust::detail::enable_if_convertible< - typename reference::pointer, - pointer - >::type * = 0) - : super_t(other) - {} - - /*! Copy assignment operator copy assigns from another \p reference of related type. - * - * \param other The other \p reference to assign from. - * \return *this - * \tparam OtherT The element type of the other \p reference. - */ - template - reference &operator=(const reference &other); - - /*! Assignment operator assigns from a \p value_type. - * - * \param x The \p value_type to assign from. - * \return *this - */ - reference &operator=(const value_type &x); -}; // end reference - -/*! Exchanges the values of two objects referred to by \p reference. - * \p x The first \p reference of interest. - * \p y The second \p reference of interest. +template +using universal_pointer = thrust::pointer< + T, + thrust::system::cpp::tag, + typename std::add_lvalue_reference::type +>; + +/*! \p reference is a wrapped reference to an object stored in memory available + * to the \p cpp system. \p reference is the type of the result of + * dereferencing a \p cpp::pointer. + * + * \tparam T Specifies the type of the referenced object. */ -template -__host__ __device__ -void swap(reference x, reference y); +template +using reference = thrust::reference; -} // end cpp +}} // namespace system::cpp -/*! \} +/*! \addtogroup system_backends Systems + * \ingroup system + * \{ */ -} // end system - +/*! \namespace thrust::cpp + * \brief \p thrust::cpp is a top-level alias for \p thrust::system::cpp. */ namespace cpp { - using thrust::system::cpp::pointer; +using thrust::system::cpp::universal_pointer; using thrust::system::cpp::reference; +} // namespace cpp -} // end cpp - -} // end thrust +} // namespace thrust -#include diff --git a/thrust/system/cpp/vector.h b/thrust/system/cpp/vector.h index ee5cfce6aa..0d328f134a 100644 --- a/thrust/system/cpp/vector.h +++ b/thrust/system/cpp/vector.h @@ -26,15 +26,7 @@ #include #include -namespace thrust -{ - -// forward declaration of host_vector -template class host_vector; - -namespace system -{ -namespace cpp +namespace thrust { namespace system { namespace cpp { /*! \p cpp::vector is a container that supports random access to elements, @@ -42,28 +34,48 @@ namespace cpp * and removal of elements at the beginning or in the middle. The number of * elements in a \p cpp::vector may vary dynamically; memory management is * automatic. The elements contained in a \p cpp::vector reside in memory - * available to the \p cpp system. + * accessible by the \p cpp system. * * \tparam T The element type of the \p cpp::vector. - * \tparam Allocator The allocator type of the \p cpp::vector. Defaults to \p cpp::allocator. + * \tparam Allocator The allocator type of the \p cpp::vector. + * Defaults to \p cpp::allocator. * - * \see http://www.sgi.com/tech/stl/Vector.html + * \see https://en.cppreference.com/w/cpp/container/vector * \see host_vector For the documentation of the complete interface which is - * shared by \p cpp::vector + * shared by \p cpp::vector. * \see device_vector + * \see universal_vector */ -template > +template > using vector = thrust::detail::vector_base; -} // end cpp -} // end system +/*! \p cpp::universal_vector is a container that supports random access to + * elements, constant time removal of elements at the end, and linear time + * insertion and removal of elements at the beginning or in the middle. The + * number of elements in a \p cpp::universal_vector may vary dynamically; + * memory management is automatic. The elements contained in a + * \p cpp::universal_vector reside in memory accessible by the \p cpp system + * and host systems. + * + * \tparam T The element type of the \p cpp::universal_vector. + * \tparam Allocator The allocator type of the \p cpp::universal_vector. + * Defaults to \p cpp::universal_allocator. + * + * \see https://en.cppreference.com/w/cpp/container/vector + * \see host_vector For the documentation of the complete interface which is + * shared by \p cpp::universal_vector + * \see device_vector + * \see universal_vector + */ +template > +using universal_vector = thrust::detail::vector_base; + +}} // namespace system::cpp -// alias system::cpp names at top-level namespace cpp { - using thrust::system::cpp::vector; - -} // end cpp +using thrust::system::cpp::universal_vector; +} } // end thrust diff --git a/thrust/system/cuda/detail/managed_memory_pointer.h b/thrust/system/cuda/detail/managed_memory_pointer.h deleted file mode 100644 index c6a4c9756b..0000000000 --- a/thrust/system/cuda/detail/managed_memory_pointer.h +++ /dev/null @@ -1,195 +0,0 @@ -/* - * Copyright 2020 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include - -#include -#include - -namespace thrust -{ -namespace system -{ -namespace cuda -{ -namespace detail -{ - -// forward decl for iterator traits: -template -class managed_memory_pointer; - -} // end namespace detail -} // end namespace cuda -} // end namespace system - -// Specialize iterator traits to define `pointer` to something meaningful. -template -struct iterator_traits > > { -private: - typedef thrust::pointer< - Element, - Tag, - Reference, - thrust::system::cuda::detail::managed_memory_pointer > - ptr; - -public: - typedef typename ptr::iterator_category iterator_category; - typedef typename ptr::value_type value_type; - typedef typename ptr::difference_type difference_type; - typedef Element* pointer; - typedef typename ptr::reference reference; -}; // end iterator_traits - -namespace system -{ -namespace cuda -{ -namespace detail -{ - -/*! A version of thrust::cuda_cub::pointer that uses c++ references instead - * of thrust::cuda::reference. This is to allow managed memory pointers to - * be used with host-side code in standard libraries that are not compatible - * with proxy references. - */ -template -class managed_memory_pointer - : public thrust::pointer< - T, - thrust::cuda_cub::tag, - typename thrust::detail::add_reference::type, - thrust::system::cuda::detail::managed_memory_pointer > -{ -private: - typedef thrust::pointer< - T, - thrust::cuda_cub::tag, - typename thrust::detail::add_reference::type, - thrust::system::cuda::detail::managed_memory_pointer > - super_t; - -public: - typedef typename super_t::raw_pointer pointer; - - /*! \p managed_memory_pointer's no-argument constructor initializes its - * encapsulated pointer to \c 0. - */ - __host__ __device__ managed_memory_pointer() - : super_t() - {} - -#if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ managed_memory_pointer(decltype(nullptr)) - : super_t(nullptr) - {} -#endif - - /*! This constructor allows construction of a from a - * T*. - * - * \param ptr A raw pointer to copy from, presumed to point to a location - * in memory accessible by the \p cuda system. \tparam OtherT \p OtherT - * shall be convertible to \p T. - */ - template - __host__ __device__ explicit managed_memory_pointer(OtherT* ptr) - : super_t(ptr) - {} - - /*! This constructor allows construction from another pointer-like object - * with related type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer - * shall be convertible to \p thrust::system::cuda::tag and its element - * type shall be convertible to \p T. - */ - template - __host__ __device__ managed_memory_pointer( - const OtherPointer& other, - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - managed_memory_pointer>::type* = 0) - : super_t(other) - {} - - /*! This constructor allows construction from another pointer-like object - * with \p void type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer - * shall be convertible to \p thrust::system::cuda::tag and its element - * type shall be \p void. - */ - template - __host__ __device__ explicit managed_memory_pointer( - const OtherPointer& other, - typename thrust::detail::enable_if_void_pointer_is_system_convertible< - OtherPointer, - managed_memory_pointer>::type* = 0) - : super_t(other) - {} - - /*! Assignment operator allows assigning from another pointer-like object - * with related type. - * - * \param other The other pointer-like object to assign from. - * \tparam OtherPointer The system tag associated with \p OtherPointer - * shall be convertible to \p thrust::system::cuda::tag and its element - * type shall be convertible to \p T. - */ - template - __host__ __device__ typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - managed_memory_pointer, - managed_memory_pointer&>::type - operator=(const OtherPointer& other) - { - return super_t::operator=(other); - } - -#if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ managed_memory_pointer& operator=(decltype(nullptr)) - { - super_t::operator=(nullptr); - return *this; - } -#endif - - __host__ __device__ - pointer operator->() const - { - return this->get(); - } - -}; // class managed_memory_pointer - -} // namespace detail -} // namespace cuda -} // namespace system -} // namespace thrust diff --git a/thrust/system/cuda/detail/pointer.inl b/thrust/system/cuda/detail/pointer.inl deleted file mode 100644 index 60f277f597..0000000000 --- a/thrust/system/cuda/detail/pointer.inl +++ /dev/null @@ -1,59 +0,0 @@ -/* - * Copyright 2008-2018 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -namespace thrust -{ - -// XXX WAR an issue with MSVC 2005 (cl v14.00) incorrectly implementing -// pointer_raw_pointer for pointer by specializing it here -// note that we specialize it here, before the use of raw_pointer_cast -// below, which causes pointer_raw_pointer's instantiation -#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (_MSC_VER <= 1400) -namespace detail -{ - -template - struct pointer_raw_pointer< thrust::cuda_cub::pointer > -{ - typedef typename thrust::cuda_cub::pointer::raw_pointer type; -}; // end pointer_raw_pointer - -} // end detail -#endif - -namespace cuda_cub { - -template -template -__host__ __device__ reference &reference::operator=( - const reference &other) { - return super_t::operator=(other); -} // end reference::operator=() - -template -__host__ __device__ reference &reference::operator=(const value_type &x) { - return super_t::operator=(x); -} // end reference::operator=() - -template -__host__ __device__ -void swap(reference a, reference b) -{ - a.swap(b); -} // end swap() - -} // end cuda_cub -} // end thrust diff --git a/thrust/system/cuda/memory.h b/thrust/system/cuda/memory.h index f20ce352a2..4d94a08850 100644 --- a/thrust/system/cuda/memory.h +++ b/thrust/system/cuda/memory.h @@ -2,7 +2,7 @@ * Copyright 2008-2018 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in ccudaliance with the License. + * you may not use this file except in compliance with the License. * You may obtain a copy of the License at * * http://www.apache.org/licenses/LICENSE-2.0 @@ -27,9 +27,8 @@ #include #include -namespace thrust +namespace thrust { namespace cuda_cub { -namespace cuda_cub { /*! Allocates an area of memory available to Thrust's cuda system. * \param n Number of bytes to allocate. @@ -64,30 +63,46 @@ inline __host__ __device__ pointer malloc(std::size_t n); */ inline __host__ __device__ void free(pointer ptr); -/*! \p cuda::allocator is the default allocator used by the \p cuda system's containers such as - * cuda::vector if no user-specified allocator is provided. \p cuda::allocator allocates - * (deallocates) storage with \p cuda::malloc (\p cuda::free). +/*! \p cuda::allocator is the default allocator used by the \p cuda system's + * containers such as cuda::vector if no user-specified allocator is + * provided. \p cuda::allocator allocates (deallocates) storage with \p + * cuda::malloc (\p cuda::free). */ template -using allocator = thrust::mr::stateless_resource_allocator; +using allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::cuda::memory_resource +>; -} // namespace cuda_cub +/*! \p cuda::universal_allocator allocates memory that can be used by the \p cuda + * system and host systems. + */ +template +using universal_allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::cuda::universal_memory_resource +>; -namespace system { -namespace cuda { +} // namespace cuda_cub + +namespace system { namespace cuda +{ using thrust::cuda_cub::malloc; using thrust::cuda_cub::free; using thrust::cuda_cub::allocator; -} // namespace cuda -} // namespace system +using thrust::cuda_cub::universal_allocator; +}} // namespace system::cuda -namespace cuda { +/*! \namespace thrust::cuda + * \brief \p thrust::cuda is a top-level alias for \p thrust::system::cuda. + */ +namespace cuda +{ using thrust::cuda_cub::malloc; using thrust::cuda_cub::free; using thrust::cuda_cub::allocator; -} // end cuda +using thrust::cuda_cub::universal_allocator; +} // namespace cuda -} // end namespace thrust +} // namespace thrust #include diff --git a/thrust/system/cuda/memory_resource.h b/thrust/system/cuda/memory_resource.h index 9110e0af45..54483a4863 100644 --- a/thrust/system/cuda/memory_resource.h +++ b/thrust/system/cuda/memory_resource.h @@ -22,7 +22,6 @@ #include #include -#include #include #include #include @@ -88,7 +87,7 @@ namespace detail thrust::cuda::pointer > device_memory_resource; typedef detail::cuda_memory_resource > + thrust::cuda::universal_pointer > managed_memory_resource; typedef detail::cuda_memory_resource @@ -107,5 +106,12 @@ typedef detail::pinned_memory_resource universal_host_pinned_memory_resource; } // end cuda } // end system +namespace cuda +{ +using thrust::system::cuda::memory_resource; +using thrust::system::cuda::universal_memory_resource; +using thrust::system::cuda::universal_host_pinned_memory_resource; +} + } // end namespace thrust diff --git a/thrust/system/cuda/pointer.h b/thrust/system/cuda/pointer.h index f198385ce2..c586eb9dc4 100644 --- a/thrust/system/cuda/pointer.h +++ b/thrust/system/cuda/pointer.h @@ -1,8 +1,8 @@ /* - * Copyright 2008-2018 NVIDIA Corporation + * Copyright 2008-2020 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in ccudaliance with the License. + * you may not use this file except in compliance with the License. * You may obtain a copy of the License at * * http://www.apache.org/licenses/LICENSE-2.0 @@ -14,76 +14,36 @@ * limitations under the License. */ +/*! \file thrust/system/cuda/memory.h + * \brief Managing memory associated with Thrust's Standard C++ system. + */ + #pragma once #include +#include #include -#include #include #include -namespace thrust -{ -namespace cuda_cub -{ - -template -class pointer; - -} // end cuda_cub -} // end thrust - - -// specialize thrust::iterator_traits to avoid problems with the name of -// pointer's constructor shadowing its nested pointer type -// do this before pointer is defined so the specialization is correctly -// used inside the definition -namespace thrust -{ - -template -struct iterator_traits > -{ -private: - typedef thrust::cuda_cub::pointer ptr; - -public: - typedef typename ptr::iterator_category iterator_category; - typedef typename ptr::value_type value_type; - typedef typename ptr::difference_type difference_type; - typedef ptr pointer; - typedef typename ptr::reference reference; -}; // end iterator_traits - -namespace cuda_cub { - -// forward declaration of reference for pointer -template -class reference; - -// XXX nvcc + msvc have trouble instantiating reference below -// this is a workaround -template -struct reference_msvc_workaround +namespace thrust { namespace cuda_cub { - typedef thrust::cuda_cub::reference type; -}; // end reference_msvc_workaround - -/*! \p pointer stores a pointer to an object allocated in memory available to the cuda system. - * This type provides type safety when dispatching standard algorithms on ranges resident - * in cuda memory. +/*! \p cuda::pointer stores a pointer to an object allocated in memory + * accessible by the \p cuda system. This type provides type safety when + * dispatching algorithms on ranges resident in \p cuda memory. * - * \p pointer has pointer semantics: it may be dereferenced and manipulated with pointer arithmetic. + * \p cuda::pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. * - * \p pointer can be created with the function \p cuda::malloc, or by explicitly calling its constructor - * with a raw pointer. + * \p cuda::pointer can be created with the function \p cuda::malloc, or by + * explicitly calling its constructor with a raw pointer. * - * The raw pointer encapsulated by a \p pointer may be obtained by eiter its get member function - * or the \p raw_pointer_cast function. + * The raw pointer encapsulated by a \p cuda::pointer may be obtained by eiter + * its get member function or the \p raw_pointer_cast function. * - * \note \p pointer is not a "smart" pointer; it is the programmer's responsibility to deallocate memory - * pointed to by \p pointer. + * \note \p cuda::pointer is not a "smart" pointer; it is the programmer's + * responsibility to deallocate memory pointed to by \p cuda::pointer. * * \tparam T specifies the type of the pointee. * @@ -92,198 +52,53 @@ struct reference_msvc_workaround * \see raw_pointer_cast */ template -class pointer - : public thrust::pointer< - T, - thrust::cuda_cub::tag, - thrust::cuda_cub::reference, - thrust::cuda_cub::pointer > -{ - -private: - typedef thrust::pointer< - T, - thrust::cuda_cub::tag, - typename reference_msvc_workaround::type, - thrust::cuda_cub::pointer > - super_t; - -public: - /*! \p pointer's no-argument constructor initializes its encapsulated pointer to \c 0. - */ - __host__ __device__ - pointer() : super_t() {} - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer(decltype(nullptr)) : super_t(nullptr) {} - #endif - - /*! This constructor allows construction of a pointer from a T*. - * - * \param ptr A raw pointer to copy from, presumed to point to a location in memory - * accessible by the \p cuda system. - * \tparam OtherT \p OtherT shall be convertible to \p T. - */ - template - __host__ __device__ explicit pointer(OtherT *ptr) : super_t(ptr) - { - } - - /*! This constructor allows construction from another pointer-like object with related type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::cuda::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer>::type * = 0) : super_t(other) - { - } - - /*! This constructor allows construction from another pointer-like object with \p void type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::cuda::tag and its element type shall be \p void. - */ - template - __host__ __device__ - explicit - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_void_pointer_is_system_convertible< - OtherPointer, - pointer>::type * = 0) : super_t(other) - { - } - - /*! Assignment operator allows assigning from another pointer-like object with related type. - * - * \param other The other pointer-like object to assign from. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::cuda::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer, - pointer &>::type - operator=(const OtherPointer &other) - { - return super_t::operator=(other); - } - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer& operator=(decltype(nullptr)) - { - super_t::operator=(nullptr); - return *this; - } - #endif -}; // struct pointer - -/*! \p reference is a wrapped reference to an object stored in memory available to the \p cuda system. - * \p reference is the type of the result of dereferencing a \p cuda::pointer. +using pointer = thrust::pointer< + T, + thrust::cuda_cub::tag, + thrust::tagged_reference +>; + +/*! \p cuda::universal_pointer stores a pointer to an object allocated in + * memory accessible by the \p cuda system and host systems. * - * \tparam T Specifies the type of the referenced object. + * \p cuda::universal_pointer has pointer semantics: it may be dereferenced + * and manipulated with pointer arithmetic. + * + * \p cuda::universal_pointer can be created with \p cuda::universal_allocator + * or by explicitly calling its constructor with a raw pointer. + * + * The raw pointer encapsulated by a \p cuda::universal_pointer may be + * obtained by eiter its get member function or the \p + * raw_pointer_cast function. + * + * \note \p cuda::universal_pointer is not a "smart" pointer; it is the + * programmer's responsibility to deallocate memory pointed to by + * \p cuda::universal_pointer. + * + * \tparam T specifies the type of the pointee. + * + * \see cuda::universal_allocator + * \see raw_pointer_cast */ template -class reference - : public thrust::reference< - T, - thrust::cuda_cub::pointer, - thrust::cuda_cub::reference > -{ - -private: - typedef thrust::reference< - T, - thrust::cuda_cub::pointer, - thrust::cuda_cub::reference > - super_t; - -public: - /*! \cond - */ - - typedef typename super_t::value_type value_type; - typedef typename super_t::pointer pointer; - - /*! \endcond - */ - - /*! This constructor initializes this \p reference to refer to an object - * pointed to by the given \p pointer. After this \p reference is constructed, - * it shall refer to the object pointed to by \p ptr. - * - * \param ptr A \p pointer to copy from. - */ - __host__ __device__ explicit reference(const pointer &ptr) - : super_t(ptr) - { - } - - /*! This constructor accepts a const reference to another \p reference of related type. - * After this \p reference is constructed, it shall refer to the same object as \p other. - * - * \param other A \p reference to copy from. - * \tparam OtherT The element type of the other \p reference. - * - * \note This constructor is templated primarily to allow initialization of reference - * from reference. - */ - template - __host__ __device__ - reference(const reference &other, - typename thrust::detail::enable_if_convertible< - typename reference::pointer, - pointer>::type * = 0) - : super_t(other) - { - } - - /*! Copy assignment operator copy assigns from another \p reference of related type. - * - * \param other The other \p reference to assign from. - * \return *this - * \tparam OtherT The element type of the other \p reference. - */ - template - __host__ __device__ - reference & - operator=(const reference &other); - - /*! Assignment operator assigns from a \p value_type. - * - * \param x The \p value_type to assign from. - * \return *this - */ - __host__ __device__ - reference & - operator=(const value_type &x); -}; // struct reference - -/*! Exchanges the values of two objects referred to by \p reference. - * \p x The first \p reference of interest. - * \p y The second \p reference of interest. +using universal_pointer = thrust::pointer< + T, + thrust::cuda_cub::tag, + typename std::add_lvalue_reference::type +>; + +/*! \p cuda::reference is a wrapped reference to an object stored in memory + * accessible by the \p cuda system. \p cuda::reference is the type of the + * result of dereferencing a \p cuda::pointer. + * + * \tparam T Specifies the type of the referenced object. + * + * \see cuda::pointer */ template -__host__ __device__ void swap(reference x, reference y); - -} // end cuda_cub - -namespace system { +using reference = thrust::tagged_reference; +} // namespace cuda_cub /*! \addtogroup system_backends Systems * \ingroup system @@ -291,31 +106,31 @@ namespace system { */ /*! \namespace thrust::system::cuda - * \brief \p thrust::system::cuda is the namespace containing functionality for allocating, manipulating, - * and deallocating memory available to Thrust's CUDA backend system. - * The identifiers are provided in a separate namespace underneath thrust::system - * for import convenience but are also aliased in the top-level thrust::cuda - * namespace for easy access. + * \brief \p thrust::system::cuda is the namespace containing functionality + * for allocating, manipulating, and deallocating memory available to Thrust's + * CUDA backend system. The identifiers are provided in a separate namespace + * underneath thrust::system for import convenience but are also + * aliased in the top-level thrust::cuda namespace for easy access. * */ - -namespace cuda { +namespace system { namespace cuda +{ using thrust::cuda_cub::pointer; +using thrust::cuda_cub::universal_pointer; using thrust::cuda_cub::reference; -} // end cuda - +}} // namespace system::cuda /*! \} */ -} // end system - /*! \namespace thrust::cuda - * \brief \p thrust::cuda is a top-level alias for \p thrust::system::cuda. */ -namespace cuda { + * \brief \p thrust::cuda is a top-level alias for \p thrust::system::cuda. + */ +namespace cuda +{ using thrust::cuda_cub::pointer; +using thrust::cuda_cub::universal_pointer; using thrust::cuda_cub::reference; -} // end cuda +} // namespace cuda -} // end thrust +} // namespace thrust -#include diff --git a/thrust/system/cuda/vector.h b/thrust/system/cuda/vector.h index 9348057a70..7a90a07fb2 100644 --- a/thrust/system/cuda/vector.h +++ b/thrust/system/cuda/vector.h @@ -2,7 +2,7 @@ * Copyright 2008-2013 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in ccudaliance with the License. + * you may not use this file except in compliance with the License. * You may obtain a copy of the License at * * http://www.apache.org/licenses/LICENSE-2.0 @@ -26,47 +26,63 @@ #include #include -namespace thrust +namespace thrust { namespace cuda_cub { -// forward declaration of host_vector -template class host_vector; - -namespace cuda_cub -{ - -/*! \p cuda_bulk::vector is a container that supports random access to elements, +/*! \p cuda::vector is a container that supports random access to elements, * constant time removal of elements at the end, and linear time insertion * and removal of elements at the beginning or in the middle. The number of - * elements in a \p cuda_bulk::vector may vary dynamically; memory management is - * automatic. The elements contained in a \p cuda_bulk::vector reside in memory - * available to the \p cuda_bulk system. + * elements in a \p cuda::vector may vary dynamically; memory management is + * automatic. The elements contained in a \p cuda::vector reside in memory + * accessible by the \p cuda system. * - * \tparam T The element type of the \p cuda_bulk::vector. - * \tparam Allocator The allocator type of the \p cuda_bulk::vector. Defaults to \p cuda_bulk::allocator. + * \tparam T The element type of the \p cuda::vector. + * \tparam Allocator The allocator type of the \p cuda::vector. + * Defaults to \p cuda::allocator. * - * \see http://www.sgi.com/tech/stl/Vector.html + * \see https://en.cppreference.com/w/cpp/container/vector * \see host_vector For the documentation of the complete interface which is - * shared by \p cuda_bulk::vector + * shared by \p cuda::vector * \see device_vector + * \see universal_vector */ -template > +template > using vector = thrust::detail::vector_base; -} // end cuda_cub +/*! \p cuda::universal_vector is a container that supports random access to + * elements, constant time removal of elements at the end, and linear time + * insertion and removal of elements at the beginning or in the middle. The + * number of elements in a \p cuda::universal_vector may vary dynamically; + * memory management is automatic. The elements contained in a + * \p cuda::universal_vector reside in memory accessible by the \p cuda system + * and host systems. + * + * \tparam T The element type of the \p cuda::universal_vector. + * \tparam Allocator The allocator type of the \p cuda::universal_vector. + * Defaults to \p cuda::universal_allocator. + * + * \see https://en.cppreference.com/w/cpp/container/vector + * \see host_vector For the documentation of the complete interface which is + * shared by \p cuda::universal_vector + * \see device_vector + * \see universal_vector + */ +template > +using universal_vector = thrust::detail::vector_base; + +} // namespace cuda_cub -// alias system::cuda_bulk names at top-level -namespace cuda +namespace system { namespace cuda { - using thrust::cuda_cub::vector; +using thrust::cuda_cub::universal_vector; +}} -} // end cuda_bulk - -namespace system { -namespace cuda { +namespace cuda +{ using thrust::cuda_cub::vector; +using thrust::cuda_cub::universal_vector; } -} -} // end thrust +} // namespace thrust + diff --git a/thrust/system/omp/detail/pointer.inl b/thrust/system/omp/detail/pointer.inl deleted file mode 100644 index 2125302e4c..0000000000 --- a/thrust/system/omp/detail/pointer.inl +++ /dev/null @@ -1,52 +0,0 @@ -/* - * Copyright 2008-2018 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -namespace thrust -{ -namespace system -{ -namespace omp -{ - - -template - template - reference & - reference - ::operator=(const reference &other) -{ - return super_t::operator=(other); -} // end reference::operator=() - -template - reference & - reference - ::operator=(const value_type &x) -{ - return super_t::operator=(x); -} // end reference::operator=() - -template -__host__ __device__ -void swap(reference a, reference b) -{ - a.swap(b); -} // end swap() - -} // end omp -} // end system -} // end thrust - diff --git a/thrust/system/omp/memory.h b/thrust/system/omp/memory.h index 9b2f070ccd..ff59036ba6 100644 --- a/thrust/system/omp/memory.h +++ b/thrust/system/omp/memory.h @@ -27,11 +27,7 @@ #include #include -namespace thrust -{ -namespace system -{ -namespace omp +namespace thrust { namespace system { namespace omp { /*! Allocates an area of memory available to Thrust's omp system. @@ -67,29 +63,38 @@ inline pointer malloc(std::size_t n); */ inline void free(pointer ptr); -/*! \p omp::allocator is the default allocator used by the \p omp system's containers such as - * omp::vector if no user-specified allocator is provided. \p omp::allocator allocates - * (deallocates) storage with \p omp::malloc (\p omp::free). +/*! \p omp::allocator is the default allocator used by the \p omp system's + * containers such as omp::vector if no user-specified allocator is + * provided. \p omp::allocator allocates (deallocates) storage with \p + * omp::malloc (\p omp::free). + */ +template +using allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::omp::memory_resource +>; + +/*! \p omp::universal_allocator allocates memory that can be used by the \p omp + * system and host systems. */ template -using allocator = thrust::mr::stateless_resource_allocator; +using universal_allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::omp::universal_memory_resource +>; -} // end omp -} // end system +}} // namespace system::omp /*! \namespace thrust::omp * \brief \p thrust::omp is a top-level alias for thrust::system::omp. */ namespace omp { - using thrust::system::omp::malloc; using thrust::system::omp::free; using thrust::system::omp::allocator; +using thrust::system::omp::universal_allocator; +} // namespace omp -} // end omp - -} // end thrust +} // namespace thrust #include diff --git a/thrust/system/omp/pointer.h b/thrust/system/omp/pointer.h index 36b6bed12a..d72069bd86 100644 --- a/thrust/system/omp/pointer.h +++ b/thrust/system/omp/pointer.h @@ -1,5 +1,5 @@ /* - * Copyright 2008-2018 NVIDIA Corporation + * Copyright 2008-2020 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,113 +21,29 @@ #pragma once #include +#include #include -#include #include #include -namespace thrust +namespace thrust { namespace system { namespace omp { -namespace system -{ -namespace omp -{ - -template class pointer; - -} // end omp -} // end system -} // end thrust - - -/*! \cond - */ - -// specialize thrust::iterator_traits to avoid problems with the name of -// pointer's constructor shadowing its nested pointer type -// do this before pointer is defined so the specialization is correctly -// used inside the definition -namespace thrust -{ - -template - struct iterator_traits > -{ - private: - typedef thrust::system::omp::pointer ptr; - - public: - typedef typename ptr::iterator_category iterator_category; - typedef typename ptr::value_type value_type; - typedef typename ptr::difference_type difference_type; - typedef ptr pointer; - typedef typename ptr::reference reference; -}; // end iterator_traits - -} // end thrust - -/*! \endcond - */ - - -namespace thrust -{ -namespace system -{ - -/*! \addtogroup system_backends Systems - * \ingroup system - * \{ - */ -/*! \namespace thrust::system::omp - * \brief \p thrust::system::omp is the namespace containing functionality for allocating, manipulating, - * and deallocating memory available to Thrust's OpenMP backend system. - * The identifiers are provided in a separate namespace underneath thrust::system - * for import convenience but are also aliased in the top-level thrust::omp - * namespace for easy access. +/*! \p omp::pointer stores a pointer to an object allocated in memory accessible + * by the \p omp system. This type provides type safety when dispatching + * algorithms on ranges resident in \p omp memory. * - */ -namespace omp -{ - -// forward declaration of reference for pointer -template class reference; - -/*! \cond - */ - -// XXX nvcc + msvc have trouble instantiating reference below -// this is a workaround -namespace detail -{ - -template - struct reference_msvc_workaround -{ - typedef thrust::system::omp::reference type; -}; // end reference_msvc_workaround - -} // end detail - -/*! \endcond - */ - - -/*! \p pointer stores a pointer to an object allocated in memory available to the omp system. - * This type provides type safety when dispatching standard algorithms on ranges resident - * in omp memory. + * \p omp::pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. * - * \p pointer has pointer semantics: it may be dereferenced and manipulated with pointer arithmetic. + * \p omp::pointer can be created with the function \p omp::malloc, or by + * explicitly calling its constructor with a raw pointer. * - * \p pointer can be created with the function \p omp::malloc, or by explicitly calling its constructor - * with a raw pointer. + * The raw pointer encapsulated by a \p omp::pointer may be obtained by eiter its + * get member function or the \p raw_pointer_cast function. * - * The raw pointer encapsulated by a \p pointer may be obtained by eiter its get member function - * or the \p raw_pointer_cast function. - * - * \note \p pointer is not a "smart" pointer; it is the programmer's responsibility to deallocate memory - * pointed to by \p pointer. + * \note \p omp::pointer is not a "smart" pointer; it is the programmer's + * responsibility to deallocate memory pointed to by \p omp::pointer. * * \tparam T specifies the type of the pointee. * @@ -135,226 +51,66 @@ template * \see omp::free * \see raw_pointer_cast */ -template - class pointer - : public thrust::pointer< - T, - thrust::system::omp::tag, - thrust::system::omp::reference, - thrust::system::omp::pointer - > -{ - /*! \cond - */ - - private: - typedef thrust::pointer< - T, - thrust::system::omp::tag, - //thrust::system::omp::reference, - typename detail::reference_msvc_workaround::type, - thrust::system::omp::pointer - > super_t; - - /*! \endcond - */ - - public: - // note that omp::pointer's member functions need __host__ __device__ - // to interoperate with nvcc + iterators' dereference member function - - /*! \p pointer's no-argument constructor initializes its encapsulated pointer to \c 0. - */ - __host__ __device__ - pointer() : super_t() {} - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer(decltype(nullptr)) : super_t(nullptr) {} - #endif - - /*! This constructor allows construction of a pointer from a T*. - * - * \param ptr A raw pointer to copy from, presumed to point to a location in memory - * accessible by the \p omp system. - * \tparam OtherT \p OtherT shall be convertible to \p T. - */ - template - __host__ __device__ - explicit pointer(OtherT *ptr) : super_t(ptr) {} - - /*! This constructor allows construction from another pointer-like object with related type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::omp::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer - >::type * = 0) : super_t(other) {} - - /*! This constructor allows construction from another pointer-like object with \p void type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::omp::tag and its element type shall be \p void. - */ - template - __host__ __device__ - explicit - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_void_pointer_is_system_convertible< - OtherPointer, - pointer - >::type * = 0) : super_t(other) {} - - /*! Assignment operator allows assigning from another pointer-like object with related type. - * - * \param other The other pointer-like object to assign from. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::omp::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer, - pointer & - >::type - operator=(const OtherPointer &other) - { - return super_t::operator=(other); - } - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer& operator=(decltype(nullptr)) - { - super_t::operator=(nullptr); - return *this; - } - #endif -}; // end pointer - - -/*! \p reference is a wrapped reference to an object stored in memory available to the \p omp system. - * \p reference is the type of the result of dereferencing a \p omp::pointer. +template +using pointer = thrust::pointer< + T, + thrust::system::omp::tag, + thrust::tagged_reference +>; + +/*! \p omp::universal_pointer stores a pointer to an object allocated in memory + * accessible by the \p omp system and host systems. * - * \tparam T Specifies the type of the referenced object. + * \p omp::universal_pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. + * + * \p omp::universal_pointer can be created with \p omp::universal_allocator + * or by explicitly calling its constructor with a raw pointer. + * + * The raw pointer encapsulated by a \p omp::universal_pointer may be obtained + * by eiter its get member function or the \p raw_pointer_cast + * function. + * + * \note \p omp::universal_pointer is not a "smart" pointer; it is the + * programmer's responsibility to deallocate memory pointed to by + * \p omp::universal_pointer. + * + * \tparam T specifies the type of the pointee. + * + * \see omp::universal_allocator + * \see raw_pointer_cast */ -template - class reference - : public thrust::reference< - T, - thrust::system::omp::pointer, - thrust::system::omp::reference - > -{ - /*! \cond - */ - - private: - typedef thrust::reference< - T, - thrust::system::omp::pointer, - thrust::system::omp::reference - > super_t; - - /*! \endcond - */ - - public: - /*! \cond - */ - - typedef typename super_t::value_type value_type; - typedef typename super_t::pointer pointer; - - /*! \endcond - */ - - /*! This constructor initializes this \p reference to refer to an object - * pointed to by the given \p pointer. After this \p reference is constructed, - * it shall refer to the object pointed to by \p ptr. - * - * \param ptr A \p pointer to copy from. - */ - __host__ __device__ - explicit reference(const pointer &ptr) - : super_t(ptr) - {} - - /*! This constructor accepts a const reference to another \p reference of related type. - * After this \p reference is constructed, it shall refer to the same object as \p other. - * - * \param other A \p reference to copy from. - * \tparam OtherT The element type of the other \p reference. - * - * \note This constructor is templated primarily to allow initialization of reference - * from reference. - */ - template - __host__ __device__ - reference(const reference &other, - typename thrust::detail::enable_if_convertible< - typename reference::pointer, - pointer - >::type * = 0) - : super_t(other) - {} - - /*! Copy assignment operator copy assigns from another \p reference of related type. - * - * \param other The other \p reference to assign from. - * \return *this - * \tparam OtherT The element type of the other \p reference. - */ - template - reference &operator=(const reference &other); - - /*! Assignment operator assigns from a \p value_type. - * - * \param x The \p value_type to assign from. - * \return *this - */ - reference &operator=(const value_type &x); -}; // end reference - -/*! Exchanges the values of two objects referred to by \p reference. - * \p x The first \p reference of interest. - * \p y The second \p reference of interest. +template +using universal_pointer = thrust::pointer< + T, + thrust::system::omp::tag, + typename std::add_lvalue_reference::type +>; + +/*! \p reference is a wrapped reference to an object stored in memory available + * to the \p omp system. \p reference is the type of the result of + * dereferencing a \p omp::pointer. + * + * \tparam T Specifies the type of the referenced object. */ -template -__host__ __device__ -void swap(reference x, reference y); +template +using reference = thrust::tagged_reference; -} // end omp +}} // namespace system::omp -/*! \} +/*! \addtogroup system_backends Systems + * \ingroup system + * \{ */ -} // end system - /*! \namespace thrust::omp - * \brief \p thrust::omp is a top-level alias for thrust::system::omp. - */ + * \brief \p thrust::omp is a top-level alias for \p thrust::system::omp. */ namespace omp { - using thrust::system::omp::pointer; +using thrust::system::omp::universal_pointer; using thrust::system::omp::reference; +} // namespace omp -} // end omp - -} // end thrust - -#include +} // namespace thrust diff --git a/thrust/system/omp/vector.h b/thrust/system/omp/vector.h index 101a22c7b2..dead9f5921 100644 --- a/thrust/system/omp/vector.h +++ b/thrust/system/omp/vector.h @@ -26,16 +26,7 @@ #include #include -namespace thrust -{ - -// forward declaration of host_vector -// XXX why is this here? it doesn't seem necessary for anything below -template class host_vector; - -namespace system -{ -namespace omp +namespace thrust { namespace system { namespace omp { /*! \p omp::vector is a container that supports random access to elements, @@ -43,28 +34,48 @@ namespace omp * and removal of elements at the beginning or in the middle. The number of * elements in a \p omp::vector may vary dynamically; memory management is * automatic. The elements contained in an \p omp::vector reside in memory - * available to the \p omp system. + * accessible by the \p omp system. * * \tparam T The element type of the \p omp::vector. - * \tparam Allocator The allocator type of the \p omp::vector. Defaults to \p omp::allocator. + * \tparam Allocator The allocator type of the \p omp::vector. + * Defaults to \p omp::allocator. * - * \see http://www.sgi.com/tech/stl/Vector.html + * \see https://en.cppreference.com/w/cpp/container/vector * \see host_vector For the documentation of the complete interface which is - * shared by \p omp::vector + * shared by \p omp::vector. * \see device_vector + * \see universal_vector */ -template > +template > using vector = thrust::detail::vector_base; -} // end omp -} // end system +/*! \p omp::universal_vector is a container that supports random access to + * elements, constant time removal of elements at the end, and linear time + * insertion and removal of elements at the beginning or in the middle. The + * number of elements in a \p omp::universal_vector may vary dynamically; + * memory management is automatic. The elements contained in a + * \p omp::universal_vector reside in memory accessible by the \p omp system + * and host systems. + * + * \tparam T The element type of the \p omp::universal_vector. + * \tparam Allocator The allocator type of the \p omp::universal_vector. + * Defaults to \p omp::universal_allocator. + * + * \see https://en.cppreference.com/w/cpp/container/vector + * \see host_vector For the documentation of the complete interface which is + * shared by \p omp::universal_vector + * \see device_vector + * \see universal_vector + */ +template > +using universal_vector = thrust::detail::vector_base; + +}} // namespace system::omp -// alias system::omp names at top-level namespace omp { - using thrust::system::omp::vector; - -} // end omp +using thrust::system::omp::universal_vector; +} } // end thrust diff --git a/thrust/system/tbb/detail/pointer.inl b/thrust/system/tbb/detail/pointer.inl deleted file mode 100644 index 2b21422bc6..0000000000 --- a/thrust/system/tbb/detail/pointer.inl +++ /dev/null @@ -1,53 +0,0 @@ -/* - * Copyright 2008-2018 NVIDIA Corporation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - - -namespace thrust -{ -namespace system -{ -namespace tbb -{ - - -template - template - reference & - reference - ::operator=(const reference &other) -{ - return super_t::operator=(other); -} // end reference::operator=() - -template - reference & - reference - ::operator=(const value_type &x) -{ - return super_t::operator=(x); -} // end reference::operator=() - -template -__host__ __device__ -void swap(reference a, reference b) -{ - a.swap(b); -} // end swap() - -} // end tbb -} // end system -} // end thrust - diff --git a/thrust/system/tbb/memory.h b/thrust/system/tbb/memory.h index a680157006..8320584741 100644 --- a/thrust/system/tbb/memory.h +++ b/thrust/system/tbb/memory.h @@ -2,7 +2,7 @@ * Copyright 2008-2018 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. + * you may not use this file except in ctbbliance with the License. * You may obtain a copy of the License at * * http://www.apache.org/licenses/LICENSE-2.0 @@ -67,33 +67,38 @@ inline pointer malloc(std::size_t n); */ inline void free(pointer ptr); -/*! \p tbb::allocator is the default allocator used by the \p tbb system's containers such as - * tbb::vector if no user-specified allocator is provided. \p tbb::allocator allocates - * (deallocates) storage with \p tbb::malloc (\p tbb::free). +/*! \p tbb::allocator is the default allocator used by the \p tbb system's + * containers such as tbb::vector if no user-specified allocator is + * provided. \p tbb::allocator allocates (deallocates) storage with \p + * tbb::malloc (\p tbb::free). */ template -using allocator = thrust::mr::stateless_resource_allocator; +using allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::tbb::memory_resource +>; -} // end tbb - -/*! \} +/*! \p tbb::universal_allocator allocates memory that can be used by the \p tbb + * system and host systems. */ +template +using universal_allocator = thrust::mr::stateless_resource_allocator< + T, thrust::system::tbb::universal_memory_resource +>; -} // end system +}} // namespace system::tbb /*! \namespace thrust::tbb * \brief \p thrust::tbb is a top-level alias for thrust::system::tbb. */ namespace tbb { - using thrust::system::tbb::malloc; using thrust::system::tbb::free; using thrust::system::tbb::allocator; +using thrust::system::tbb::universal_allocator; +} // namsespace tbb -} // end tbb - -} // end thrust +} // namespace thrust #include diff --git a/thrust/system/tbb/pointer.h b/thrust/system/tbb/pointer.h index d2912508a5..ad01f44a72 100644 --- a/thrust/system/tbb/pointer.h +++ b/thrust/system/tbb/pointer.h @@ -1,5 +1,5 @@ /* - * Copyright 2008-2018 NVIDIA Corporation + * Copyright 2008-2020 NVIDIA Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,114 +14,36 @@ * limitations under the License. */ +/*! \file thrust/system/tbb/memory.h + * \brief Managing memory associated with Thrust's TBB system. + */ + +#pragma once + #include +#include #include -#include #include #include -namespace thrust -{ -namespace system -{ -namespace tbb -{ - -template class pointer; - -} // end tbb -} // end system -} // end thrust - - -/*! \cond - */ - -// specialize thrust::iterator_traits to avoid problems with the name of -// pointer's constructor shadowing its nested pointer type -// do this before pointer is defined so the specialization is correctly -// used inside the definition -namespace thrust -{ - -template - struct iterator_traits > -{ - private: - typedef thrust::system::tbb::pointer ptr; - - public: - typedef typename ptr::iterator_category iterator_category; - typedef typename ptr::value_type value_type; - typedef typename ptr::difference_type difference_type; - typedef ptr pointer; - typedef typename ptr::reference reference; -}; // end iterator_traits - -} // end thrust - -/*! \endcond - */ - - -namespace thrust +namespace thrust { namespace system { namespace tbb { -namespace system -{ - -/*! \addtogroup system_backends Systems - * \ingroup system - * \{ - */ - -/*! \namespace thrust::system::tbb - * \brief \p thrust::system::tbb is the namespace containing functionality for allocating, manipulating, - * and deallocating memory available to Thrust's TBB backend system. - * The identifiers are provided in a separate namespace underneath thrust::system - * for import convenience but are also aliased in the top-level thrust::tbb - * namespace for easy access. - * - */ -namespace tbb -{ - -// forward declaration of reference for pointer -template class reference; - -/*! \cond - */ - -// XXX nvcc + msvc have trouble instantiating reference below -// this is a workaround -namespace detail -{ - -template - struct reference_msvc_workaround -{ - typedef thrust::system::tbb::reference type; -}; // end reference_msvc_workaround - -} // end detail - -/*! \endcond - */ - -/*! \p pointer stores a pointer to an object allocated in memory available to the tbb system. - * This type provides type safety when dispatching standard algorithms on ranges resident - * in tbb memory. +/*! \p tbb::pointer stores a pointer to an object allocated in memory accessible + * by the \p tbb system. This type provides type safety when dispatching + * algorithms on ranges resident in \p tbb memory. * - * \p pointer has pointer semantics: it may be dereferenced and manipulated with pointer arithmetic. + * \p tbb::pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. * - * \p pointer can be created with the function \p tbb::malloc, or by explicitly calling its constructor - * with a raw pointer. + * \p tbb::pointer can be created with the function \p tbb::malloc, or by + * explicitly calling its constructor with a raw pointer. * - * The raw pointer encapsulated by a \p pointer may be obtained by eiter its get member function - * or the \p raw_pointer_cast function. + * The raw pointer encapsulated by a \p tbb::pointer may be obtained by eiter its + * get member function or the \p raw_pointer_cast function. * - * \note \p pointer is not a "smart" pointer; it is the programmer's responsibility to deallocate memory - * pointed to by \p pointer. + * \note \p tbb::pointer is not a "smart" pointer; it is the programmer's + * responsibility to deallocate memory pointed to by \p tbb::pointer. * * \tparam T specifies the type of the pointee. * @@ -129,226 +51,66 @@ template * \see tbb::free * \see raw_pointer_cast */ -template - class pointer - : public thrust::pointer< - T, - thrust::system::tbb::tag, - thrust::system::tbb::reference, - thrust::system::tbb::pointer - > -{ - /*! \cond - */ - - private: - typedef thrust::pointer< - T, - thrust::system::tbb::tag, - //thrust::system::tbb::reference, - typename detail::reference_msvc_workaround::type, - thrust::system::tbb::pointer - > super_t; - - /*! \endcond - */ - - public: - // note that tbb::pointer's member functions need __host__ __device__ - // to interoperate with nvcc + iterators' dereference member function - - /*! \p pointer's no-argument constructor initializes its encapsulated pointer to \c 0. - */ - __host__ __device__ - pointer() : super_t() {} - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer(decltype(nullptr)) : super_t(nullptr) {} - #endif - - /*! This constructor allows construction of a pointer from a T*. - * - * \param ptr A raw pointer to copy from, presumed to point to a location in memory - * accessible by the \p tbb system. - * \tparam OtherT \p OtherT shall be convertible to \p T. - */ - template - __host__ __device__ - explicit pointer(OtherT *ptr) : super_t(ptr) {} - - /*! This constructor allows construction from another pointer-like object with related type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::tbb::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer - >::type * = 0) : super_t(other) {} - - /*! This constructor allows construction from another pointer-like object with \p void type. - * - * \param other The \p OtherPointer to copy. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::tbb::tag and its element type shall be \p void. - */ - template - __host__ __device__ - explicit - pointer(const OtherPointer &other, - typename thrust::detail::enable_if_void_pointer_is_system_convertible< - OtherPointer, - pointer - >::type * = 0) : super_t(other) {} - - /*! Assignment operator allows assigning from another pointer-like object with related type. - * - * \param other The other pointer-like object to assign from. - * \tparam OtherPointer The system tag associated with \p OtherPointer shall be convertible - * to \p thrust::system::tbb::tag and its element type shall be convertible to \p T. - */ - template - __host__ __device__ - typename thrust::detail::enable_if_pointer_is_convertible< - OtherPointer, - pointer, - pointer & - >::type - operator=(const OtherPointer &other) - { - return super_t::operator=(other); - } - - #if THRUST_CPP_DIALECT >= 2011 - // NOTE: This is needed so that Thrust smart pointers can be used in - // `std::unique_ptr`. - __host__ __device__ - pointer& operator=(decltype(nullptr)) - { - super_t::operator=(nullptr); - return *this; - } - #endif -}; // end pointer - - -/*! \p reference is a wrapped reference to an object stored in memory available to the \p tbb system. - * \p reference is the type of the result of dereferencing a \p tbb::pointer. +template +using pointer = thrust::pointer< + T, + thrust::system::tbb::tag, + thrust::tagged_reference +>; + +/*! \p tbb::universal_pointer stores a pointer to an object allocated in memory + * accessible by the \p tbb system and host systems. * - * \tparam T Specifies the type of the referenced object. + * \p tbb::universal_pointer has pointer semantics: it may be dereferenced and + * manipulated with pointer arithmetic. + * + * \p tbb::universal_pointer can be created with \p tbb::universal_allocator + * or by explicitly calling its constructor with a raw pointer. + * + * The raw pointer encapsulated by a \p tbb::universal_pointer may be obtained + * by eiter its get member function or the \p raw_pointer_cast + * function. + * + * \note \p tbb::universal_pointer is not a "smart" pointer; it is the + * programmer's responsibility to deallocate memory pointed to by + * \p tbb::universal_pointer. + * + * \tparam T specifies the type of the pointee. + * + * \see tbb::universal_allocator + * \see raw_pointer_cast */ -template - class reference - : public thrust::reference< - T, - thrust::system::tbb::pointer, - thrust::system::tbb::reference - > -{ - /*! \cond - */ - - private: - typedef thrust::reference< - T, - thrust::system::tbb::pointer, - thrust::system::tbb::reference - > super_t; - - /*! \endcond - */ - - public: - /*! \cond - */ - - typedef typename super_t::value_type value_type; - typedef typename super_t::pointer pointer; - - /*! \endcond - */ - - /*! This constructor initializes this \p reference to refer to an object - * pointed to by the given \p pointer. After this \p reference is constructed, - * it shall refer to the object pointed to by \p ptr. - * - * \param ptr A \p pointer to copy from. - */ - __host__ __device__ - explicit reference(const pointer &ptr) - : super_t(ptr) - {} - - /*! This constructor accepts a const reference to another \p reference of related type. - * After this \p reference is constructed, it shall refer to the same object as \p other. - * - * \param other A \p reference to copy from. - * \tparam OtherT The element type of the other \p reference. - * - * \note This constructor is templated primarily to allow initialization of reference - * from reference. - */ - template - __host__ __device__ - reference(const reference &other, - typename thrust::detail::enable_if_convertible< - typename reference::pointer, - pointer - >::type * = 0) - : super_t(other) - {} - - /*! Copy assignment operator copy assigns from another \p reference of related type. - * - * \param other The other \p reference to assign from. - * \return *this - * \tparam OtherT The element type of the other \p reference. - */ - template - reference &operator=(const reference &other); - - /*! Assignment operator assigns from a \p value_type. - * - * \param x The \p value_type to assign from. - * \return *this - */ - reference &operator=(const value_type &x); -}; // end reference - -/*! Exchanges the values of two objects referred to by \p reference. - * \p x The first \p reference of interest. - * \p y The second \p reference ot interest. +template +using universal_pointer = thrust::pointer< + T, + thrust::system::tbb::tag, + typename std::add_lvalue_reference::type +>; + +/*! \p reference is a wrapped reference to an object stored in memory available + * to the \p tbb system. \p reference is the type of the result of + * dereferencing a \p tbb::pointer. + * + * \tparam T Specifies the type of the referenced object. */ -template -__host__ __device__ -void swap(reference x, reference y); +template +using reference = thrust::tagged_reference; -} // end tbb +}} // namespace system::tbb -/*! \} +/*! \addtogroup system_backends Systems + * \ingroup system + * \{ */ -} // end system - /*! \namespace thrust::tbb - * \brief \p thrust::tbb is a top-level alias for thrust::system::tbb. - */ + * \brief \p thrust::tbb is a top-level alias for \p thrust::system::tbb. */ namespace tbb { - using thrust::system::tbb::pointer; +using thrust::system::tbb::universal_pointer; using thrust::system::tbb::reference; +} // namespace tbb -} // end tbb - -} // end thrust - -#include +} // namespace thrust diff --git a/thrust/system/tbb/vector.h b/thrust/system/tbb/vector.h index 0e08c8cf0d..e5d148416c 100644 --- a/thrust/system/tbb/vector.h +++ b/thrust/system/tbb/vector.h @@ -26,11 +26,7 @@ #include #include -namespace thrust -{ -namespace system -{ -namespace tbb +namespace thrust { namespace system { namespace tbb { /*! \p tbb::vector is a container that supports random access to elements, @@ -38,28 +34,48 @@ namespace tbb * and removal of elements at the beginning or in the middle. The number of * elements in a \p tbb::vector may vary dynamically; memory management is * automatic. The elements contained in a \p tbb::vector reside in memory - * available to the \p tbb system. + * accessible by the \p tbb system. * * \tparam T The element type of the \p tbb::vector. - * \tparam Allocator The allocator type of the \p tbb::vector. Defaults to \p tbb::allocator. + * \tparam Allocator The allocator type of the \p tbb::vector. + * Defaults to \p tbb::allocator. * - * \see http://www.sgi.com/tech/stl/Vector.html + * \see https://en.cppreference.com/w/cpp/container/vector * \see host_vector For the documentation of the complete interface which is - * shared by \p tbb::vector + * shared by \p tbb::vector. * \see device_vector + * \see universal_vector */ -template > +template > using vector = thrust::detail::vector_base; -} // end tbb -} // end system +/*! \p tbb::universal_vector is a container that supports random access to + * elements, constant time removal of elements at the end, and linear time + * insertion and removal of elements at the beginning or in the middle. The + * number of elements in a \p tbb::universal_vector may vary dynamically; + * memory management is automatic. The elements contained in a + * \p tbb::universal_vector reside in memory accessible by the \p tbb system + * and host systems. + * + * \tparam T The element type of the \p tbb::universal_vector. + * \tparam Allocator The allocator type of the \p tbb::universal_vector. + * Defaults to \p tbb::universal_allocator. + * + * \see https://en.cppreference.com/w/cpp/container/vector + * \see host_vector For the documentation of the complete interface which is + * shared by \p tbb::universal_vector + * \see device_vector + * \see universal_vector + */ +template > +using universal_vector = thrust::detail::vector_base; + +}} // namespace system::tbb -// alias system::tbb names at top-level namespace tbb { - using thrust::system::tbb::vector; +using thrust::system::tbb::universal_vector; +} -} // end tbb - -} // end thrust +} // namespace thrust diff --git a/thrust/universal_allocator.h b/thrust/universal_allocator.h new file mode 100644 index 0000000000..dcd08d8d40 --- /dev/null +++ b/thrust/universal_allocator.h @@ -0,0 +1,79 @@ +/* + * Copyright 2008-2020 NVIDIA Corporation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + + +/*! \file universal_allocator.h + * \brief An allocator which creates new elements in memory accessible to both + * hosts and devices. + */ + +#pragma once + +#include + +// #include the device system's vector header +#define __THRUST_DEVICE_SYSTEM_MEMORY_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/memory.h> +#include __THRUST_DEVICE_SYSTEM_MEMORY_HEADER +#undef __THRUST_DEVICE_SYSTEM_MEMORY_HEADER + +namespace thrust +{ + +/** \addtogroup memory_resources Memory Resources + * \ingroup memory_management_classes + * \{ + */ + +/*! \brief An allocator which creates new elements in memory accessible by + * both hosts and devices. + * + * \see https://en.cppreference.com/w/cpp/named_req/Allocator + */ +using thrust::system::__THRUST_DEVICE_SYSTEM_NAMESPACE::universal_allocator; + +/*! \p universal_ptr stores a pointer to an object allocated in memory accessible + * to both hosts and devices. + * + * Algorithms dispatched with this type of pointer will be dispatched to + * either host or device, depending on which backend you are using. Explicit + * policies (\p thrust::device, etc) can be used to specify where an algorithm + * should be run. + * + * \p universal_ptr has pointer semantics: it may be dereferenced safely from + * both hosts and devices and may be manipulated with pointer arithmetic. + * + * \p universal_ptr can be created with \p universal_allocator or by explicitly + * calling its constructor with a raw pointer. + * + * The raw pointer encapsulated by a \p universal_ptr may be obtained by + * either its get method or the \p raw_pointer_cast free function. + * + * \note \p universal_ptr is not a smart pointer; it is the programmer's + * responsibility to deallocate memory pointed to by \p universal_ptr. + * + * \see host_ptr For the documentation of the complete interface which is + * shared by \p universal_ptr. + * \see raw_pointer_cast + */ +template +using universal_ptr = + thrust::system::__THRUST_DEVICE_SYSTEM_NAMESPACE::universal_pointer; + +/*! \} + */ + +} // end thrust + diff --git a/thrust/universal_ptr.h b/thrust/universal_ptr.h new file mode 100644 index 0000000000..9d1de19d5f --- /dev/null +++ b/thrust/universal_ptr.h @@ -0,0 +1,26 @@ +/* + * Copyright 2008-2020 NVIDIA Corporation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + + +/*! \file universal_ptr.h + * \brief A pointer to a variable which resides memory accessible to both + * hosts and devices. + */ + +#pragma once + +#include + diff --git a/thrust/universal_vector.h b/thrust/universal_vector.h new file mode 100644 index 0000000000..485f4815bb --- /dev/null +++ b/thrust/universal_vector.h @@ -0,0 +1,59 @@ +/* + * Copyright 2008-2020 NVIDIA Corporation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + + +/*! \file universal_vector.h + * \brief A dynamically-sizable array of elements which resides in memory + * accessible to both hosts and devices. + */ + +#pragma once + +#include +#include + +// #include the device system's vector header +#define __THRUST_DEVICE_SYSTEM_VECTOR_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/vector.h> +#include __THRUST_DEVICE_SYSTEM_VECTOR_HEADER +#undef __THRUST_DEVICE_SYSTEM_VECTOR_HEADER + +namespace thrust +{ + +/** \addtogroup memory_resources Memory Resources + * \ingroup memory_management_classes + * \{ + */ + +/*! A \p universal_vector is a container that supports random access to elements, + * constant time removal of elements at the end, and linear time insertion + * and removal of elements at the beginning or in the middle. The number of + * elements in a \p universal_vector may vary dynamically; memory management is + * automatic. The memory associated with a \p universal_vector resides in memory + * accessible to hosts and devices. + * + * \see https://en.cppreference.com/w/cpp/container/vector + * \see host_vector For the documentation of the complete interface which is + * shared by \p universal_vector. + * \see device_vector + */ +using thrust::system::__THRUST_DEVICE_SYSTEM_NAMESPACE::universal_vector; + +/*! \} + */ + +} // end thrust +