HIP backend: fix __THRUST_HAS_HIPRT__, rename HIP macros
MODIFY
thrust/system/hip/config.h
thrust/system/hip/config.h
#define THRUST_UNUSED_VAR(expr) do { (void)(expr); } while (0)
-#ifdef __HIPCC__
- #ifndef __HIP_DEVICE_COMPILE__
- #define __THRUST_HAS_HIPRT__ 1
- #else
+#ifdef __HCC__
+ #ifdef __HCC_ACCELERATOR__
#define __THRUST_HAS_HIPRT__ 0
+ #define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
+ #else
+ #define __THRUST_HAS_HIPRT__ 1
+ #define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
#endif
- #define THRUST_RUNTIME_FUNCTION __host__ __forceinline__
#else
#define __THRUST_HAS_HIPRT__ 0
+ #define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
#endif
-
-#ifdef __HIP_DEVICE_COMPILE__
-#define THRUST_DEVICE_CODE
-#endif
-
-#ifdef THRUST_AGENT_ENTRY_NOINLINE
-#define THRUST_AGENT_ENTRY_INLINE_ATTR __noinline__
-#else
-#define THRUST_AGENT_ENTRY_INLINE_ATTR __forceinline__
+#ifdef __HCC_ACCELERATOR__
+#define THRUST_HIP_DEVICE_CODE
#endif
-#define THRUST_DEVICE_FUNCTION __device__ __forceinline__
-#define THRUST_HOST_FUNCTION __host__ __forceinline__
-#define THRUST_FUNCTION __host__ __device__ __forceinline__
-#if 0
-#define THRUST_ARGS(...) __VA_ARGS__
-#define THRUST_STRIP_PARENS(X) X
-#define THRUST_AGENT_ENTRY(ARGS) THRUST_FUNCTION static void entry(THRUST_STRIP_PARENS(THRUST_ARGS ARGS))
-#else
-#define THRUST_AGENT_ENTRY(...) THRUST_AGENT_ENTRY_INLINE_ATTR __device__ static void entry(__VA_ARGS__)
-#endif
+#define THRUST_HIP_DEVICE_FUNCTION __device__ __forceinline__
+#define THRUST_HIP_HOST_FUNCTION __host__ __forceinline__
+#define THRUST_HIP_FUNCTION __host__ __device__ __forceinline__
-#ifdef THRUST_DEBUG_SYNC
-#define THRUST_DEBUG_SYNC_FLAG true
+#ifdef THRUST_HIP_DEBUG_SYNC
+#define THRUST_HIP_DEBUG_SYNC_FLAG true
#define DEBUG
#else
-#define THRUST_DEBUG_SYNC_FLAG false
+#define THRUST_HIP_DEBUG_SYNC_FLAG false
#endif
#define THRUST_ROCPRIM_NS_PREFIX namespace thrust { namespace hip_rocprim {

MODIFY
thrust/system/hip/detail/for_each.h
thrust/system/hip/detail/for_each.h
@@ -47,12 +47,12 @@ namespace hip_rocprim {
Input input;
UnaryOp op;
- THRUST_FUNCTION
+ THRUST_HIP_FUNCTION
for_each_f(Input input, UnaryOp op)
: input(input), op(op) {}
template <class Size>
- THRUST_FUNCTION
+ THRUST_HIP_FUNCTION
void operator()(Size idx)
{
op(raw_reference_cast(input[idx]));
@@ -68,7 +68,7 @@ namespace hip_rocprim {
class Input,
class Size,
class UnaryOp>
- Input THRUST_FUNCTION
+ Input THRUST_HIP_FUNCTION
for_each_n(execution_policy<Derived> &policy,
Input first,
Size count,
@@ -87,7 +87,7 @@ namespace hip_rocprim {
template <class Derived,
class Input,
class UnaryOp>
- Input THRUST_FUNCTION
+ Input THRUST_HIP_FUNCTION
for_each(execution_policy<Derived> &policy,
Input first,
Input last,

MODIFY
thrust/system/hip/detail/par.h
thrust/system/hip/detail/par.h
@@ -47,7 +47,7 @@ get_stream(execution_policy<Derived> &)
}
template <class Derived>
-hipError_t THRUST_RUNTIME_FUNCTION
+hipError_t THRUST_HIP_RUNTIME_FUNCTION
synchronize_stream(execution_policy<Derived> &)
{
hipDeviceSynchronize();
@@ -82,7 +82,7 @@ private:
return exec.stream;
}
- friend hipError_t THRUST_RUNTIME_FUNCTION
+ friend hipError_t THRUST_HIP_RUNTIME_FUNCTION
synchronize_stream(execute_on_stream_base &exec)
{
#ifdef __HIP_DEVICE_COMPILE__

MODIFY
thrust/system/hip/detail/parallel_for.h
thrust/system/hip/detail/parallel_for.h
@@ -54,8 +54,8 @@ namespace __parallel_for {
unsigned int ItemsPerThread,
class F,
class Size>
- __global__ void
- kernel(F f, Size num_items)
+ __global__
+ void kernel(F f, Size num_items)
{
constexpr auto items_per_block = BlockSize * ItemsPerThread;
Size tile_base = blockIdx.x * items_per_block;
@@ -87,14 +87,13 @@ namespace __parallel_for {
template <class F,
class Size>
- THRUST_RUNTIME_FUNCTION hipError_t
+ THRUST_HIP_RUNTIME_FUNCTION hipError_t
parallel_for(Size num_items,
F f,
hipStream_t stream)
{
using config = kernel_config<256, 1>;
-
- bool debug_sync = THRUST_DEBUG_SYNC_FLAG;
+ bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
// STREAMHPC Use debug_sync
(void) debug_sync;
@@ -104,9 +103,7 @@ namespace __parallel_for {
const auto number_of_blocks = (num_items + items_per_block - 1)/items_per_block;
hipLaunchKernelGGL(
- HIP_KERNEL_NAME(kernel<
- block_size, items_per_thread, F, Size
- >),
+ HIP_KERNEL_NAME(kernel<block_size, items_per_thread, F, Size>),
dim3(number_of_blocks), dim3(block_size), 0, stream,
f, num_items
);
@@ -131,7 +128,6 @@ parallel_for(execution_policy<Derived> &policy,
#if __THRUST_HAS_HIPRT__
{
- std::cout << "HOST PATH" << std::endl;
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");

MODIFY
thrust/system/hip/detail/swap_ranges.h
thrust/system/hip/detail/swap_ranges.h
@@ -52,12 +52,12 @@ namespace __swap_ranges {
typedef typename iterator_traits<ItemsIt1>::value_type value1_type;
typedef typename iterator_traits<ItemsIt2>::value_type value2_type;
- THRUST_FUNCTION
+ THRUST_HIP_FUNCTION
swap_f(ItemsIt1 items1_, ItemsIt2 items2_)
: items1(items1_), items2(items2_) {}
template<class Size>
- void THRUST_DEVICE_FUNCTION operator()(Size idx)
+ void THRUST_HIP_DEVICE_FUNCTION operator()(Size idx)
{
value1_type item1 = items1[idx];
value2_type item2 = items2[idx];

MODIFY
thrust/system/hip/detail/util.h
thrust/system/hip/detail/util.h
@@ -56,7 +56,7 @@ stream(execution_policy<Derived> &policy)
}
template <class Type>
-THRUST_HOST_FUNCTION hipError_t
+THRUST_HIP_HOST_FUNCTION hipError_t
trivial_copy_from_device(Type * dst,
Type const * src,
size_t count,
@@ -75,7 +75,7 @@ trivial_copy_from_device(Type * dst,
}
template <class Type>
-THRUST_HOST_FUNCTION hipError_t
+THRUST_HIP_HOST_FUNCTION hipError_t
trivial_copy_to_device(Type * dst,
Type const * src,
size_t count,

HIP backend: fix __THRUST_HAS_HIPRT__, rename HIP macros