HIP backend: copy assing_value.h, get_value.h from CUDA backend
ADD
thrust/system/hip/detail/assign_value.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#include <thrust/detail/config.h>
+#include <thrust/system/cuda/config.h>
+#include <thrust/system/cuda/detail/execution_policy.h>
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/system/cuda/detail/copy.h>
+
+
+BEGIN_NS_THRUST
+namespace cuda_cub {
+
+
+template<typename DerivedPolicy, typename Pointer1, typename Pointer2>
+inline __host__ __device__
+ void assign_value(thrust::cuda::execution_policy<DerivedPolicy> &exec, Pointer1 dst, Pointer2 src)
+{
+ // XXX war nvbugs/881631
+ struct war_nvbugs_881631
+ {
+ __host__ inline static void host_path(thrust::cuda::execution_policy<DerivedPolicy> &exec, Pointer1 dst, Pointer2 src)
+ {
+ cuda_cub::copy(exec, src, src + 1, dst);
+ }
+
+ __device__ inline static void device_path(thrust::cuda::execution_policy<DerivedPolicy> &, Pointer1 dst, Pointer2 src)
+ {
+ *thrust::raw_pointer_cast(dst) = *thrust::raw_pointer_cast(src);
+ }
+ };
+
+#ifndef __CUDA_ARCH__
+ war_nvbugs_881631::host_path(exec,dst,src);
+#else
+ war_nvbugs_881631::device_path(exec,dst,src);
+#endif // __CUDA_ARCH__
+} // end assign_value()
+
+
+template<typename System1, typename System2, typename Pointer1, typename Pointer2>
+inline __host__ __device__
+ void assign_value(cross_system<System1,System2> &systems, Pointer1 dst, Pointer2 src)
+{
+ // XXX war nvbugs/881631
+ struct war_nvbugs_881631
+ {
+ __host__ inline static void host_path(cross_system<System1,System2> &systems, Pointer1 dst, Pointer2 src)
+ {
+ // 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();
+ cuda_cub::copy(rotated_systems, src, src + 1, dst);
+ }
+
+ __device__ inline static void device_path(cross_system<System1,System2> &, Pointer1 dst, Pointer2 src)
+ {
+ // XXX forward the true cuda::execution_policy inside systems here
+ // instead of materializing a tag
+ thrust::cuda::tag cuda_tag;
+ thrust::cuda_cub::assign_value(cuda_tag, dst, src);
+ }
+ };
+
+#if __CUDA_ARCH__
+ war_nvbugs_881631::device_path(systems,dst,src);
+#else
+ war_nvbugs_881631::host_path(systems,dst,src);
+#endif
+} // end assign_value()
+
+
+
+
+} // end cuda_cub
+END_NS_THRUST
+#endif

ADD
thrust/system/hip/detail/get_value.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#include <thrust/detail/config.h>
+#include <thrust/system/cuda/config.h>
+#include <thrust/system/cuda/detail/cross_system.h>
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/iterator/iterator_traits.h>
+
+BEGIN_NS_THRUST
+namespace cuda_cub {
+
+
+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)
+{
+ 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)
+ {
+ // when called from host code, implement with assign_value
+ // note that this requires a type with default constructor
+ result_type result;
+
+ thrust::host_system_tag host_tag;
+ cross_system<thrust::host_system_tag, DerivedPolicy> systems(host_tag, exec);
+ 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 __CUDA_ARCH__
+ return war_nvbugs_881631::host_path(exec, ptr);
+#else
+ return war_nvbugs_881631::device_path(exec, ptr);
+#endif // __CUDA_ARCH__
+} // 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);
+} // end get_value()
+
+
+} // end cuda_cub
+END_NS_THRUST
+
+#endif

HIP backend: copy assing_value.h, get_value.h from CUDA backend