Implemented tuple tests and fixed transform
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -81,6 +81,9 @@ add_hipthrust_test("thrust.hip.sort" test_sort.cpp)
add_hipthrust_test("thrust.hip.tabulate" test_tabulate.cpp)
add_hipthrust_test("thrust.hip.transform" test_transform.cpp)
add_hipthrust_test("thrust.hip.transform_reduce" test_transform_reduce.cpp)
+add_hipthrust_test("thrust.hip.tuple_reduce" test_tuple_reduce.cpp)
+add_hipthrust_test("thrust.hip.tuple_sort" test_tuple_sort.cpp)
+add_hipthrust_test("thrust.hip.tuple_transform" test_tuple_transform.cpp)
add_hipthrust_test("thrust.hip.tuple" test_tuple.cpp)
add_hipthrust_test("thrust.hip.uninitialized_copy" test_uninitialized_copy.cpp)
add_hipthrust_test("thrust.hip.uninitialized_fill" test_uninitialized_fill.cpp)

MODIFY
test/test_pair_transform.cpp
test/test_pair_transform.cpp
@@ -121,12 +121,12 @@ TYPED_TEST(PairTransformTests, TestPairTransform)
// TODO: add_pairs operator problem.
// add pairs on the host
- /*thrust::transform(h_result.begin(), h_result.end(), h_result.begin(), h_result.begin(), add_pairs());
+ thrust::transform(h_result.begin(), h_result.end(), h_result.begin(), h_result.begin(), add_pairs());
// add pairs on the device
thrust::transform(d_result.begin(), d_result.end(), d_result.begin(), d_result.begin(), add_pairs());
- ASSERT_EQ_QUIET(h_result, d_result);*/
+ ASSERT_EQ_QUIET(h_result, d_result);
}
}

ADD
test/test_tuple_reduce.cpp
+// MIT License
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/tuple.h>
+#include <thrust/sort.h>
+#include <thrust/transform.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include "test_utils.hpp"
+#include "test_assertations.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class TupleReduceTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<short>,
+ Params<int>,
+ Params<long long>,
+ Params<unsigned short>,
+ Params<unsigned int>,
+ Params<unsigned long long>
+> TupleReduceTestsParams;
+
+TYPED_TEST_CASE(TupleReduceTests, TupleReduceTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+struct SumTupleFunctor
+{
+ template <typename Tuple>
+ __host__ __device__
+ Tuple operator()(const Tuple &lhs, const Tuple &rhs)
+ {
+ using thrust::get;
+
+ return thrust::make_tuple(get<0>(lhs) + get<0>(rhs),
+ get<1>(lhs) + get<1>(rhs));
+ }
+};
+
+struct MakeTupleFunctor
+{
+ template<typename T1, typename T2>
+ __host__ __device__
+ thrust::tuple<T1,T2> operator()(T1 &lhs, T2 &rhs)
+ {
+ return thrust::make_tuple(lhs, rhs);
+ }
+};
+
+TYPED_TEST(TupleReduceTests, TestTupleReduce)
+{
+ using T = typename TestFixture::input_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_t1 = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::host_vector<T> h_t2 = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ // zip up the data
+ thrust::host_vector< thrust::tuple<T,T> > h_tuples(size);
+ thrust::transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_tuples.begin(), MakeTupleFunctor());
+
+ // copy to device
+ thrust::device_vector< thrust::tuple<T,T> > d_tuples = h_tuples;
+
+ thrust::tuple<T,T> zero(0,0);
+
+ // sum on host
+ thrust::tuple<T,T> h_result = thrust::reduce(h_tuples.begin(), h_tuples.end(), zero, SumTupleFunctor());
+
+ // sum on device
+ thrust::tuple<T,T> d_result = thrust::reduce(d_tuples.begin(), d_tuples.end(), zero, SumTupleFunctor());
+
+ ASSERT_EQ_QUIET(h_result, d_result);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_tuple_sort.cpp
+// MIT License
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/tuple.h>
+#include <thrust/sort.h>
+#include <thrust/transform.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include "test_utils.hpp"
+#include "test_assertations.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class TupleSortTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<short>,
+ Params<int>,
+ Params<long long>,
+ Params<unsigned short>,
+ Params<unsigned int>,
+ Params<unsigned long long>
+> TupleSortTestsParams;
+
+TYPED_TEST_CASE(TupleSortTests, TupleSortTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+struct MakeTupleFunctor
+{
+ template<typename T1, typename T2>
+ __host__ __device__
+ thrust::tuple<T1,T2> operator()(T1 &lhs, T2 &rhs)
+ {
+ return thrust::make_tuple(lhs, rhs);
+ }
+};
+
+template<int N>
+struct GetFunctor
+{
+ template<typename Tuple>
+ __host__ __device__
+ typename thrust::access_traits<
+ typename thrust::tuple_element<N, Tuple>::type
+ >::const_type
+ operator()(const Tuple &t)
+ {
+ return thrust::get<N>(t);
+ }
+};
+
+TYPED_TEST(TupleSortTests, TestTupleStableSort)
+{
+ using T = typename TestFixture::input_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_keys = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::host_vector<T> h_values = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::host_vector< thrust::tuple<T,T> > h_tuples(size);
+ transform(h_keys.begin(), h_keys.end(),
+ h_values.begin(), h_tuples.begin(),
+ MakeTupleFunctor());
+
+ // copy to device
+ thrust::device_vector< thrust::tuple<T,T> > d_tuples = h_tuples;
+
+ // sort on host
+ thrust::stable_sort(h_tuples.begin(), h_tuples.end());
+
+ // sort on device
+ thrust::stable_sort(d_tuples.begin(), d_tuples.end());
+
+ ASSERT_EQ(true, is_sorted(d_tuples.begin(), d_tuples.end()));
+
+ // select keys
+ thrust::transform(h_tuples.begin(), h_tuples.end(), h_keys.begin(), GetFunctor<0>());
+
+ thrust::device_vector<T> d_keys(h_keys.size());
+ thrust::transform(d_tuples.begin(), d_tuples.end(), d_keys.begin(), GetFunctor<0>());
+
+ // select values
+ thrust::transform(h_tuples.begin(), h_tuples.end(), h_values.begin(), GetFunctor<1>());
+
+ thrust::device_vector<T> d_values(h_values.size());
+ thrust::transform(d_tuples.begin(), d_tuples.end(), d_values.begin(), GetFunctor<1>());
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

ADD
test/test_tuple_transform.cpp
+// MIT License
+//
+// Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/tuple.h>
+#include <thrust/transform.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+#include <hip/hip_runtime.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include "test_utils.hpp"
+#include "test_assertations.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class TupleTransformTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<short>,
+ Params<int>,
+ Params<long long>
+> TupleTransformTestsParams;
+
+TYPED_TEST_CASE(TupleTransformTests, TupleTransformTestsParams);
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+struct MakeTupleFunctor
+{
+ template<typename T1, typename T2>
+ __host__ __device__
+ thrust::tuple<T1,T2> operator()(T1 &lhs, T2 &rhs)
+ {
+ return thrust::make_tuple(lhs, rhs);
+ }
+};
+
+template<int N>
+struct GetFunctor
+{
+ template<typename Tuple>
+ __host__ __device__
+ typename thrust::access_traits<
+ typename thrust::tuple_element<N, Tuple>::type
+ >::const_type
+ operator()(const Tuple &t)
+ {
+ return thrust::get<N>(t);
+ }
+};
+
+TYPED_TEST(TupleTransformTests, TestTupleTransform)
+{
+ using T = typename TestFixture::input_type;
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<T> h_t1 = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ thrust::host_vector<T> h_t2 = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ // zip up the data
+ thrust::host_vector< thrust::tuple<T,T> > h_tuples(size);
+ thrust::transform(h_t1.begin(), h_t1.end(),
+ h_t2.begin(), h_tuples.begin(),
+ MakeTupleFunctor());
+
+ // copy to device
+ thrust::device_vector< thrust::tuple<T,T> > d_tuples = h_tuples;
+
+ thrust::device_vector<T> d_t1(size), d_t2(size);
+
+ // select 0th
+ thrust::transform(d_tuples.begin(), d_tuples.end(), d_t1.begin(), GetFunctor<0>());
+
+ // select 1st
+ thrust::transform(d_tuples.begin(), d_tuples.end(), d_t2.begin(), GetFunctor<1>());
+
+ ASSERT_EQ(h_t1, d_t1);
+ ASSERT_EQ(h_t2, d_t2);
+
+ ASSERT_EQ_QUIET(h_tuples, d_tuples);
+ }
+}
+
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
thrust/system/hip/detail/transform.h
thrust/system/hip/detail/transform.h
@@ -370,7 +370,7 @@ transform(execution_policy<Derived> &policy,
(void) policy;
while (first != last)
{
- *result++ = transform_op(*first++);
+ *result++ = transform_op(raw_reference_cast(*first++));
}
return result;
}
@@ -453,7 +453,7 @@ transform(execution_policy<Derived> &policy,
(void) policy;
while (first1 != last1)
{
- *result++ = transform_op(*first1++, *first2++);
+ *result++ = transform_op(raw_reference_cast(*first1++), raw_reference_cast(*first2++));
}
return result;
}

Implemented tuple tests and fixed transform