Changes necessary to support Feta.

Reviewed-by: Bryce Adelstein Lelbach aka wash <[email protected]>

thrust/detail/allocator/temporary_allocator.inl

@@ -20,7 +20,7 @@
 #include <thrust/system/detail/bad_alloc.h>
 #include <cassert>
 
-#if defined(__CUDA_ARCH__) && THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+#if (defined(__PGI_CUDA__) || defined(__CUDA_ARCH__)) && THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
 #include <thrust/system/cuda/detail/terminate.h>
 #endif
 
@@ -45,11 +45,15 @@ __host__ __device__
     // note that we pass cnt to deallocate, not a value derived from result.second
     deallocate(result.first, cnt);
 
-#if !defined(__CUDA_ARCH__)
-    throw thrust::system::detail::bad_alloc("temporary_buffer::allocate: get_temporary_buffer failed");
-#elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
-    thrust::system::cuda::detail::terminate_with_message("temporary_buffer::allocate: get_temporary_buffer failed");
-#endif
+    if (THRUST_IS_HOST_CODE) {
+      #if THRUST_INCLUDE_HOST_CODE
+        throw thrust::system::detail::bad_alloc("temporary_buffer::allocate: get_temporary_buffer failed");
+      #endif
+    } else {
+      #if THRUST_INCLUDE_DEVICE_CODE
+        thrust::system::cuda::detail::terminate_with_message("temporary_buffer::allocate: get_temporary_buffer failed");
+      #endif
+    }
   } // end if
 
   return result.first;

thrust/detail/config/cpp_compatibility.h

@@ -34,6 +34,10 @@
 #  define THRUST_DEFAULT = default;
 #  define THRUST_NOEXCEPT noexcept
 #  define THRUST_FINAL final
+// THRUST_STATIC_CONSTANT is a holdover from an earlier Thrust version, and is
+// here only because we are using a hybrid of Thrust versions.  Don't push this
+// back to Thrust.
+#  define THRUST_STATIC_CONSTANT static constexpr
 #else
 #  define THRUST_CONSTEXPR
 #  define THRUST_OVERRIDE
@@ -49,13 +53,14 @@
 // FIXME: Combine THRUST_INLINE_CONSTANT and
 // THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT into one macro when NVCC properly
 // supports `constexpr` globals in host and device code.
-#ifdef __CUDA_ARCH__
+// NVC++ uses the same definitions as NVCC does for device code.
+#if defined(__CUDA_ARCH__) || defined(__PGI_CUDA__)
 // FIXME: Add this when NVCC supports inline variables.
 //#  if   THRUST_CPP_DIALECT >= 2017
 //#    define THRUST_INLINE_CONSTANT                 inline constexpr
 //#    define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT inline constexpr
 #  if THRUST_CPP_DIALECT >= 2011
-#    define THRUST_INLINE_CONSTANT                 static constexpr
+#    define THRUST_INLINE_CONSTANT                 static const __device__
 #    define THRUST_INLINE_INTEGRAL_MEMBER_CONSTANT static constexpr
 #  else
 #    define THRUST_INLINE_CONSTANT                 static const __device__
@@ -75,3 +80,20 @@
 #  endif
 #endif
 
+#if defined(__PGI_CUDA__)
+#  define THRUST_IS_DEVICE_CODE __builtin_is_device_code()
+#  define THRUST_IS_HOST_CODE (!__builtin_is_device_code())
+#  define THRUST_INCLUDE_DEVICE_CODE 1
+#  define THRUST_INCLUDE_HOST_CODE 1
+#elif defined(__CUDA_ARCH__)
+#  define THRUST_IS_DEVICE_CODE 1
+#  define THRUST_IS_HOST_CODE 0
+#  define THRUST_INCLUDE_DEVICE_CODE 1
+#  define THRUST_INCLUDE_HOST_CODE 0
+#else
+#  define THRUST_IS_DEVICE_CODE 0
+#  define THRUST_IS_HOST_CODE 1
+#  define THRUST_INCLUDE_DEVICE_CODE 0
+#  define THRUST_INCLUDE_HOST_CODE 1
+#endif
+

thrust/detail/config/exec_check_disable.h

@@ -22,7 +22,11 @@
 
 #include <thrust/detail/config.h>
 
-#if defined(__CUDACC__) && !(defined(__CUDA__) && defined(__clang__))
+/* pragma nv_exec_check_disable only works with NVCC, not with PGI or Clang.
+   Having a macro expand to a #pragma (rather than _Pragma) only works with
+   NVCC's compilation model, not with other compilers. */
+#if defined(__CUDACC__) && !defined(__PGI_CUDA__) && \
+    !(defined(__CUDA__) && defined(__clang__))
 
 #define __thrust_exec_check_disable__ #pragma nv_exec_check_disable
 

thrust/detail/contiguous_storage.inl

@@ -428,15 +428,19 @@ __host__ __device__
   void contiguous_storage<T,Alloc>
     ::swap_allocators(false_type, Alloc &other)
 {
-#ifdef __CUDA_ARCH__
-  // allocators must be equal when swapping containers with allocators that propagate on swap
-  assert(!is_allocator_not_equal(other));
-#else
-  if (is_allocator_not_equal(other))
-  {
-    throw allocator_mismatch_on_swap();
+  if (THRUST_IS_DEVICE_CODE) {
+    #if THRUST_INCLUDE_DEVICE_CODE
+      // allocators must be equal when swapping containers with allocators that propagate on swap
+      assert(!is_allocator_not_equal(other));
+    #endif
+  } else {
+    #if THRUST_INCLUDE_HOST_CODE
+      if (is_allocator_not_equal(other))
+      {
+        throw allocator_mismatch_on_swap();
+      }
+    #endif
   }
-#endif
   thrust::swap(m_allocator, other);
 } // end contiguous_storage::swap_allocators()
 

thrust/detail/functional/actor.h

@@ -52,7 +52,7 @@ template<typename Eval>
   typedef Eval eval_type;
 
   __host__ __device__
-  actor(void);
+  THRUST_CONSTEXPR actor(void);
 
   __host__ __device__
   actor(const Eval &base);

thrust/detail/functional/actor.inl

@@ -38,7 +38,7 @@ namespace functional
 
 template<typename Eval>
   __host__ __device__
-  actor<Eval>
+  THRUST_CONSTEXPR actor<Eval>
     ::actor(void)
       : eval_type()
 {}

thrust/detail/functional/argument.h

@@ -59,7 +59,7 @@ template<unsigned int i>
     };
 
     __host__ __device__
-    argument(void){}
+    THRUST_CONSTEXPR argument(void){}
 
     template<typename Env>
     __host__ __device__

thrust/detail/integer_math.h

@@ -32,22 +32,27 @@ template <typename Integer>
 __host__ __device__ __thrust_forceinline__
 Integer clz(Integer x)
 {
-#if __CUDA_ARCH__
-  return ::__clz(x);
-#else
-  int num_bits = 8 * sizeof(Integer);
-  int num_bits_minus_one = num_bits - 1;
-
-  for (int i = num_bits_minus_one; i >= 0; --i)
-  {
-    if ((Integer(1) << i) & x)
-    {
-      return num_bits_minus_one - i;
-    }
+  Integer result;
+  if (THRUST_IS_DEVICE_CODE) {
+    #if THRUST_INCLUDE_DEVICE_CODE
+      result = ::__clz(x);
+    #endif
+  } else {
+    #if THRUST_INCLUDE_HOST_CODE
+      int num_bits = 8 * sizeof(Integer);
+      int num_bits_minus_one = num_bits - 1;
+      result = num_bits;
+      for (int i = num_bits_minus_one; i >= 0; --i)
+      {
+        if ((Integer(1) << i) & x)
+        {
+          result = num_bits_minus_one - i;
+          break;
+        }
+      }
+    #endif
   }
-
-  return num_bits;
-#endif
+  return result;
 }
 
 template <typename Integer>

thrust/detail/seq.h

@@ -31,7 +31,7 @@ struct seq_t : thrust::system::detail::sequential::execution_policy<seq_t>,
     thrust::system::detail::sequential::execution_policy>
 {
   __host__ __device__
-  seq_t() : thrust::system::detail::sequential::execution_policy<seq_t>() {}
+  THRUST_CONSTEXPR seq_t() : thrust::system::detail::sequential::execution_policy<seq_t>() {}
 
   // allow any execution_policy to convert to seq_t
   template<typename DerivedPolicy>
@@ -45,11 +45,7 @@ struct seq_t : thrust::system::detail::sequential::execution_policy<seq_t>,
 } // end detail
 
 
-#ifdef __CUDA_ARCH__
-static const __device__ detail::seq_t seq;
-#else
-static const detail::seq_t seq;
-#endif
+THRUST_INLINE_CONSTANT detail::seq_t seq;
 
 
 } // end thrust

thrust/execution_policy.h

@@ -27,6 +27,8 @@
 
 //! \cond
 
+//! \cond
+
 // #include the host system's execution_policy header
 #define __THRUST_HOST_SYSTEM_EXECUTION_POLICY_HEADER <__THRUST_HOST_SYSTEM_ROOT/execution_policy.h>
 #include __THRUST_HOST_SYSTEM_EXECUTION_POLICY_HEADER
@@ -39,6 +41,8 @@
 
 //! \endcond
 
+//! \endcond
+
 namespace thrust
 {
 
@@ -344,11 +348,7 @@ static const detail::host_t host;
  *  \see host_execution_policy
  *  \see thrust::device
  */
-#ifdef __CUDA_ARCH__
-static const __device__ detail::device_t device;
-#else
-static const detail::device_t device;
-#endif
+THRUST_INLINE_CONSTANT detail::device_t device;
 
 
 // define seq for the purpose of Doxygenating it

thrust/functional.h

@@ -1448,92 +1448,52 @@ namespace placeholders
 
 /*! \p thrust::placeholders::_1 is the placeholder for the first function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<0>::type _1;
-#else
-static const thrust::detail::functional::placeholder<0>::type _1;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<0>::type _1;
 
 
 /*! \p thrust::placeholders::_2 is the placeholder for the second function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<1>::type _2;
-#else
-static const thrust::detail::functional::placeholder<1>::type _2;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<1>::type _2;
 
 
 /*! \p thrust::placeholders::_3 is the placeholder for the third function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<2>::type _3;
-#else
-static const thrust::detail::functional::placeholder<2>::type _3;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<2>::type _3;
 
 
 /*! \p thrust::placeholders::_4 is the placeholder for the fourth function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<3>::type _4;
-#else
-static const thrust::detail::functional::placeholder<3>::type _4;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<3>::type _4;
 
 
 /*! \p thrust::placeholders::_5 is the placeholder for the fifth function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<4>::type _5;
-#else
-static const thrust::detail::functional::placeholder<4>::type _5;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<4>::type _5;
 
 
 /*! \p thrust::placeholders::_6 is the placeholder for the sixth function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<5>::type _6;
-#else
-static const thrust::detail::functional::placeholder<5>::type _6;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<5>::type _6;
 
 
 /*! \p thrust::placeholders::_7 is the placeholder for the seventh function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<6>::type _7;
-#else
-static const thrust::detail::functional::placeholder<6>::type _7;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<6>::type _7;
 
 
 /*! \p thrust::placeholders::_8 is the placeholder for the eighth function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<7>::type _8;
-#else
-static const thrust::detail::functional::placeholder<7>::type _8;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<7>::type _8;
 
 
 /*! \p thrust::placeholders::_9 is the placeholder for the ninth function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<8>::type _9;
-#else
-static const thrust::detail::functional::placeholder<8>::type _9;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<8>::type _9;
 
 
 /*! \p thrust::placeholders::_10 is the placeholder for the tenth function parameter.
  */
-#ifdef __CUDA_ARCH__
-static const __device__ thrust::detail::functional::placeholder<9>::type _10;
-#else
-static const thrust::detail::functional::placeholder<9>::type _10;
-#endif
+THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<9>::type _10;
 
 
 } // end placeholders

thrust/system/cpp/detail/par.h

@@ -35,14 +35,14 @@ struct par_t : thrust::system::cpp::detail::execution_policy<par_t>,
     thrust::system::cpp::detail::execution_policy>
 {
   __host__ __device__
-  par_t() : thrust::system::cpp::detail::execution_policy<par_t>() {}
+  THRUST_CONSTEXPR par_t() : thrust::system::cpp::detail::execution_policy<par_t>() {}
 };
 
 
 } // end detail
 
 
-static const detail::par_t par;
+THRUST_INLINE_CONSTANT detail::par_t par;
 
 
 } // end cpp