Adjust some functions for HIP-clang
* fix __host__ and __device__ spacifiers.
* improve workarounds in transforma and parallel_for,
so kernels are preserved in binaries on both HCC and HIP-clang.
MODIFY
test/test_transform_reduce.cpp
test/test_transform_reduce.cpp
@@ -105,7 +105,7 @@ TYPED_TEST(TransformReduceIntegerPrimitiveTests, TestTransformReduce)
T gpu_result = thrust::transform_reduce(
d_data.begin(), d_data.end(), thrust::negate<T>(), init, thrust::plus<T>());
- ASSERT_NEAR(cpu_result, gpu_result, std::abs(T(0.01 * cpu_result)));
+ ASSERT_EQ(cpu_result, gpu_result);
}
}
@@ -128,7 +128,7 @@ TYPED_TEST(TransformReduceIntegerPrimitiveTests, TestTransformReduceFromConst)
T gpu_result = thrust::transform_reduce(
d_data.cbegin(), d_data.cend(), thrust::negate<T>(), init, thrust::plus<T>());
- ASSERT_NEAR(cpu_result, gpu_result, std::abs(T(0.01 * cpu_result)));
+ ASSERT_EQ(cpu_result, gpu_result);
}
}

MODIFY
test/test_uninitialized_copy.cpp
test/test_uninitialized_copy.cpp
@@ -141,6 +141,7 @@ TYPED_TEST(UninitializedCopyTests, TestUninitializedCopyNSimplePOD)
struct CopyConstructTest
{
+ __host__ __device__
CopyConstructTest(void)
: copy_constructed_on_host(false)
, copy_constructed_on_device(false)

MODIFY
thrust/complex.h
thrust/complex.h
@@ -878,7 +878,6 @@ operator<<(std::basic_ostream<CharT, Traits>& os, const complex<T>& z);
* \param z The \p complex number to set.
*/
template <typename T, typename CharT, typename Traits>
-__host__ __device__
std::basic_istream<CharT, Traits>&
operator>>(std::basic_istream<CharT, Traits>& is, complex<T>& z);

MODIFY
thrust/system/hip/detail/fill.h
thrust/system/hip/detail/fill.h
@@ -49,7 +49,7 @@ fill_n(execution_policy<Derived>& policy,
thrust::make_counting_iterator<Size>(0),
thrust::make_counting_iterator<Size>(count),
first,
- [value](Size) { return value; });
+ [value] __host__ __device__ (Size) { return value; });
} // func fill_n
template <class Derived, class ForwardIterator, class T>

MODIFY
thrust/system/hip/detail/internal/copy_device_to_device.h
thrust/system/hip/detail/internal/copy_device_to_device.h
@@ -40,7 +40,7 @@ namespace hip_rocprim
namespace __copy
{
template <class Derived, class InputIt, class OutputIt>
- OutputIt THRUST_HIP_HOST_FUNCTION
+ OutputIt THRUST_HIP_FUNCTION
device_to_device(execution_policy<Derived>& policy,
InputIt first,
InputIt last,

MODIFY
thrust/system/hip/detail/parallel_for.h
thrust/system/hip/detail/parallel_for.h
@@ -112,19 +112,38 @@ parallel_for(execution_policy<Derived>& policy, F f, Size count)
if(count == 0)
return;
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((__parallel_for::parallel_for<F, Size>));
-#if __THRUST_HAS_HIPRT__
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
{
+ __host__
+ static void par(execution_policy<Derived>& policy, F f, Size count)
+ {
+#if __HCC__ && __HIP_DEVICE_COMPILE__
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((__parallel_for::parallel_for<F, Size>));
+ (void)policy;
+ (void)f;
+ (void)count;
+#else
hipStream_t stream = hip_rocprim::stream(policy);
hipError_t status = __parallel_for::parallel_for(count, f, stream);
hip_rocprim::throw_on_error(status, "parallel_for failed");
+#endif
}
-#else
+
+ __device__
+ static void seq(execution_policy<Derived>& policy, F f, Size count)
{
(void)policy;
for(Size idx = 0; idx != count; ++idx)
f(idx);
}
+ };
+
+#if __THRUST_HAS_HIPRT__
+ workaround::par(policy, f, count);
+#else
+ workaround::seq(policy, f, count);
#endif
}

MODIFY
thrust/system/hip/detail/transform.h
thrust/system/hip/detail/transform.h
@@ -319,20 +319,40 @@ transform(execution_policy<Derived>& policy,
if(num_items == 0)
return result;
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
+ {
+ __host__
+ static OutputIt par(execution_policy<Derived>& policy,
+ InputIt first,
+ InputIt last,
+ OutputIt result,
+ TransformOp transform_op)
+ {
+ size_type num_items = static_cast<size_type>(thrust::distance(first, last));
+#if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(rocprim::transform<rocprim::default_config, InputIt, OutputIt, TransformOp>)
);
-#if __THRUST_HAS_HIPRT__
- {
+ (void)policy;
+ (void)transform_op;
+#else
hipStream_t stream = hip_rocprim::stream(policy);
bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
hipError_t status
= rocprim::transform(first, result, num_items, transform_op, stream, debug_sync);
hip_rocprim::throw_on_error(status, "transform failed");
-
+#endif
return result + num_items;
}
-#else
+
+ __device__
+ static OutputIt seq(execution_policy<Derived>& policy,
+ InputIt first,
+ InputIt last,
+ OutputIt result,
+ TransformOp transform_op)
{
(void)policy;
while(first != last)
@@ -341,6 +361,12 @@ transform(execution_policy<Derived>& policy,
}
return result;
}
+ };
+
+#if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, first, last, result, transform_op);
+#else
+ return workaround::seq(policy, first, last, result, transform_op);
#endif
} // func transform

Adjust some functions for HIP-clang

* fix __host__ and __device__ spacifiers.
* improve workarounds in transforma and parallel_for,
so kernels are preserved in binaries on both HCC and HIP-clang.