HIP: hipify for_each, parallel_for (TODO: implement), swap_ranges
MODIFY
thrust/detail/config/compiler.h
thrust/detail/config/compiler.h
#if defined(__CUDA__)
#define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_NVCC
#elif defined(__HCC__)
-// STREAMHPC
-// We don't set THRUST_DEVICE_COMPILER to THRUST_DEVICE_COMPILER_HCC to limit
-// number of headers that need to be ported to HCC before testing simple features like malloc.
-#define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_NVCC //THRUST_DEVICE_COMPILER_HCC
+#define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_HCC
#else
#define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_CLANG
#endif

MODIFY
thrust/system/hip/config.h
thrust/system/hip/config.h
#define THRUST_UNUSED_VAR(expr) do { (void)(expr); } while (0)
-#define __THRUST_HAS_CUDART__ 0
+#ifdef __HIPCC__
+ #ifndef __HIP_DEVICE_COMPILE__
+ #define __THRUST_HAS_HIPRT__ 1
+ #else
+ #define __THRUST_HAS_HIPRT__ 0
+ #endif
#define THRUST_RUNTIME_FUNCTION __host__ __forceinline__
+#else
+ #define __THRUST_HAS_HIPRT__ 0
+#endif
+
#ifdef __HIP_DEVICE_COMPILE__
#define THRUST_DEVICE_CODE

MODIFY
thrust/system/hip/detail/assign_value.h
thrust/system/hip/detail/assign_value.h
@@ -37,6 +37,7 @@ inline __host__ __device__
{
__host__ inline static void host_path(thrust::hip::execution_policy<DerivedPolicy> &exec, Pointer1 dst, Pointer2 src)
{
+ (void) exec; (void) dst; (void) src;
// hip_rocprim::copy(exec, src, src + 1, dst);
}
@@ -66,6 +67,7 @@ inline __host__ __device__
// rotate the systems so that they are ordered the same as (src, dst)
// for the call to thrust::copy
cross_system<System2,System1> rotated_systems = systems.rotate();
+ (void) rotated_systems; (void) dst; (void) src;
// hip_rocprim::copy(rotated_systems, src, src + 1, dst);
}
@@ -73,8 +75,8 @@ inline __host__ __device__
{
// XXX forward the true hip::execution_policy inside systems here
// instead of materializing a tag
- thrust::hip::tag cuda_tag;
- thrust::hip_rocprim::assign_value(cuda_tag, dst, src);
+ thrust::hip::tag hip_tag;
+ thrust::hip_rocprim::assign_value(hip_tag, dst, src);
}
};

MODIFY
thrust/system/hip/detail/for_each.h
thrust/system/hip/detail/for_each.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
#include <iterator>
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
-#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/parallel_for.h>
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/parallel_for.h>
#include <thrust/detail/function.h>
#include <thrust/distance.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
// for_each functor
template <class Input, class UnaryOp>
@@ -76,7 +76,7 @@ namespace cuda_cub {
typedef detail::wrapped_function<UnaryOp, void> wrapped_t;
wrapped_t wrapped_op(op);
- cuda_cub::parallel_for(policy,
+ hip_rocprim::parallel_for(policy,
for_each_f<Input, wrapped_t>(first, wrapped_op),
count);
return first + count;
@@ -94,9 +94,9 @@ namespace cuda_cub {
{
typedef typename iterator_traits<Input>::difference_type size_type;
size_type count = static_cast<size_type>(thrust::distance(first,last));
- return cuda_cub::for_each_n(policy, first, count, op);
+ return hip_rocprim::for_each_n(policy, first, count, op);
}
-} // namespace cuda_cub
+} // namespace hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/detail/par.h
thrust/system/hip/detail/par.h
#include <thrust/detail/config.h>
#include <thrust/detail/allocator/allocator_traits.h>
#include <thrust/detail/execute_with_allocator.h>
-#include <thrust/system/cuda/detail/execution_policy.h>
+#include <thrust/system/hip/detail/execution_policy.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
-__host__ __device__ inline cudaStream_t default_stream()
+__host__ __device__ inline hipStream_t default_stream()
{
- return cudaStreamLegacy;
+ return hipStreamDefault; // STREAMHPC There's not hipStreamLegacy
}
template <class Derived>
-cudaStream_t __host__ __device__
+hipStream_t __host__ __device__
get_stream(execution_policy<Derived> &)
{
return default_stream();
}
template <class Derived>
-cudaError_t THRUST_RUNTIME_FUNCTION
+hipError_t THRUST_RUNTIME_FUNCTION
synchronize_stream(execution_policy<Derived> &)
{
- cudaDeviceSynchronize();
- return cudaGetLastError();
+ hipDeviceSynchronize();
+ return hipGetLastError();
}
@@ -59,16 +59,16 @@ template <class Derived>
struct execute_on_stream_base : execution_policy<Derived>
{
private:
- cudaStream_t stream;
+ hipStream_t stream;
public:
__host__ __device__
- execute_on_stream_base(cudaStream_t stream_ = default_stream())
+ execute_on_stream_base(hipStream_t stream_ = default_stream())
: stream(stream_) {}
__host__ __device__
Derived
- on(cudaStream_t const &s) const
+ on(hipStream_t const &s) const
{
Derived result = derived_cast(*this);
result.stream = s;
@@ -76,24 +76,24 @@ public:
}
private:
- friend cudaStream_t __host__ __device__
+ friend hipStream_t __host__ __device__
get_stream(execute_on_stream_base &exec)
{
return exec.stream;
}
- friend cudaError_t THRUST_RUNTIME_FUNCTION
+ friend hipError_t THRUST_RUNTIME_FUNCTION
synchronize_stream(execute_on_stream_base &exec)
{
-#ifdef __CUDA_ARCH__
-#ifdef __THRUST_HAS_CUDART__
+#ifdef __HIP_DEVICE_COMPILE__
+#ifdef __THRUST_HAS_HIPRT__
THRUST_UNUSED_VAR(exec);
- cudaDeviceSynchronize();
+ hipDeviceSynchronize();
#endif
#else
- cudaStreamSynchronize(exec.stream);
+ hipStreamSynchronize(exec.stream);
#endif
- return cudaGetLastError();
+ return hipGetLastError();
}
};
@@ -104,7 +104,7 @@ struct execute_on_stream : execute_on_stream_base<execute_on_stream>
__host__ __device__
execute_on_stream() : base_t(){};
__host__ __device__
- execute_on_stream(cudaStream_t stream) : base_t(stream){};
+ execute_on_stream(hipStream_t stream) : base_t(stream){};
};
@@ -135,31 +135,31 @@ struct par_t : execution_policy<par_t>
}
execute_on_stream __device__ __host__
- on(cudaStream_t const &stream) const
+ on(hipStream_t const &stream) const
{
return execute_on_stream(stream);
}
};
-#ifdef __CUDA_ARCH__
+#ifdef __HIP_DEVICE_COMPILE__
static const __device__ par_t par;
#else
static const par_t par;
#endif
-} // namespace cuda_
+} // namespace hip_rocprim
namespace system {
-namespace cuda {
- using thrust::cuda_cub::par;
+namespace hip {
+ using thrust::hip_rocprim::par;
namespace detail {
- using thrust::cuda_cub::par_t;
+ using thrust::hip_rocprim::par_t;
}
-} // namesapce cuda
+} // namesapce hip
} // namespace system
-namespace cuda {
-using thrust::cuda_cub::par;
-} // namespace cuda
+namespace hip {
+using thrust::hip_rocprim::par;
+} // namespace hip
END_NS_THRUST

MODIFY
thrust/system/hip/detail/par_to_seq.h
thrust/system/hip/detail/par_to_seq.h
#pragma once
#include <thrust/detail/seq.h>
-#include <thrust/system/cuda/detail/par.h>
+#include <thrust/system/hip/detail/par.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
template <int PAR>
struct has_par : thrust::detail::true_type {};
@@ -80,11 +80,11 @@ cvt_to_seq(Policy& policy)
return cvt_to_seq_impl<Policy>::doit(policy);
}
-#if __THRUST_HAS_CUDART__
-#define THRUST_CUDART_DISPATCH par
+#if __THRUST_HAS_HIPRT__
+#define THRUST_HIPRT_DISPATCH par
#else
-#define THRUST_CUDART_DISPATCH seq
+#define THRUST_HIPRT_DISPATCH seq
#endif
-} // namespace cuda_
+} // namespace hip_rocprim
END_NS_THRUST

MODIFY
thrust/system/hip/detail/parallel_for.h
thrust/system/hip/detail/parallel_for.h
#pragma once
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <thrust/system/hip/config.h>
-#include <thrust/system/cuda/detail/util.h>
+#include <thrust/system/hip/detail/util.h>
#include <thrust/detail/type_traits/result_of_adaptable_function.h>
-#include <thrust/system/cuda/detail/par_to_seq.h>
-#include <thrust/system/cuda/detail/core/agent_launcher.h>
-#include <thrust/system/cuda/detail/par_to_seq.h>
+#include <thrust/system/hip/detail/par_to_seq.h>
+// #include <thrust/system/cuda/detail/core/agent_launcher.h>
+#include <thrust/system/hip/detail/par_to_seq.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
namespace __parallel_for {
- template <int _BLOCK_THREADS,
- int _ITEMS_PER_THREAD = 1,
- int _MIN_BLOCKS = 1>
- struct PtxPolicy
- {
- enum
- {
- BLOCK_THREADS = _BLOCK_THREADS,
- ITEMS_PER_THREAD = _ITEMS_PER_THREAD,
- ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD,
- MIN_BLOCKS = _MIN_BLOCKS
- };
- }; // struct PtxPolicy
-
- template <class Arch, class F>
- struct Tuning;
-
- template <class F>
- struct Tuning<sm30, F>
- {
- typedef PtxPolicy<256, 2> type;
- };
-
-
- template <class F,
- class Size>
- struct ParallelForAgent
- {
- template <class Arch>
- struct PtxPlan : Tuning<Arch, F>::type
- {
- typedef Tuning<Arch, F> tuning;
- };
- typedef core::specialize_plan<PtxPlan> ptx_plan;
-
- enum
- {
- ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD,
- ITEMS_PER_TILE = ptx_plan::ITEMS_PER_TILE,
- BLOCK_THREADS = ptx_plan::BLOCK_THREADS
- };
-
- template <bool IS_FULL_TILE>
- static void THRUST_DEVICE_FUNCTION
- consume_tile(F f,
- Size tile_base,
- int items_in_tile)
- {
-#pragma unroll
- for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
- {
- int idx = BLOCK_THREADS * ITEM + threadIdx.x;
- if (IS_FULL_TILE || idx < items_in_tile)
- f(tile_base + idx);
- }
- }
-
- THRUST_AGENT_ENTRY(F f,
- Size num_items,
- char * /*shmem*/ )
- {
- Size tile_base = blockIdx.x * ITEMS_PER_TILE;
- Size num_remaining = num_items - tile_base;
- int items_in_tile = static_cast<int>(
- num_remaining < ITEMS_PER_TILE ? num_remaining : ITEMS_PER_TILE);
-
- if (items_in_tile == ITEMS_PER_TILE)
- {
- // full tile
- consume_tile<true>(f, tile_base, ITEMS_PER_TILE);
- }
- else
- {
- // partial tile
- consume_tile<false>(f, tile_base, items_in_tile);
- }
- }
- }; // struct ParallelForEagent
-
template <class F,
class Size>
- THRUST_RUNTIME_FUNCTION cudaError_t
+ THRUST_RUNTIME_FUNCTION hipError_t
parallel_for(Size num_items,
F f,
- cudaStream_t stream)
+ hipStream_t stream)
{
- if (num_items == 0)
- return cudaErrorNotSupported;
- using core::AgentLauncher;
- using core::AgentPlan;
-
bool debug_sync = THRUST_DEBUG_SYNC_FLAG;
-
- typedef AgentLauncher<ParallelForAgent<F, Size> > parallel_for_agent;
- AgentPlan parallel_for_plan = parallel_for_agent::get_plan(stream);
-
- parallel_for_agent pfa(parallel_for_plan, num_items, stream, "transform::agent", debug_sync);
- pfa.launch(f, num_items);
- CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
-
- return cudaSuccess;
+ // STREAMHPC TODO implement parallel_for
+ (void) debug_sync;
+ (void) num_items; (void) f; (void) stream;
+ return hipSuccess;
}
} // __parallel_for
@@ -158,22 +69,22 @@ parallel_for(execution_policy<Derived> &policy,
if (count == 0)
return;
- if (__THRUST_HAS_CUDART__)
+ if(__THRUST_HAS_HIPRT__)
{
- cudaStream_t stream = cuda_cub::stream(policy);
- cudaError_t status = __parallel_for::parallel_for(count, f, stream);
- cuda_cub::throw_on_error(status, "parallel_for failed");
+ 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");
}
else
{
-#if !__THRUST_HAS_CUDART__
+#if !__THRUST_HAS_HIPRT__
for (Size idx = 0; idx != count; ++idx)
f(idx);
#endif
}
}
-} // namespace cuda_cub
+} // namespace hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/detail/swap_ranges.h
thrust/system/hip/detail/swap_ranges.h
#pragma once
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <iterator>
-#include <thrust/system/cuda/detail/transform.h>
-#include <thrust/system/cuda/detail/par_to_seq.h>
+// STREAMHPC this include is not required
+// #include <thrust/system/hip/detail/transform.h>
+#include <thrust/system/hip/detail/par_to_seq.h>
#include <thrust/swap.h>
-#include <thrust/system/cuda/detail/parallel_for.h>
+#include <thrust/system/hip/detail/parallel_for.h>
#include <thrust/distance.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
namespace __swap_ranges {
@@ -86,16 +87,16 @@ swap_ranges(execution_policy<Derived> &policy,
size_type num_items = static_cast<size_type>(thrust::distance(first1, last1));
- cuda_cub::parallel_for(policy,
- __swap_ranges::swap_f<ItemsIt1,
- ItemsIt2>(first1, first2),
- num_items);
-
+ hip_rocprim::parallel_for(
+ policy,
+ __swap_ranges::swap_f<ItemsIt1, ItemsIt2>(first1, first2),
+ num_items
+ );
return first2 + num_items;
}
-} // namespace cuda_
+} // namespace hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/detail/util.h
thrust/system/hip/detail/util.h
@@ -132,7 +132,7 @@ throw_on_error(hipError_t status, char const *msg)
#if !defined(__HIP_DEVICE_COMPILE__)
throw thrust::system_error(status, thrust::hip_category(), msg);
#else
-#if __THRUST_HAS_CUDART__
+#if __THRUST_HAS_HIPRT__
printf("Error after %s: %s\n",
msg,
hipGetErrorString(status));

HIP: hipify for_each, parallel_for (TODO: implement), swap_ranges