Miscellaneous warning cleanup.

This patch gets a clean build using recent internal nvcc and
gcc 7.4.0, with the exception of Github NVIDIA#1049.

We were adding `dependencies/cub/cub` to the include line, but
including headers as `#include <cub/xxx.cuh>`. This caused builds to
pull from the system cub headers installed with the toolkit.

Templated utilites `has_member_construct1` and `contiguous_storage`
were calling `__host__` only methods from `__host__ __device__` members.

This was also calling some tricky to find "host from h/d" warnings.

Author: alliepiper

CMakeLists.txt

@@ -491,7 +491,7 @@ foreach (THRUST_TEST_SOURCE IN LISTS THRUST_TESTS)
 
   target_include_directories(
     ${THRUST_TEST}
-    PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub/cub
+    PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub
     PRIVATE ${PROJECT_SOURCE_DIR}/testing
   )
 
@@ -518,7 +518,7 @@ foreach (THRUST_TEST_SOURCE IN LISTS THRUST_TESTS)
 
     target_include_directories(
       ${THRUST_TEST_RDC}
-      PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub/cub
+      PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub
       PRIVATE ${PROJECT_SOURCE_DIR}/testing
     )
 
@@ -617,7 +617,7 @@ foreach (THRUST_EXAMPLE_SOURCE IN LISTS THRUST_EXAMPLES)
 
   target_include_directories(
     ${THRUST_EXAMPLE}
-    PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub/cub
+    PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub
     PRIVATE ${PROJECT_SOURCE_DIR}/examples
   )
 
@@ -640,7 +640,7 @@ foreach (THRUST_EXAMPLE_SOURCE IN LISTS THRUST_EXAMPLES)
 
     target_include_directories(
       ${THRUST_EXAMPLE_RDC}
-      PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub/cub
+      PUBLIC ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/dependencies/cub
       PRIVATE ${PROJECT_SOURCE_DIR}/examples
     )
 

dependencies/cub

@@ -87,6 +87,7 @@ template<typename Alloc, typename T>
   a.construct(p);
 }
 
+__thrust_exec_check_disable__
 template<typename Alloc, typename T>
   inline __host__ __device__
     typename disable_if<

thrust/detail/allocator/allocator_traits.inl

@@ -186,6 +186,7 @@ __host__ __device__
   return m_begin[n];
 } // end contiguous_storage::operator[]()
 
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   typename contiguous_storage<T,Alloc>::allocator_type
@@ -340,6 +341,7 @@ __host__ __device__
   destroy_on_allocator_mismatch_dispatch(c, other, first, last);
 } // end contiguous_storage::destroy_on_allocator_mismatch
 
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   void contiguous_storage<T,Alloc>
@@ -448,6 +450,7 @@ __host__ __device__
   return false;
 } // end contiguous_storage::is_allocator_not_equal_dispatch()
 
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   bool contiguous_storage<T,Alloc>
@@ -456,6 +459,7 @@ __host__ __device__
   return m_allocator != other;
 } // end contiguous_storage::is_allocator_not_equal_dispatch()
 
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   void contiguous_storage<T,Alloc>
@@ -474,6 +478,7 @@ __host__ __device__
 {
 } // end contiguous_storage::deallocate_on_allocator_mismatch()
 
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   void contiguous_storage<T,Alloc>
@@ -494,6 +499,7 @@ __host__ __device__
 {
 } // end contiguous_storage::destroy_on_allocator_mismatch()
 
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   void contiguous_storage<T,Alloc>
@@ -510,6 +516,7 @@ __host__ __device__
 } // end contiguous_storage::propagate_allocator()
 
 #if __cplusplus >= 201103L
+__thrust_exec_check_disable__
 template<typename T, typename Alloc>
 __host__ __device__
   void contiguous_storage<T,Alloc>

thrust/detail/contiguous_storage.inl

@@ -68,29 +68,25 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
 
     /*! This constructor creates an empty \p device_vector.
      */
-    __host__
     device_vector(void)
       :Parent() {}
 
     /*! This constructor creates an empty \p device_vector.
      *  \param alloc The allocator to use by this device_vector.
      */
-    __host__
     device_vector(const Alloc &alloc)
       :Parent(alloc) {}
 
     /*! The destructor erases the elements.
      */
     //  Define an empty destructor to explicitly specify
     //  its execution space qualifier, as a workaround for nvcc warning
-    __host__
     ~device_vector(void) {}
 
     /*! This constructor creates a \p device_vector with the given
      *  size.
      *  \param n The number of elements to initially create.
      */
-    __host__
     explicit device_vector(size_type n)
       :Parent(n) {}
 
@@ -99,7 +95,6 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
      *  \param n The number of elements to initially create.
      *  \param alloc The allocator to use by this device_vector.
      */
-    __host__
     explicit device_vector(size_type n, const Alloc &alloc)
       :Parent(n,alloc) {}
 
@@ -108,7 +103,6 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
      *  \param n The number of elements to initially create.
      *  \param value An element to copy.
      */
-    __host__
     explicit device_vector(size_type n, const value_type &value)
       :Parent(n,value) {}
 
@@ -118,54 +112,47 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
      *  \param value An element to copy.
      *  \param alloc The allocator to use by this device_vector.
      */
-    __host__
     explicit device_vector(size_type n, const value_type &value, const Alloc &alloc)
       :Parent(n,value,alloc) {}
 
     /*! Copy constructor copies from an exemplar \p device_vector.
      *  \param v The \p device_vector to copy.
      */
-    __host__
     device_vector(const device_vector &v)
       :Parent(v) {}
 
     /*! Copy constructor copies from an exemplar \p device_vector.
      *  \param v The \p device_vector to copy.
      *  \param alloc The allocator to use by this device_vector.
      */
-    __host__
     device_vector(const device_vector &v, const Alloc &alloc)
       :Parent(v,alloc) {}
 
   #if THRUST_CPP_DIALECT >= 2011
     /*! Move constructor moves from another \p device_vector.
      *  \param v The device_vector to move.
      */
-    __host__
     device_vector(device_vector &&v)
       :Parent(std::move(v)) {}
 
     /*! Move constructor moves from another \p device_vector.
      *  \param v The device_vector to move.
      *  \param alloc The allocator to use by this device_vector.
      */
-    __host__
     device_vector(device_vector &&v, const Alloc &alloc)
       :Parent(std::move(v), alloc) {}
   #endif // THRUST_CPP_DIALECT >= 2011
 
     /*! Copy assign operator copies another \p device_vector with the same type.
      *  \param v The \p device_vector to copy.
      */
-    __host__
     device_vector &operator=(const device_vector &v)
     { Parent::operator=(v); return *this; }
 
   #if THRUST_CPP_DIALECT >= 2011
     /*! Move assign operator moves from another \p device_vector.
      *  \param v The device_vector to move.
      */
-     __host__
      device_vector &operator=(device_vector &&v)
      { Parent::operator=(std::move(v)); return *this; }
   #endif // THRUST_CPP_DIALECT >= 2011
@@ -174,47 +161,40 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
      *  \param v The \p device_vector to copy.
      */
     template<typename OtherT, typename OtherAlloc>
-    __host__ explicit
-    __device__
-    device_vector(const device_vector<OtherT,OtherAlloc> &v)
+    explicit device_vector(const device_vector<OtherT,OtherAlloc> &v)
       :Parent(v) {}
 
     /*! Assign operator copies from an exemplar \p device_vector with different type.
      *  \param v The \p device_vector to copy.
      */
     template<typename OtherT, typename OtherAlloc>
-    __host__
     device_vector &operator=(const device_vector<OtherT,OtherAlloc> &v)
     { Parent::operator=(v); return *this; }
 
     /*! Copy constructor copies from an exemplar \c std::vector.
      *  \param v The <tt>std::vector</tt> to copy.
      */
     template<typename OtherT, typename OtherAlloc>
-    __host__
     device_vector(const std::vector<OtherT,OtherAlloc> &v)
       :Parent(v) {}
 
     /*! Assign operator copies from an exemplar <tt>std::vector</tt>.
      *  \param v The <tt>std::vector</tt> to copy.
      */
     template<typename OtherT, typename OtherAlloc>
-    __host__
     device_vector &operator=(const std::vector<OtherT,OtherAlloc> &v)
     { Parent::operator=(v); return *this;}
 
     /*! Copy constructor copies from an exemplar \p host_vector with possibly different type.
      *  \param v The \p host_vector to copy.
      */
     template<typename OtherT, typename OtherAlloc>
-    __host__
     device_vector(const host_vector<OtherT,OtherAlloc> &v);
 
     /*! Assign operator copies from an examplar \p host_vector.
      *  \param v The \p host_vector to copy.
      */
     template<typename OtherT, typename OtherAlloc>
-    __host__
     device_vector &operator=(const host_vector<OtherT,OtherAlloc> &v)
     { Parent::operator=(v); return *this; }
 
@@ -223,7 +203,6 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
      *  \param last The end of the range.
      */
     template<typename InputIterator>
-    __host__
     device_vector(InputIterator first, InputIterator last)
       :Parent(first,last) {}
 
@@ -233,7 +212,6 @@ template<typename T, typename Alloc = thrust::device_allocator<T> >
      *  \param alloc The allocator to use by this device_vector.
      */
     template<typename InputIterator>
-    __host__
     device_vector(InputIterator first, InputIterator last, const Alloc &alloc)
       :Parent(first,last,alloc) {}