HIP backend: copy for_each from CUDA backend
ADD
thrust/system/hip/detail/for_each.h
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#include <iterator>
+#include <thrust/system/cuda/config.h>
+
+#include <thrust/system/cuda/detail/util.h>
+#include <thrust/system/cuda/detail/parallel_for.h>
+#include <thrust/detail/function.h>
+#include <thrust/distance.h>
+
+BEGIN_NS_THRUST
+
+namespace cuda_cub {
+
+ // for_each functor
+ template <class Input, class UnaryOp>
+ struct for_each_f
+ {
+ Input input;
+ UnaryOp op;
+
+ THRUST_FUNCTION
+ for_each_f(Input input, UnaryOp op)
+ : input(input), op(op) {}
+
+ template <class Size>
+ THRUST_DEVICE_FUNCTION void operator()(Size idx)
+ {
+ op(raw_reference_cast(input[idx]));
+ }
+ };
+
+ //-------------------------
+ // Thrust API entry points
+ //-------------------------
+
+ // for_each_n
+ template <class Derived,
+ class Input,
+ class Size,
+ class UnaryOp>
+ Input THRUST_FUNCTION
+ for_each_n(execution_policy<Derived> &policy,
+ Input first,
+ Size count,
+ UnaryOp op)
+ {
+ typedef detail::wrapped_function<UnaryOp, void> wrapped_t;
+ wrapped_t wrapped_op(op);
+
+ cuda_cub::parallel_for(policy,
+ for_each_f<Input, wrapped_t>(first, wrapped_op),
+ count);
+ return first + count;
+ }
+
+ // for_each
+ template <class Derived,
+ class Input,
+ class UnaryOp>
+ Input THRUST_FUNCTION
+ for_each(execution_policy<Derived> &policy,
+ Input first,
+ Input last,
+ UnaryOp op)
+ {
+ 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);
+ }
+} // namespace cuda_cub
+
+END_NS_THRUST
+#endif

ADD
thrust/system/hip/detail/par.h
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+#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>
+
+BEGIN_NS_THRUST
+namespace cuda_cub {
+
+__host__ __device__ inline cudaStream_t default_stream()
+{
+ return cudaStreamLegacy;
+}
+
+template <class Derived>
+cudaStream_t __host__ __device__
+get_stream(execution_policy<Derived> &)
+{
+ return default_stream();
+}
+
+template <class Derived>
+cudaError_t THRUST_RUNTIME_FUNCTION
+synchronize_stream(execution_policy<Derived> &)
+{
+ cudaDeviceSynchronize();
+ return cudaGetLastError();
+}
+
+
+template <class Derived>
+struct execute_on_stream_base : execution_policy<Derived>
+{
+private:
+ cudaStream_t stream;
+
+public:
+ __host__ __device__
+ execute_on_stream_base(cudaStream_t stream_ = default_stream())
+ : stream(stream_) {}
+
+ __host__ __device__
+ Derived
+ on(cudaStream_t const &s) const
+ {
+ Derived result = derived_cast(*this);
+ result.stream = s;
+ return result;
+ }
+
+private:
+ friend cudaStream_t __host__ __device__
+ get_stream(execute_on_stream_base &exec)
+ {
+ return exec.stream;
+ }
+
+ friend cudaError_t THRUST_RUNTIME_FUNCTION
+ synchronize_stream(execute_on_stream_base &exec)
+ {
+#ifdef __CUDA_ARCH__
+#ifdef __THRUST_HAS_CUDART__
+ THRUST_UNUSED_VAR(exec);
+ cudaDeviceSynchronize();
+#endif
+#else
+ cudaStreamSynchronize(exec.stream);
+#endif
+ return cudaGetLastError();
+ }
+};
+
+struct execute_on_stream : execute_on_stream_base<execute_on_stream>
+{
+ typedef execute_on_stream_base<execute_on_stream> base_t;
+
+ __host__ __device__
+ execute_on_stream() : base_t(){};
+ __host__ __device__
+ execute_on_stream(cudaStream_t stream) : base_t(stream){};
+};
+
+
+struct par_t : execution_policy<par_t>
+{
+ typedef execution_policy<par_t> base_t;
+
+ __device__ __host__
+ par_t() : base_t() {}
+
+ template <class Allocator>
+ struct enable_alloc
+ {
+ typedef typename thrust::detail::enable_if<
+ thrust::detail::is_allocator<Allocator>::value,
+ thrust::detail::execute_with_allocator<Allocator,
+ execute_on_stream_base> >::type
+ type;
+ };
+
+ template <class Allocator>
+ __host__ __device__ typename enable_alloc<Allocator>::type
+ operator()(Allocator &alloc) const
+ {
+ return thrust::detail::execute_with_allocator<
+ Allocator,
+ execute_on_stream_base>(alloc);
+ }
+
+ execute_on_stream __device__ __host__
+ on(cudaStream_t const &stream) const
+ {
+ return execute_on_stream(stream);
+ }
+};
+
+#ifdef __CUDA_ARCH__
+static const __device__ par_t par;
+#else
+static const par_t par;
+#endif
+} // namespace cuda_
+
+namespace system {
+namespace cuda {
+ using thrust::cuda_cub::par;
+ namespace detail {
+ using thrust::cuda_cub::par_t;
+ }
+} // namesapce cuda
+} // namespace system
+
+namespace cuda {
+using thrust::cuda_cub::par;
+} // namespace cuda
+
+END_NS_THRUST
+

ADD
thrust/system/hip/detail/par_to_seq.h
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+#include <thrust/detail/seq.h>
+#include <thrust/system/cuda/detail/par.h>
+
+BEGIN_NS_THRUST
+namespace cuda_cub {
+
+template <int PAR>
+struct has_par : thrust::detail::true_type {};
+
+template <>
+struct has_par<0> : thrust::detail::false_type {};
+
+template<class Policy>
+struct cvt_to_seq_impl
+{
+ typedef thrust::detail::seq_t seq_t;
+
+ static seq_t __host__ __device__
+ doit(Policy&)
+ {
+ return seq_t();
+ }
+}; // cvt_to_seq_impl
+
+#if 0
+template <class Allocator>
+struct cvt_to_seq_impl<
+ thrust::detail::execute_with_allocator<Allocator,
+ execute_on_stream_base> >
+{
+ typedef thrust::detail::execute_with_allocator<Allocator,
+ execute_on_stream_base>
+ Policy;
+ typedef thrust::detail::execute_with_allocator<
+ Allocator,
+ thrust::system::detail::sequential::execution_policy>
+ seq_t;
+
+
+ static seq_t __host__ __device__
+ doit(Policy& policy)
+ {
+ return seq_t(policy.m_alloc);
+ }
+}; // specialization of struct cvt_to_seq_impl
+#endif
+
+template <class Policy>
+typename cvt_to_seq_impl<Policy>::seq_t __host__ __device__
+cvt_to_seq(Policy& policy)
+{
+ return cvt_to_seq_impl<Policy>::doit(policy);
+}
+
+#if __THRUST_HAS_CUDART__
+#define THRUST_CUDART_DISPATCH par
+#else
+#define THRUST_CUDART_DISPATCH seq
+#endif
+
+} // namespace cuda_
+END_NS_THRUST

ADD
thrust/system/hip/detail/parallel_for.h
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#include <thrust/system/cuda/config.h>
+
+#include <thrust/system/cuda/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>
+
+BEGIN_NS_THRUST
+
+namespace cuda_cub {
+
+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
+ parallel_for(Size num_items,
+ F f,
+ cudaStream_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;
+ }
+} // __parallel_for
+
+__thrust_exec_check_disable__
+template <class Derived,
+ class F,
+ class Size>
+void __host__ __device__
+parallel_for(execution_policy<Derived> &policy,
+ F f,
+ Size count)
+{
+ if (count == 0)
+ return;
+
+ if (__THRUST_HAS_CUDART__)
+ {
+ 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");
+ }
+ else
+ {
+#if !__THRUST_HAS_CUDART__
+ for (Size idx = 0; idx != count; ++idx)
+ f(idx);
+#endif
+ }
+}
+
+} // namespace cuda_cub
+
+END_NS_THRUST
+#endif

HIP backend: copy for_each from CUDA backend