HIP backend: Adjust #ifdefs to support HIP backend
MODIFY
thrust/detail/allocator/no_throw_allocator.h
thrust/detail/allocator/no_throw_allocator.h
@@ -44,7 +44,7 @@ template<typename BaseAllocator>
__host__ __device__
void deallocate(typename super_t::pointer p, typename super_t::size_type n)
{
-#ifndef __CUDA_ARCH__
+#if !defined(__CUDA_ARCH__) && !defined(__HCC_ACCELERATOR__)
try
{
super_t::deallocate(p, n);

MODIFY
thrust/detail/allocator/temporary_allocator.inl
thrust/detail/allocator/temporary_allocator.inl
#include <thrust/system/cuda/detail/terminate.h>
#endif
+#ifdef __HCC__
+#include <thrust/system/hip/detail/terminate.h>
+#endif
+
namespace thrust
{
namespace detail
@@ -45,10 +49,12 @@ __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");
-#else
+#if defined(__CUDA_ARCH__)
thrust::system::cuda::detail::terminate_with_message("temporary_buffer::allocate: get_temporary_buffer failed");
+#elif defined(__HCC_ACCELERATOR__)
+ thrust::system::hip::detail::terminate_with_message("temporary_buffer::allocate: get_temporary_buffer failed");
+#else
+ throw thrust::system::detail::bad_alloc("temporary_buffer::allocate: get_temporary_buffer failed");
#endif
} // end if

MODIFY
thrust/detail/seq.h
thrust/detail/seq.h
@@ -51,7 +51,7 @@ struct seq_t : thrust::system::detail::sequential::execution_policy<seq_t>
} // end detail
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ detail::seq_t seq;
#else
static const detail::seq_t seq;

MODIFY
thrust/execution_policy.h
thrust/execution_policy.h
@@ -340,7 +340,7 @@ static const detail::host_t host;
* \see host_execution_policy
* \see thrust::device
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ detail::device_t device;
#else
static const detail::device_t device;

MODIFY
thrust/functional.h
thrust/functional.h
@@ -1377,7 +1377,7 @@ namespace placeholders
/*! \p thrust::placeholders::_1 is the placeholder for the first function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<0>::type _1;
#else
static const thrust::detail::functional::placeholder<0>::type _1;
@@ -1386,7 +1386,7 @@ static const thrust::detail::functional::placeholder<0>::type _1;
/*! \p thrust::placeholders::_2 is the placeholder for the second function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<1>::type _2;
#else
static const thrust::detail::functional::placeholder<1>::type _2;
@@ -1395,7 +1395,7 @@ static const thrust::detail::functional::placeholder<1>::type _2;
/*! \p thrust::placeholders::_3 is the placeholder for the third function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<2>::type _3;
#else
static const thrust::detail::functional::placeholder<2>::type _3;
@@ -1404,7 +1404,7 @@ static const thrust::detail::functional::placeholder<2>::type _3;
/*! \p thrust::placeholders::_4 is the placeholder for the fourth function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<3>::type _4;
#else
static const thrust::detail::functional::placeholder<3>::type _4;
@@ -1413,7 +1413,7 @@ static const thrust::detail::functional::placeholder<3>::type _4;
/*! \p thrust::placeholders::_5 is the placeholder for the fifth function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<4>::type _5;
#else
static const thrust::detail::functional::placeholder<4>::type _5;
@@ -1422,7 +1422,7 @@ static const thrust::detail::functional::placeholder<4>::type _5;
/*! \p thrust::placeholders::_6 is the placeholder for the sixth function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<5>::type _6;
#else
static const thrust::detail::functional::placeholder<5>::type _6;
@@ -1431,7 +1431,7 @@ static const thrust::detail::functional::placeholder<5>::type _6;
/*! \p thrust::placeholders::_7 is the placeholder for the seventh function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<6>::type _7;
#else
static const thrust::detail::functional::placeholder<6>::type _7;
@@ -1440,7 +1440,7 @@ static const thrust::detail::functional::placeholder<6>::type _7;
/*! \p thrust::placeholders::_8 is the placeholder for the eighth function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<7>::type _8;
#else
static const thrust::detail::functional::placeholder<7>::type _8;
@@ -1449,7 +1449,7 @@ static const thrust::detail::functional::placeholder<7>::type _8;
/*! \p thrust::placeholders::_9 is the placeholder for the ninth function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<8>::type _9;
#else
static const thrust::detail::functional::placeholder<8>::type _9;
@@ -1458,7 +1458,7 @@ static const thrust::detail::functional::placeholder<8>::type _9;
/*! \p thrust::placeholders::_10 is the placeholder for the tenth function parameter.
*/
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ thrust::detail::functional::placeholder<9>::type _10;
#else
static const thrust::detail::functional::placeholder<9>::type _10;

MODIFY
thrust/system/detail/sequential/execution_policy.h
thrust/system/detail/sequential/execution_policy.h
@@ -66,7 +66,7 @@ template<typename Derived>
};
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
static const __device__ tag seq;
#else
static const tag seq;

MODIFY
thrust/system/detail/sequential/sort.inl
thrust/system/detail/sequential/sort.inl
@@ -162,7 +162,7 @@ void stable_sort(sequential::execution_policy<DerivedPolicy> &exec,
{
// the compilation time of stable_primitive_sort is too expensive to use within a single CUDA thread
-#ifndef __CUDA_ARCH__
+#if !defined(__CUDA_ARCH__) && !defined(__HCC_ACCELERATOR__)
typedef typename thrust::iterator_traits<RandomAccessIterator>::value_type KeyType;
sort_detail::use_primitive_sort<KeyType,StrictWeakOrdering> use_primitive_sort;
#else
@@ -186,7 +186,7 @@ void stable_sort_by_key(sequential::execution_policy<DerivedPolicy> &exec,
{
// the compilation time of stable_primitive_sort_by_key is too expensive to use within a single CUDA thread
-#ifndef __CUDA_ARCH__
+#if !defined(__CUDA_ARCH__) && !defined(__HCC_ACCELERATOR__)
typedef typename thrust::iterator_traits<RandomAccessIterator1>::value_type KeyType;
sort_detail::use_primitive_sort<KeyType,StrictWeakOrdering> use_primitive_sort;
#else

MODIFY
thrust/system/detail/sequential/stable_merge_sort.inl
thrust/system/detail/sequential/stable_merge_sort.inl
@@ -354,7 +354,7 @@ void stable_merge_sort(sequential::execution_policy<DerivedPolicy> &exec,
StrictWeakOrdering comp)
{
// avoid recursion in CUDA threads
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
stable_merge_sort_detail::iterative_stable_merge_sort(exec, first, last, comp);
#else
stable_merge_sort_detail::recursive_stable_merge_sort(exec, first, last, comp);
@@ -374,7 +374,7 @@ void stable_merge_sort_by_key(sequential::execution_policy<DerivedPolicy> &exec,
StrictWeakOrdering comp)
{
// avoid recursion in CUDA threads
-#ifdef __CUDA_ARCH__
+#if defined(__CUDA_ARCH__) || defined(__HCC_ACCELERATOR__)
stable_merge_sort_detail::iterative_stable_merge_sort_by_key(exec, first1, last1, first2, comp);
#else
stable_merge_sort_detail::recursive_stable_merge_sort_by_key(exec, first1, last1, first2, comp);

MODIFY
thrust/system/detail/sequential/trivial_copy.h
thrust/system/detail/sequential/trivial_copy.h
@@ -40,7 +40,7 @@ __host__ __device__
std::ptrdiff_t n,
T *result)
{
-#ifndef __CUDA_ARCH__
+#if !defined(__CUDA_ARCH__) && !defined(__HCC_ACCELERATOR__)
std::memmove(result, first, n * sizeof(T));
return result + n;
#else

HIP backend: Adjust #ifdefs to support HIP backend