Extend device_reference tests
MODIFY
test/test_device_reference.cpp
test/test_device_reference.cpp
@@ -144,7 +144,7 @@ TYPED_TEST(DeviceReferenceTests, TestDeviceReferenceConstructorFromDevicePointer
}
// TODO: No device code available for function: _ZN7rocprim6detail16transform_kernelILj256ELj16ERKjN6thrust10device_ptrIjEES6_NS4_8identityIjEEEEvT2_mT3_T4_
-/*TEST(DeviceReferenceTests, TestDeviceReferenceAssignmentFromDeviceReference)
+TEST(DeviceReferenceTests, TestDeviceReferenceAssignmentFromDeviceReference)
{
// test same types
using T0 = int;
@@ -171,7 +171,7 @@ TYPED_TEST(DeviceReferenceTests, TestDeviceReferenceConstructorFromDevicePointer
ASSERT_EQ(13.0f, ref2);
ASSERT_EQ(ref0, ref2);
ASSERT_EQ(ref1, ref2);
-}*/
+}
TYPED_TEST(DeviceReferenceTests,TestDeviceReferenceManipulation)
{
@@ -308,8 +308,7 @@ TYPED_TEST(DeviceReferenceIntegerTests,TestDeviceReferenceIntegerManipulation)
ASSERT_EQ(true, ref1 == ref);
}
-// TODO: Implement swap_ranges
-/*TYPED_TEST(DeviceReferenceTests,TestDeviceReferenceSwap)
+TYPED_TEST(DeviceReferenceTests,TestDeviceReferenceSwap)
{
using T = typename TestFixture::input_type;
@@ -329,6 +328,6 @@ TYPED_TEST(DeviceReferenceIntegerTests,TestDeviceReferenceIntegerManipulation)
ref1.swap(ref2);
ASSERT_EQ(T(7), ref1);
ASSERT_EQ(T(13), ref2);
-}*/
+}
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
thrust/system/detail/adl/iter_swap.h
thrust/system/detail/adl/iter_swap.h
#if 0
#include <thrust/system/cpp/detail/iter_swap.h>
#include <thrust/system/cuda/detail/iter_swap.h>
+#include <thrust/system/hip/detail/iter_swap.h>
#include <thrust/system/omp/detail/iter_swap.h>
#include <thrust/system/tbb/detail/iter_swap.h>
#endif
#define __THRUST_DEVICE_SYSTEM_ITER_SWAP_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/iter_swap.h>
#include __THRUST_DEVICE_SYSTEM_ITER_SWAP_HEADER
#undef __THRUST_DEVICE_SYSTEM_ITER_SWAP_HEADER
-

MODIFY
thrust/system/detail/adl/swap_ranges.h
thrust/system/detail/adl/swap_ranges.h
#if 0
#include <thrust/system/cpp/detail/swap_ranges.h>
#include <thrust/system/cuda/detail/swap_ranges.h>
+#include <thrust/system/hip/detail/swap_ranges.h>
#include <thrust/system/omp/detail/swap_ranges.h>
#include <thrust/system/tbb/detail/swap_ranges.h>
#endif
#define __THRUST_DEVICE_SYSTEM_SWAP_RANGES_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/swap_ranges.h>
#include __THRUST_DEVICE_SYSTEM_SWAP_RANGES_HEADER
#undef __THRUST_DEVICE_SYSTEM_SWAP_RANGES_HEADER
-

MODIFY
thrust/system/hip/detail/iter_swap.h
thrust/system/hip/detail/iter_swap.h
#include <thrust/system/hip/config.h>
#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/system/hip/detail/execution_policy.h>
#include <thrust/swap.h>
BEGIN_NS_THRUST
namespace hip_rocprim {
-template<typename Pointer1, typename Pointer2>
+template<typename DerivedPolicy, typename Pointer1, typename Pointer2>
inline __host__ __device__
-void iter_swap(tag, Pointer1 a, Pointer2 b)
+void iter_swap(thrust::hip::execution_policy<DerivedPolicy> &, Pointer1 a, Pointer2 b)
{
- // XXX war nvbugs/881631
- struct war_nvbugs_881631
- {
- __host__ inline static void host_path(Pointer1 a, Pointer2 b)
- {
- hip_rocprim::swap_ranges(a, a + 1, b);
- }
+#if defined(THRUST_HIP_DEVICE_CODE)
+
+ /*THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((
+ thrust::swap<typename detail::pointer_element<Pointer1>::type, typename detail::pointer_element<Pointer2>::type>
+ ));*/
- __device__ inline static void device_path(Pointer1 a, Pointer2 b)
- {
using thrust::swap;
+ void (*fptr)(typename detail::pointer_element<Pointer1>::type&, typename detail::pointer_element<Pointer2>::type&) = swap;
+ (void) fptr;
+
swap(*thrust::raw_pointer_cast(a),
*thrust::raw_pointer_cast(b));
- }
- };
-#ifndef __HIP_DEVICE_COMPILE__
- return war_nvbugs_881631::host_path(a,b);
#else
- return war_nvbugs_881631::device_path(a,b);
-#endif // __HIP_DEVICE_COMPILE__
+
+ thrust::swap_ranges(a, a + 1, b);
+
+#endif
} // end iter_swap()

Extend device_reference tests