Add workaround for "No device code available" to get_value, assign_value
MODIFY
thrust/system/detail/adl/assign_value.h
thrust/system/detail/adl/assign_value.h
#if 0
#include <thrust/system/cpp/detail/assign_value.h>
#include <thrust/system/cuda/detail/assign_value.h>
+#include <thrust/system/hip/detail/assign_value.h>
#include <thrust/system/omp/detail/assign_value.h>
#include <thrust/system/tbb/detail/assign_value.h>
#endif
#define __THRUST_DEVICE_SYSTEM_ASSIGN_VALUE_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/assign_value.h>
#include __THRUST_DEVICE_SYSTEM_ASSIGN_VALUE_HEADER
#undef __THRUST_DEVICE_SYSTEM_ASSIGN_VALUE_HEADER
-

MODIFY
thrust/system/detail/adl/get_value.h
thrust/system/detail/adl/get_value.h
#if 0
#include <thrust/system/cpp/detail/get_value.h>
#include <thrust/system/cuda/detail/get_value.h>
+#include <thrust/system/hip/detail/get_value.h>
#include <thrust/system/omp/detail/get_value.h>
#include <thrust/system/tbb/detail/get_value.h>
#endif
#define __THRUST_DEVICE_SYSTEM_GET_VALUE_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/get_value.h>
#include __THRUST_DEVICE_SYSTEM_GET_VALUE_HEADER
#undef __THRUST_DEVICE_SYSTEM_GET_VALUE_HEADER
-

MODIFY
thrust/system/hip/detail/assign_value.h
thrust/system/hip/detail/assign_value.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <thrust/detail/config.h>
+#include <thrust/detail/raw_pointer_cast.h>
#include <thrust/system/hip/config.h>
#include <thrust/system/hip/detail/execution_policy.h>
-#include <thrust/detail/raw_pointer_cast.h>
#include <thrust/system/hip/detail/copy.h>
@@ -32,6 +32,7 @@ template<typename DerivedPolicy, typename Pointer1, typename Pointer2>
inline __host__ __device__
void assign_value(thrust::hip::execution_policy<DerivedPolicy> &exec, Pointer1 dst, Pointer2 src)
{
+ // STREAMHPC TODO add workaround?
#ifdef __HIP_DEVICE_COMPILE__
(void) exec;
*thrust::raw_pointer_cast(dst) = *thrust::raw_pointer_cast(src);
@@ -45,13 +46,21 @@ template<typename System1, typename System2, typename Pointer1, typename Pointer
inline __host__ __device__
void assign_value(cross_system<System1, System2> &systems, Pointer1 dst, Pointer2 src)
{
-#ifdef __HIP_DEVICE_COMPILE__
- (void) systems;
+ // STREAMHPC WORKAROUND
+#if defined(THRUST_HIP_DEVICE_CODE)
+
+ THRUST_UNUSED_VAR(systems);
+ Pointer1 (*fptr)(cross_system<System2, System1>, Pointer2, Pointer2, Pointer1) = hip_rocprim::copy;
+ (void) fptr;
+
thrust::hip::tag hip_tag;
thrust::hip_rocprim::assign_value(hip_tag, dst, src);
+
#else
+
cross_system<System2, System1> rotated_systems = systems.rotate();
hip_rocprim::copy(rotated_systems, src, src + 1, dst);
+
#endif
} // end assign_value()

MODIFY
thrust/system/hip/detail/copy.h
thrust/system/hip/detail/copy.h
@@ -63,7 +63,7 @@ template <class System1,
class System2,
class InputIterator,
class OutputIterator>
-OutputIterator __host__
+OutputIterator __host__ /* STREAMHPC WORKAROUND */ __device__
copy(cross_system<System1, System2> systems,
InputIterator first,
InputIterator last,
@@ -167,7 +167,7 @@ template <class System1,
class System2,
class InputIterator,
class OutputIterator>
-OutputIterator __host__
+OutputIterator __host__ /* STREAMHPC WORKAROUND */ __device__
copy(cross_system<System1, System2> systems,
InputIterator first,
InputIterator last,
@@ -187,21 +187,7 @@ copy_n(cross_system<System1, System2> systems,
Size n,
OutputIterator result)
{
- // STREAMHPC WORKAROUND
-#if defined(THRUST_HIP_DEVICE_CODE)
-
- THRUST_UNUSED_VAR(systems);
- THRUST_UNUSED_VAR(first);
- THRUST_UNUSED_VAR(n);
- OutputIterator (*ptr)(cross_system<System1, System2>, InputIterator, Size, OutputIterator) = __copy::cross_system_copy_n;
- (void) ptr;
- return result;
-
-#else
-
return __copy::cross_system_copy_n(systems, first, n, result);
-
-#endif
} // end copy_n()

MODIFY
thrust/system/hip/detail/get_value.h
thrust/system/hip/detail/get_value.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <thrust/detail/config.h>
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/system/detail/adl/assign_value.h>
#include <thrust/system/hip/config.h>
#include <thrust/system/hip/detail/cross_system.h>
-#include <thrust/detail/raw_pointer_cast.h>
#include <thrust/iterator/iterator_traits.h>
BEGIN_NS_THRUST
namespace hip_rocprim {
-namespace
-{
-
-
template<typename DerivedPolicy, typename Pointer>
inline __host__ __device__
typename thrust::iterator_value<Pointer>::type
- get_value_msvc2005_war(execution_policy<DerivedPolicy> &exec, Pointer ptr)
+ get_value(execution_policy<DerivedPolicy> &exec, Pointer ptr)
{
typedef typename thrust::iterator_value<Pointer>::type result_type;
- // XXX war nvbugs/881631
- struct war_nvbugs_881631
- {
- __host__ inline static result_type host_path(execution_policy<DerivedPolicy> &exec, Pointer ptr)
- {
+ // STREAMHPC WORKAROUND
+#if defined(THRUST_HIP_DEVICE_CODE)
+
+ THRUST_UNUSED_VAR(exec);
+ void (*fptr)(cross_system<thrust::host_system_tag, DerivedPolicy> &, result_type *, Pointer) = assign_value;
+ (void) fptr;
+
+ return *thrust::raw_pointer_cast(ptr);
+
+#else
+
// when called from host code, implement with assign_value
// note that this requires a type with default constructor
result_type result;
@@ -52,32 +55,8 @@ inline __host__ __device__
assign_value(systems, &result, ptr);
return result;
- }
- __device__ inline static result_type device_path(execution_policy<DerivedPolicy> &, Pointer ptr)
- {
- // when called from device code, just do simple deref
- return *thrust::raw_pointer_cast(ptr);
- }
- };
-
-#ifndef __HIP_DEVICE_COMPILE__
- return war_nvbugs_881631::host_path(exec, ptr);
-#else
- return war_nvbugs_881631::device_path(exec, ptr);
-#endif // __HIP_DEVICE_COMPILE__
-} // end get_value_msvc2005_war()
-
-
-} // end anon namespace
-
-
-template<typename DerivedPolicy, typename Pointer>
-inline __host__ __device__
- typename thrust::iterator_value<Pointer>::type
- get_value(execution_policy<DerivedPolicy> &exec, Pointer ptr)
-{
- return get_value_msvc2005_war(exec,ptr);
+#endif
} // end get_value()

MODIFY
thrust/system/hip/detail/internal/copy_cross_system.h
thrust/system/hip/detail/internal/copy_cross_system.h
@@ -306,7 +306,7 @@ namespace __copy {
class System2,
class InputIterator,
class OutputIterator>
- OutputIterator __host__
+ OutputIterator __host__ /* STREAMHPC WORKAROUND */ __device__
cross_system_copy(cross_system<System1, System2> systems,
InputIterator begin,
InputIterator end,

Add workaround for "No device code available" to get_value, assign_value