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..1c6dde9498 100644 --- a/testing/unittest/testframework.h +++ b/testing/unittest/testframework.h @@ -13,8 +13,9 @@ #include #include -#include -#include +#include +#include +#include #include // define some common lists of types @@ -359,7 +360,7 @@ class NAME##UnitTest : public UnitTest { \ public: \ NAME##UnitTest() : UnitTest(#NAME) {} \ void run(){ \ - TEST(); \ + TEST(); \ } \ }; \ NAME##UnitTest NAME##Instance @@ -388,15 +389,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 +412,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/caching_allocator.h b/thrust/detail/caching_allocator.h index bb98f815f7..13df1d33f9 100644 --- a/thrust/detail/caching_allocator.h +++ b/thrust/detail/caching_allocator.h @@ -19,7 +19,7 @@ #include #include #include -#include +#include namespace thrust { diff --git a/thrust/mr/detail/config.h b/thrust/detail/config/memory_resource.h similarity index 100% rename from thrust/mr/detail/config.h rename to thrust/detail/config/memory_resource.h diff --git a/thrust/detail/pointer.h b/thrust/detail/pointer.h index e9204978f5..0f48de11a8 100644 --- a/thrust/detail/pointer.h +++ b/thrust/detail/pointer.h @@ -19,6 +19,7 @@ #include #include #include +#include #include #include #include @@ -52,7 +53,7 @@ template typedef typename ptr::value_type value_type; typedef typename ptr::difference_type difference_type; // XXX implement this type (the result of operator->) later - typedef void pointer; + typedef void pointer; typedef typename ptr::reference reference; }; // end iterator_traits @@ -72,7 +73,7 @@ template // void pointers should have no element type // note that we remove_cv from the Element type to get the value_type typedef typename thrust::detail::eval_if< - thrust::detail::is_void::type>::value, + thrust::detail::is_void::type>::value, thrust::detail::identity_, thrust::detail::remove_cv >::type value_type; @@ -87,14 +88,14 @@ template // void pointers should have no reference type // if no Reference type is given, just use reference typedef typename thrust::detail::eval_if< - thrust::detail::is_void::type>::value, + thrust::detail::is_void::type>::value, thrust::detail::identity_, thrust::detail::eval_if< thrust::detail::is_same::value, 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 +103,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 @@ -146,12 +147,10 @@ template __host__ __device__ pointer(); - #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)); - #endif // OtherValue shall be convertible to Value // XXX consider making the pointer implementation a template parameter which defaults to Element * @@ -182,12 +181,10 @@ template // assignment - #if THRUST_CPP_DIALECT >= 2011 // NOTE: This is needed so that Thrust smart pointers can be used in // `std::unique_ptr`. __host__ __device__ derived_type& operator=(decltype(nullptr)); - #endif // OtherPointer's element_type shall be convertible to Element // OtherPointer's system shall be convertible to Tag @@ -205,12 +202,13 @@ template __host__ __device__ Element *get() const; - #if THRUST_CPP_DIALECT >= 2011 + __host__ __device__ + Element *operator->() const; + // NOTE: This is needed so that Thrust smart pointers can be used in // `std::unique_ptr`. __host__ __device__ explicit operator bool() const; - #endif __host__ __device__ static derived_type pointer_to(typename thrust::detail::pointer_traits_detail::pointer_to_param::type r) @@ -227,7 +225,6 @@ std::basic_ostream & operator<<(std::basic_ostream &os, const pointer &p); -#if THRUST_CPP_DIALECT >= 2011 // NOTE: This is needed so that Thrust smart pointers can be used in // `std::unique_ptr`. template @@ -245,7 +242,6 @@ bool operator!=(decltype(nullptr), pointer p); template __host__ __device__ bool operator!=(pointer p, decltype(nullptr)); -#endif } // end thrust diff --git a/thrust/detail/pointer.inl b/thrust/detail/pointer.inl index 464c3579ed..bd5e340db2 100644 --- a/thrust/detail/pointer.inl +++ b/thrust/detail/pointer.inl @@ -27,24 +27,16 @@ template __host__ __device__ pointer ::pointer() - : super_t(static_cast( - #if THRUST_CPP_DIALECT >= 2011 - nullptr - #else - 0 - #endif - )) + : super_t(static_cast(nullptr)) {} // end pointer::pointer -#if THRUST_CPP_DIALECT >= 2011 template __host__ __device__ pointer ::pointer(decltype(nullptr)) : super_t(static_cast(nullptr)) {} // end pointer::pointer -#endif template @@ -82,7 +74,6 @@ template {} // end pointer::pointer -#if THRUST_CPP_DIALECT >= 2011 template __host__ __device__ typename pointer::derived_type & @@ -92,7 +83,6 @@ template super_t::base_reference() = nullptr; return static_cast(*this); } // end pointer::operator= -#endif template @@ -159,7 +149,15 @@ template } // end pointer::get -#if THRUST_CPP_DIALECT >= 2011 +template + __host__ __device__ + Element *pointer + ::operator->() const +{ + return super_t::base(); +} // end pointer::operator-> + + template __host__ __device__ pointer @@ -167,7 +165,6 @@ template { return bool(get()); } // end pointer::operator bool -#endif template &os, return os << p.get(); } -#if THRUST_CPP_DIALECT >= 2011 // NOTE: These are needed so that Thrust smart pointers work with // `std::unique_ptr`. template @@ -209,65 +205,6 @@ bool operator!=(pointer p, decltype(nullptr)) { return !(nullptr == p); } -#endif - -namespace detail -{ - -#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 > -{ - 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 > -{ - 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 - - - -// XXX WAR g++-4.1 problem with correctly implementing -// rebind_pointer for pointer by specializing it here -template - struct rebind_pointer, NewElement> -{ - // XXX note we don't attempt to rebind the pointer's Reference type (or Derived) - typedef thrust::pointer type; -}; - -template - struct rebind_pointer, NewElement> - : rebind_pointer, NewElement> -{}; - -template - struct rebind_pointer, NewElement> - : rebind_pointer, NewElement> -{}; -#endif - -} // end namespace detail - -} // end thrust +} // namespace thrust diff --git a/thrust/detail/reference.h b/thrust/detail/reference.h index 89bcf63ca7..100bfc6b6b 100644 --- a/thrust/detail/reference.h +++ b/thrust/detail/reference.h @@ -17,162 +17,519 @@ #pragma once #include -#include -#include #include +#include +#include +#include +#include +#include +#include +#include +#include #include - namespace thrust { + namespace detail { - -template struct is_wrapped_reference; - +template +struct is_wrapped_reference; } -// the base type for all of thrust's system-annotated references. -// for reasonable reference-like semantics, derived types must reimplement the following: -// 1. constructor from pointer -// 2. copy constructor -// 3. templated copy constructor from other reference -// 4. templated assignment from other reference -// 5. assignment from value_type -template - class reference +/*! \p reference acts as a reference-like wrapper for an object residing in + * memory that a \p pointer refers to. + */ +template +class reference { - private: - typedef typename thrust::detail::eval_if< - thrust::detail::is_same::value, - thrust::detail::identity_, - thrust::detail::identity_ - >::type derived_type; - - // hint for is_wrapped_reference lets it know that this type (or a derived type) - // is a wrapped reference - struct wrapped_reference_hint {}; - template friend struct thrust::detail::is_wrapped_reference; - - public: - typedef Pointer pointer; - typedef typename thrust::detail::remove_const::type value_type; - - __host__ __device__ - explicit reference(const pointer &ptr); - -#if THRUST_CPP_DIALECT >= 2011 - reference(const reference &) = default; -#endif - - template - __host__ __device__ - reference(const reference &other, - typename thrust::detail::enable_if_convertible< - typename reference::pointer, - pointer - >::type * = 0); - - __host__ __device__ - derived_type &operator=(const reference &other); - - // XXX this may need an enable_if - template - __host__ __device__ - derived_type &operator=(const reference &other); - - __host__ __device__ - derived_type &operator=(const value_type &x); - - __host__ __device__ - pointer operator&() const; - - __host__ __device__ - operator value_type () const; - - __host__ __device__ - void swap(derived_type &other); - - derived_type &operator++(); - - value_type operator++(int); - - // XXX parameterize the type of rhs - derived_type &operator+=(const value_type &rhs); - - derived_type &operator--(); - - value_type operator--(int); - - // XXX parameterize the type of rhs - derived_type &operator-=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator*=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator/=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator%=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator<<=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator>>=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator&=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator|=(const value_type &rhs); - - // XXX parameterize the type of rhs - derived_type &operator^=(const value_type &rhs); - - private: - const pointer m_ptr; - - // allow access to m_ptr for other references - template friend class reference; - - template - __host__ __device__ - inline value_type strip_const_get_value(const System &system) const; - - template - __host__ __device__ - inline void assign_from(OtherPointer src); - - // XXX this helper exists only to avoid warnings about null references from the other assign_from - template - inline __host__ __device__ - void assign_from(System1 *system1, System2 *system2, OtherPointer src); - - template - __host__ __device__ - inline void strip_const_assign_value(const System &system, OtherPointer src); - - // XXX this helper exists only to avoid warnings about null references from the other swap - template - inline __host__ __device__ - void swap(System *system, derived_type &other); - - // XXX this helper exists only to avoid warnings about null references from operator value_type () - template - inline __host__ __device__ - value_type convert_to_value_type(System *system) const; -}; // end reference +private: + using derived_type = typename std::conditional< + std::is_same::value, reference, Derived + >::type; + +public: + using pointer = Pointer; + using value_type = typename thrust::remove_cvref::type; + + reference(reference const&) = default; + + reference(reference&&) = default; + + /*! Construct a \p reference from another \p reference of a 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 OtherElement The element type of the other \p reference. + * \tparam OtherPointer The pointer type of the other \p reference. + * \tparam OtherDerived The derived type of the other \p reference. + */ + template + __host__ __device__ + reference( + reference const& other + , typename std::enable_if< + std::is_convertible< + typename reference::pointer + , pointer + >::value + >::type* = nullptr + ) + : ptr(other.ptr) + {} + + /*! Construct a \p reference that refers 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 construct from. + */ + __host__ __device__ + explicit reference(pointer const& p) : ptr(p) {} + + /*! Assign the object referred to \p other to the object referred to by + * this \p reference. + * + * \param other The other \p reference to assign from. + * + * \return *this. + */ + __host__ __device__ + derived_type& operator=(reference const& other) + { + assign_from(&other); + return derived(); + } + + /*! Assign the object referred to by this \p reference with the object + * referred to by another \p reference of related type. + * + * \param other The other \p reference to assign from. + * \tparam OtherElement The element type of the other \p reference. + * \tparam OtherPointer The pointer type of the other \p reference. + * \tparam OtherDerived The derived type of the other \p reference. + * + * \return *this. + */ + template + __host__ __device__ + typename std::enable_if< + std::is_convertible< + typename reference::pointer + , pointer + >::value + , derived_type& + >::type + operator=(reference const& other) + { + assign_from(&other); + return derived(); + } + + /*! Assign \p rhs to the object referred to by this \p tagged_reference. + * + * \param rhs The \p value_type to assign from. + * + * \return *this. + */ + __host__ __device__ + derived_type& operator=(value_type const& rhs) + { + assign_from(&rhs); + return derived(); + } + + /*! Assign \p rhs to the object referred to by this \p reference. + * + * \param rhs The \p value_type to assign from. + * + * \return *this. + */ + __host__ __device__ + derived_type& operator=(value_type&& rhs) + { + assign_from(std::move(rhs)); + return derived(); + } + + /*! Exchanges the value of the object referred to by this \p tagged_reference + * with the object referred to by \p other. + * + * \param other The \p tagged_reference to swap with. + */ + __host__ __device__ + void swap(derived_type& other) + { + // Avoid default-constructing a system; instead, just use a null pointer + // for dispatch. This assumes that `get_value` will not access any system + // state. + typename thrust::iterator_system::type* system = nullptr; + swap(system, other); + } + + __host__ __device__ pointer operator&() const { return ptr; } + + // This is inherently hazardous, as it discards the strong type information + // about what system the object is on. + __host__ __device__ operator value_type() const + { + // Avoid default-constructing a system; instead, just use a null pointer + // for dispatch. This assumes that `get_value` will not access any system + // state. + typename thrust::iterator_system::type* system = nullptr; + return convert_to_value_type(system); + } + + __host__ __device__ + derived_type& operator++() + { + // Sadly, this has to make a copy. The only mechanism we have for + // modifying the value, which may be in memory inaccessible to this + // system, is to get a copy of it, modify the copy, and then update it. + value_type tmp = *this; + ++tmp; + *this = tmp; + return derived(); + } + + __host__ __device__ + value_type operator++(int) + { + value_type tmp = *this; + value_type result = tmp++; + *this = std::move(tmp); + return result; + } + + derived_type& operator--() + { + // Sadly, this has to make a copy. The only mechanism we have for + // modifying the value, which may be in memory inaccessible to this + // system, is to get a copy of it, modify the copy, and then update it. + value_type tmp = *this; + --tmp; + *this = std::move(tmp); + return derived(); + } + + value_type operator--(int) + { + value_type tmp = *this; + value_type result = tmp--; + *this = std::move(tmp); + return derived(); + } + + __host__ __device__ + derived_type& operator+=(value_type const& rhs) + { + value_type tmp = *this; + tmp += rhs; + *this = tmp; + return derived(); + } + + derived_type& operator-=(value_type const& rhs) + { + value_type tmp = *this; + tmp -= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator*=(value_type const& rhs) + { + value_type tmp = *this; + tmp *= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator/=(value_type const& rhs) + { + value_type tmp = *this; + tmp /= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator%=(value_type const& rhs) + { + value_type tmp = *this; + tmp %= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator<<=(value_type const& rhs) + { + value_type tmp = *this; + tmp <<= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator>>=(value_type const& rhs) + { + value_type tmp = *this; + tmp >>= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator&=(value_type const& rhs) + { + value_type tmp = *this; + tmp &= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator|=(value_type const& rhs) + { + value_type tmp = *this; + tmp |= rhs; + *this = tmp; + return derived(); + } + + derived_type& operator^=(value_type const& rhs) + { + value_type tmp = *this; + tmp ^= rhs; + *this = tmp; + return derived(); + } + +private: + pointer const ptr; + + // `thrust::detail::is_wrapped_reference` is a trait that indicates whether + // a type is a fancy reference. It detects such types by loooking for a + // nested `wrapped_reference_hint` type. + struct wrapped_reference_hint {}; + template + friend struct thrust::detail::is_wrapped_reference; + + template + friend class reference; + + derived_type& derived() { return static_cast(*this); } + + template + __host__ __device__ + value_type convert_to_value_type(System* system) const + { + using thrust::system::detail::generic::select_system; + return strip_const_get_value(select_system(*system)); + } + + template + __host__ __device__ + value_type strip_const_get_value(System const& system) const + { + System &non_const_system = const_cast(system); + + using thrust::system::detail::generic::get_value; + return get_value(thrust::detail::derived_cast(non_const_system), ptr); + } + + template + __host__ __device__ + void assign_from(OtherPointer src) + { + // Avoid default-constructing systems; instead, just use a null pointer + // for dispatch. This assumes that `get_value` will not access any system + // state. + typename thrust::iterator_system::type* system0 = nullptr; + typename thrust::iterator_system::type* system1 = nullptr; + assign_from(system0, system1, src); + } + + template + __host__ __device__ + void assign_from(System1* system0, System1* system1, OtherPointer src) + { + using thrust::system::detail::generic::select_system; + strip_const_assign_value(select_system(*system0, *system1), src); + } + + template + __host__ __device__ + void strip_const_assign_value(System const& system, OtherPointer src) + { + System& non_const_system = const_cast(system); + + using thrust::system::detail::generic::assign_value; + assign_value(thrust::detail::derived_cast(non_const_system), ptr, src); + } + + template + __host__ __device__ + void swap(System* system, derived_type& other) + { + using thrust::system::detail::generic::select_system; + using thrust::system::detail::generic::iter_swap; + + iter_swap(select_system(*system, *system), ptr, other.ptr); + } +}; + +template +class reference {}; + +template +class reference {}; + +template < + typename Element, typename Pointer, typename Derived +, typename CharT, typename Traits +> +std::basic_ostream& operator<<( + std::basic_ostream&os +, reference const& r +) { + using value_type = typename reference::value_type; + return os << static_cast(r); +} -// Output stream operator -template -std::basic_ostream & -operator<<(std::basic_ostream &os, - const reference &y); +template +class tagged_reference; -} // end thrust +template +class tagged_reference + : public thrust::reference< + Element + , thrust::pointer> + , tagged_reference + > +{ +private: + using base_type = thrust::reference< + Element + , thrust::pointer> + , tagged_reference + >; + +public: + using value_type = typename base_type::value_type; + using pointer = typename base_type::pointer; + + tagged_reference(tagged_reference const&) = default; + + tagged_reference(tagged_reference&&) = default; + + /*! Construct a \p tagged_reference from another \p tagged_reference of a + * 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 OtherElement The element type of the other \p tagged_reference. + * \tparam OtherTag The tag type of the other \p tagged_reference. + */ + template + __host__ __device__ + tagged_reference( + tagged_reference const& other + , typename std::enable_if< + std::is_convertible< + typename tagged_reference::pointer + , pointer + >::value + >::type * = nullptr + ) + : base_type(other) + {} + + /*! Construct a \p tagged_reference that refers 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 construct from. + */ + __host__ __device__ explicit tagged_reference(pointer const& p) + : base_type(p) + {} + + /*! Assign the object referred to \p other to the object referred to by + * this \p tagged_reference. + * + * \param other The other \p tagged_reference to assign from. + * + * \return *this. + */ + __host__ __device__ + tagged_reference& operator=(tagged_reference const& other) + { + return base_type::operator=(other); + } + + /*! Assign the object referred to by this \p tagged_reference with the object + * referred to by another \p tagged_reference of related type. + * + * \param other The other \p tagged_reference to assign from. + * \tparam OtherElement The element type of the other \p tagged_reference. + * \tparam OtherTag The tag type of the other \p tagged_reference. + * + * \return *this. + */ + template + __host__ __device__ + typename std::enable_if< + std::is_convertible< + typename tagged_reference::pointer + , pointer + >::value + , tagged_reference& + >::type + operator=(tagged_reference const& other) + { + return base_type::operator=(other); + } + + /*! Assign \p rhs to the object referred to by this \p tagged_reference. + * + * \param rhs The \p value_type to assign from. + * + * \return *this. + */ + __host__ __device__ + tagged_reference& operator=(value_type const& rhs) + { + return base_type::operator=(rhs); + } + + /*! Assign \p rhs to the object referred to by this \p tagged_reference. + * + * \param rhs The \p value_type to assign from. + * + * \return *this. + */ + __host__ __device__ + tagged_reference& operator=(value_type&& rhs) + { + return base_type::operator=(std::move(rhs)); + } +}; + +template +class tagged_reference {}; + +template +class tagged_reference {}; + +/*! Exchanges the values of two objects referred to by \p tagged_reference. + * + * \param x The first \p tagged_reference of interest. + * \param y The second \p tagged_reference of interest. + */ +template +__host__ __device__ +void swap(tagged_reference& x, tagged_reference& y) +{ + x.swap(y); +} -#include +} // namespace thrust diff --git a/thrust/detail/reference.inl b/thrust/detail/reference.inl deleted file mode 100644 index 91f2b97368..0000000000 --- a/thrust/detail/reference.inl +++ /dev/null @@ -1,382 +0,0 @@ -/* - * Copyright 2008-2013 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. - */ - -#include -#include -#include -#include -#include -#include -#include -#include -#include - - -namespace thrust -{ - - -template - template - __host__ __device__ - reference - ::reference(const reference &other, - typename thrust::detail::enable_if_convertible< - typename reference::pointer, - pointer - >::type *) - : m_ptr(other.m_ptr) -{} - - -template - __host__ __device__ - reference - ::reference(const pointer &ptr) - : m_ptr(ptr) -{} - - -template - __host__ __device__ - typename reference::pointer - reference - ::operator&() const -{ - return m_ptr; -} // end reference::operator&() - - -template - __host__ __device__ - typename reference::derived_type & - reference - ::operator=(const value_type &v) -{ - assign_from(&v); - return static_cast(*this); -} // end reference::operator=() - - -template - __host__ __device__ - typename reference::derived_type & - reference - ::operator=(const reference &other) -{ - assign_from(&other); - return static_cast(*this); -} // end reference::operator=() - - -template - template - __host__ __device__ - typename reference::derived_type & - reference - ::operator=(const reference &other) -{ - assign_from(&other); - return static_cast(*this); -} // end reference::operator=() - - -template - template - __host__ __device__ - typename reference::value_type - reference - ::convert_to_value_type(System *system) const -{ - using thrust::system::detail::generic::select_system; - return strip_const_get_value(select_system(*system)); -} // end convert_to_value_type() - - -template - __host__ __device__ - reference - ::operator typename reference::value_type () const -{ - typedef typename thrust::iterator_system::type System; - - // XXX avoid default-constructing a system - // XXX use null a reference for dispatching - // XXX this assumes that the eventual invocation of - // XXX get_value will not access system state - System *system = 0; - - return convert_to_value_type(system); -} // end reference::operator value_type () - - -template - template - __host__ __device__ - typename reference::value_type - reference - ::strip_const_get_value(const System &system) const -{ - System &non_const_system = const_cast(system); - - using thrust::system::detail::generic::get_value; - - return get_value(thrust::detail::derived_cast(non_const_system), m_ptr); -} // end reference::strip_const_get_value() - - -template - template - __host__ __device__ - void reference - ::assign_from(System1 *system1, System2 *system2, OtherPointer src) -{ - using thrust::system::detail::generic::select_system; - - strip_const_assign_value(select_system(*system1, *system2), src); -} // end assign_from() - - -template - template - __host__ __device__ - void reference - ::assign_from(OtherPointer src) -{ - typedef typename thrust::iterator_system::type System1; - typedef typename thrust::iterator_system::type System2; - - // XXX avoid default-constructing a system - // XXX use null references for dispatching - // XXX this assumes that the eventual invocation of - // XXX assign_value will not access system state - System1 *system1 = 0; - System2 *system2 = 0; - - assign_from(system1, system2, src); -} // end assign_from() - - -template - template - __host__ __device__ - void reference - ::strip_const_assign_value(const System &system, OtherPointer src) -{ - System &non_const_system = const_cast(system); - - using thrust::system::detail::generic::assign_value; - - assign_value(thrust::detail::derived_cast(non_const_system), m_ptr, src); -} // end strip_const_assign_value() - - -template - template - __host__ __device__ - void reference - ::swap(System *system, derived_type &other) -{ - using thrust::system::detail::generic::select_system; - using thrust::system::detail::generic::iter_swap; - - iter_swap(select_system(*system, *system), m_ptr, other.m_ptr); -} // end reference::swap() - - -template - __host__ __device__ - void reference - ::swap(derived_type &other) -{ - typedef typename thrust::iterator_system::type System; - - // XXX avoid default-constructing a system - // XXX use null references for dispatching - // XXX this assumes that the eventual invocation - // XXX of iter_swap will not access system state - System *system = 0; - - swap(system, other); -} // end reference::swap() - - -template - typename reference::derived_type & - reference - ::operator++(void) -{ - value_type temp = *this; - ++temp; - *this = temp; - return static_cast(*this); -} // end reference::operator++() - - -template - typename reference::value_type - reference - ::operator++(int) -{ - value_type temp = *this; - value_type result = temp++; - *this = temp; - return result; -} // end reference::operator++() - - -template - typename reference::derived_type & - reference - ::operator+=(const value_type &rhs) -{ - value_type temp = *this; - temp += rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator+=() - -template - typename reference::derived_type & - reference - ::operator--(void) -{ - value_type temp = *this; - --temp; - *this = temp; - return static_cast(*this); -} // end reference::operator--() - -template - typename reference::value_type - reference - ::operator--(int) -{ - value_type temp = *this; - value_type result = temp--; - *this = temp; - return result; -} // end reference::operator--() - -template - typename reference::derived_type & - reference - ::operator-=(const value_type &rhs) -{ - value_type temp = *this; - temp -= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator-=() - -template - typename reference::derived_type & - reference - ::operator*=(const value_type &rhs) -{ - value_type temp = *this; - temp *= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator*=() - -template - typename reference::derived_type & - reference - ::operator/=(const value_type &rhs) -{ - value_type temp = *this; - temp /= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator/=() - -template - typename reference::derived_type & - reference - ::operator%=(const value_type &rhs) -{ - value_type temp = *this; - temp %= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator%=() - -template - typename reference::derived_type & - reference - ::operator<<=(const value_type &rhs) -{ - value_type temp = *this; - temp <<= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator<<=() - -template - typename reference::derived_type & - reference - ::operator>>=(const value_type &rhs) -{ - value_type temp = *this; - temp >>= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator>>=() - -template - typename reference::derived_type & - reference - ::operator&=(const value_type &rhs) -{ - value_type temp = *this; - temp &= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator&=() - -template - typename reference::derived_type & - reference - ::operator|=(const value_type &rhs) -{ - value_type temp = *this; - temp |= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator|=() - -template - typename reference::derived_type & - reference - ::operator^=(const value_type &rhs) -{ - value_type temp = *this; - temp ^= rhs; - *this = temp; - return static_cast(*this); -} // end reference::operator^=() - -template -std::basic_ostream & -operator<<(std::basic_ostream &os, - const reference &y) { - typedef typename reference::value_type value_type; - return os << static_cast(y); -} // end operator<<() - -} // end thrust diff --git a/thrust/detail/reference_forward_declaration.h b/thrust/detail/reference_forward_declaration.h index a8912ca43a..aa0168e535 100644 --- a/thrust/detail/reference_forward_declaration.h +++ b/thrust/detail/reference_forward_declaration.h @@ -1,5 +1,5 @@ /* - * Copyright 2008-2013 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. @@ -22,7 +22,8 @@ namespace thrust { -template class reference; +template +class reference; -} // end thrust +} // namespace thrust diff --git a/thrust/detail/type_traits/pointer_traits.h b/thrust/detail/type_traits/pointer_traits.h index 48ac7d6dc4..15f5851274 100644 --- a/thrust/detail/type_traits/pointer_traits.h +++ b/thrust/detail/type_traits/pointer_traits.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. @@ -22,6 +22,7 @@ #include #include #include +#include namespace thrust { @@ -83,34 +84,58 @@ template struct rebind_pointer; template struct rebind_pointer { - typedef U* type; + using type = U*; }; -template class Ptr, typename Arg, typename T> - struct rebind_pointer,T> +// Rebind generic fancy pointers. +template class Ptr, typename OldT, typename... Tail, typename T> + struct rebind_pointer,T> { - typedef Ptr type; + using type = Ptr; }; -template class Ptr, typename Arg1, typename Arg2, typename T> - struct rebind_pointer,T> +// Rebind `thrust::pointer`-like things with `thrust::reference`-like references. +template class Ptr, typename OldT, typename Tag, + template class Ref, typename... RefTail, + typename... PtrTail, typename T> + struct rebind_pointer,PtrTail...>,T> { - typedef Ptr type; +// static_assert(std::is_same::value, "0"); + using type = Ptr,PtrTail...>; }; -template class Ptr, typename Arg1, typename Arg2, typename Arg3, typename T> - struct rebind_pointer,T> +// Rebind `thrust::pointer`-like things with `thrust::reference`-like references +// and templated derived types. +template class Ptr, typename OldT, typename Tag, + template class Ref, typename... RefTail, + template class DerivedPtr, typename... DerivedPtrTail, + typename T> + struct rebind_pointer,DerivedPtr>,T> { - typedef Ptr type; +// static_assert(std::is_same::value, "1"); + using type = Ptr,DerivedPtr>; }; -template class Ptr, typename Arg1, typename Arg2, typename Arg3, typename Arg4, typename T> - struct rebind_pointer,T> +// Rebind `thrust::pointer`-like things with native reference types. +template class Ptr, typename OldT, typename Tag, + typename... PtrTail, typename T> + struct rebind_pointer::type,PtrTail...>,T> { - typedef Ptr type; +// static_assert(std::is_same::value, "2"); + using type = Ptr::type,PtrTail...>; +}; + +// Rebind `thrust::pointer`-like things with native reference types and templated +// derived types. +template class Ptr, typename OldT, typename Tag, + template class DerivedPtr, typename... DerivedPtrTail, + typename T> + struct rebind_pointer::type,DerivedPtr>,T> +{ +// static_assert(std::is_same::value, "3"); + using type = Ptr::type,DerivedPtr>; }; -// XXX this should probably be renamed native_type or similar __THRUST_DEFINE_HAS_NESTED_TYPE(has_raw_pointer, raw_pointer) namespace pointer_traits_detail @@ -179,7 +204,7 @@ template typedef typename pointer_difference::type difference_type; template - struct rebind + struct rebind { typedef typename rebind_pointer::type other; }; @@ -189,7 +214,7 @@ template { // XXX this is supposed to be pointer::pointer_to(&r); (i.e., call a static member function of pointer called pointer_to) // assume that pointer has a constructor from raw pointer instead - + return pointer(&r); } diff --git a/thrust/device_allocator.h b/thrust/device_allocator.h index f5ff0d9654..7b8100fe0b 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 @@ -24,7 +25,7 @@ #include #include #include -#include +#include #include #include @@ -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_reference.h b/thrust/device_reference.h index 6d8538b2fb..9a5b722799 100644 --- a/thrust/device_reference.h +++ b/thrust/device_reference.h @@ -38,7 +38,7 @@ namespace thrust * \p device_reference is not intended to be used directly; rather, this type * is the result of deferencing a \p device_ptr. Similarly, taking the address of * a \p device_reference yields a \p device_ptr. - * + * * \p device_reference may often be used from host code in place of operations defined on * its associated \c value_type. For example, when \p device_reference refers to an * arithmetic type, arithmetic operations on it are legal: @@ -158,7 +158,7 @@ namespace thrust * return 0; * } * \endcode - * + * * Another common case where a \p device_reference cannot directly be used in place of * its referent object occurs when passing them as parameters to functions like \c printf * which have varargs parameters. Because varargs parameters must be Plain Old Data, a @@ -209,7 +209,7 @@ template /*! This copy constructor accepts a const reference to another * \p device_reference. After this \p device_reference is constructed, * it shall refer to the same object as \p other. - * + * * \param other A \p device_reference to copy from. * * The following code snippet demonstrates the semantics of this @@ -233,7 +233,7 @@ template * assert(ref == 13); * \endcode * - * \note This constructor is templated primarily to allow initialization of + * \note This constructor is templated primarily to allow initialization of * device_reference from device_reference. */ template @@ -293,7 +293,7 @@ template /*! Assignment operator assigns the value of the given value to the * value referenced by this \p device_reference. - * + * * \param x The value to assign from. * \return *this */ @@ -332,7 +332,7 @@ template * \p device_reference. * * \return *this - * + * * The following code snippet demonstrates the semantics of * \p device_reference's prefix increment operator. * @@ -467,7 +467,7 @@ template * \p device_reference. * * \return *this - * + * * The following code snippet demonstrates the semantics of * \p device_reference's prefix decrement operator. * @@ -979,5 +979,3 @@ operator<<(std::basic_ostream &os, const device_reference &y); } // end thrust -#include - 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/mr/allocator.h b/thrust/mr/allocator.h index 4c6c328860..e51d46e63e 100644 --- a/thrust/mr/allocator.h +++ b/thrust/mr/allocator.h @@ -23,9 +23,9 @@ #include #include +#include #include -#include #include #include diff --git a/thrust/memory/detail/device_system_resource.h b/thrust/mr/device_memory_resource.h similarity index 96% rename from thrust/memory/detail/device_system_resource.h rename to thrust/mr/device_memory_resource.h index 9e94991d61..2230843094 100644 --- a/thrust/memory/detail/device_system_resource.h +++ b/thrust/mr/device_memory_resource.h @@ -1,5 +1,5 @@ /* - * Copyright 2018 NVIDIA Corporation + * Copyright 2018-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. diff --git a/thrust/memory/detail/host_system_resource.h b/thrust/mr/host_memory_resource.h similarity index 95% rename from thrust/memory/detail/host_system_resource.h rename to thrust/mr/host_memory_resource.h index ded1c4d0bf..755c1b3197 100644 --- a/thrust/memory/detail/host_system_resource.h +++ b/thrust/mr/host_memory_resource.h @@ -1,5 +1,5 @@ /* - * Copyright 2018 NVIDIA Corporation + * Copyright 2018-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. diff --git a/thrust/mr/memory_resource.h b/thrust/mr/memory_resource.h index 048ca24059..ea958f5fab 100644 --- a/thrust/mr/memory_resource.h +++ b/thrust/mr/memory_resource.h @@ -21,7 +21,7 @@ #pragma once -#include "detail/config.h" +#include #ifdef THRUST_MR_STD_MR_HEADER # include THRUST_MR_STD_MR_HEADER #endif diff --git a/thrust/mr/polymorphic_adaptor.h b/thrust/mr/polymorphic_adaptor.h index d5d98bf838..67c581a068 100644 --- a/thrust/mr/polymorphic_adaptor.h +++ b/thrust/mr/polymorphic_adaptor.h @@ -16,7 +16,7 @@ #pragma once -#include "memory_resource.h" +#include namespace thrust { diff --git a/thrust/mr/pool_options.h b/thrust/mr/pool_options.h index 60430b7d27..7994e914ab 100644 --- a/thrust/mr/pool_options.h +++ b/thrust/mr/pool_options.h @@ -24,7 +24,7 @@ #include -#include +#include namespace thrust { diff --git a/thrust/mr/universal_memory_resource.h b/thrust/mr/universal_memory_resource.h new file mode 100644 index 0000000000..b7f1ebd6f4 --- /dev/null +++ b/thrust/mr/universal_memory_resource.h @@ -0,0 +1,22 @@ +/* + * Copyright 2018-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 + diff --git a/thrust/mr/validator.h b/thrust/mr/validator.h index 9376ae870b..8f8676d11f 100644 --- a/thrust/mr/validator.h +++ b/thrust/mr/validator.h @@ -16,8 +16,8 @@ #pragma once -#include "detail/config.h" -#include "memory_resource.h" +#include +#include namespace thrust { 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..376b8f4f5a 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/async/customization.h b/thrust/system/cuda/detail/async/customization.h index eb52c2cf02..aead7b12bb 100644 --- a/thrust/system/cuda/detail/async/customization.h +++ b/thrust/system/cuda/detail/async/customization.h @@ -42,7 +42,7 @@ #include #include #include -#include +#include #include #include #include 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..26993f0bf8 100644 --- a/thrust/system/cuda/memory_resource.h +++ b/thrust/system/cuda/memory_resource.h @@ -22,13 +22,12 @@ #include #include -#include #include #include #include #include -#include +#include namespace thrust { @@ -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/type_traits/remove_cvref.h b/thrust/type_traits/remove_cvref.h index 4079bfe8e8..fe407f9e22 100644 --- a/thrust/type_traits/remove_cvref.h +++ b/thrust/type_traits/remove_cvref.h @@ -17,7 +17,7 @@ #pragma once #include -#include +#include namespace thrust { @@ -32,9 +32,9 @@ using std::remove_cvref_t; template struct remove_cvref { - typedef typename detail::remove_cv< - typename detail::remove_reference::type - >::type type; + using type = typename std::remove_cv< + typename std::remove_reference::type + >::type; }; #if THRUST_CPP_DIALECT >= 2011 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 +