Implemented dereference tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -65,6 +65,7 @@ add_hipthrust_test("thrust.hip.copy" test_copy.cpp)
add_hipthrust_test("thrust.hip.copy_n" test_copy_n.cpp)
add_hipthrust_test("thrust.hip.count" test_count.cpp)
add_hipthrust_test("thrust.hip.counting_iterator" test_counting_iterator.cpp)
+add_hipthrust_test("thrust.hip.dereference" test_dereference.cpp)
add_hipthrust_test("thrust.hip.device_delete" test_device_delete.cpp)
add_hipthrust_test("thrust.hip.device_ptr" test_device_ptr.cpp)
add_hipthrust_test("thrust.hip.device_reference" test_device_reference.cpp)

ADD
test/test_dereference.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>
+
+#include <thrust/device_vector.h>
+#include <thrust/device_ptr.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/transform_iterator.h>
+#include <thrust/iterator/counting_iterator.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"
+
+template<
+ class Input
+>
+struct Params
+{
+ using input_type = Input;
+};
+
+template<class Params>
+class DereferenceTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> DereferenceTestsParams;
+
+TYPED_TEST_CASE(DereferenceTests, DereferenceTestsParams);
+
+
+template <typename Iterator1, typename Iterator2>
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
+__global__
+#endif
+void simple_copy_on_device(Iterator1 first1, Iterator1 last1, Iterator2 first2)
+{
+ while(first1 != last1)
+ *(first2++) = *(first1++);
+}
+
+template <typename Iterator1, typename Iterator2>
+void simple_copy(Iterator1 first1, Iterator1 last1, Iterator2 first2)
+{
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP
+ hipLaunchKernelGGL(
+ HIP_KERNEL_NAME(simple_copy_on_device<Iterator1, Iterator2>),
+ dim3(1), dim3(1), 0, 0,
+ first1, last1, first2
+ );
+#else
+ simple_copy_on_device(first1, last1, first2);
+#endif
+}
+
+TEST(DereferenceTests, TestDeviceDereferenceDeviceVectorIterator)
+{
+ thrust::device_vector<int> input = get_random_data<int>(100,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+ thrust::device_vector<int> output(input.size(), 0);
+
+ simple_copy(input.begin(), input.end(), output.begin());
+
+ ASSERT_EQ(input, output);
+}
+
+TEST(DereferenceTests, TestDeviceDereferenceDevicePtr)
+{
+ thrust::device_vector<int> input = get_random_data<int>(100,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+ thrust::device_vector<int> output(input.size(), 0);
+
+ thrust::device_ptr<int> _first1 = &input[0];
+ thrust::device_ptr<int> _last1 = _first1 + input.size();
+ thrust::device_ptr<int> _first2 = &output[0];
+
+ simple_copy(_first1, _last1, _first2);
+
+ ASSERT_EQ(input, output);
+}
+
+TEST(DereferenceTests, TestDeviceDereferenceTransformIterator)
+{
+ thrust::device_vector<int> input = get_random_data<int>(100,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+ thrust::device_vector<int> output(input.size(), 0);
+
+ simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity<int>()),
+ thrust::make_transform_iterator(input.end (), thrust::identity<int>()),
+ output.begin());
+
+ ASSERT_EQ(input, output);
+}
+
+TEST(DereferenceTests, TestDeviceDereferenceCountingIterator)
+{
+ thrust::counting_iterator<int> first(1);
+ thrust::counting_iterator<int> last(6);
+
+ thrust::device_vector<int> output(5);
+
+ simple_copy(first, last, output.begin());
+
+ ASSERT_EQ(output[0], 1);
+ ASSERT_EQ(output[1], 2);
+ ASSERT_EQ(output[2], 3);
+ ASSERT_EQ(output[3], 4);
+ ASSERT_EQ(output[4], 5);
+}
+
+TEST(DereferenceTests, TestDeviceDereferenceTransformedCountingIterator)
+{
+ thrust::counting_iterator<int> first(1);
+ thrust::counting_iterator<int> last(6);
+
+ thrust::device_vector<int> output(5);
+
+ simple_copy(thrust::make_transform_iterator(first, thrust::negate<int>()),
+ thrust::make_transform_iterator(last, thrust::negate<int>()),
+ output.begin());
+
+ ASSERT_EQ(output[0], -1);
+ ASSERT_EQ(output[1], -2);
+ ASSERT_EQ(output[2], -3);
+ ASSERT_EQ(output[3], -4);
+ ASSERT_EQ(output[4], -5);
+}

Implemented dereference tests