Add tests to test_reduce_by_key
MODIFY
test/test_reduce_by_key.cpp
test/test_reduce_by_key.cpp
// Thrust
#include <thrust/reduce.h>
#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/retag.h>
// HIP API
@@ -186,5 +187,161 @@ TYPED_TEST(ReduceByKeysTests, TestReduceByKeySimple)
ASSERT_EQ(output_values[4], 15);
}
+TYPED_TEST(ReduceByKeysTests, TestReduceByKey)
+{
+ using Vector = typename TestFixture::input_type;
+ using V = typename Vector::value_type; // value type
+ typedef unsigned int K; // key type
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<K> h_keys = get_random_data<K>(size,
+ std::numeric_limits<K>::min(),
+ std::numeric_limits<K>::max());
+ thrust::host_vector<V> h_vals = get_random_data<V>(size,
+ std::numeric_limits<V>::min(),
+ std::numeric_limits<V>::max());
+ thrust::device_vector<K> d_keys = h_keys;
+ thrust::device_vector<V> d_vals = h_vals;
+
+ thrust::host_vector<K> h_keys_output(size);
+ thrust::host_vector<V> h_vals_output(size);
+ thrust::device_vector<K> d_keys_output(size);
+ thrust::device_vector<V> d_vals_output(size);
+
+ typedef typename thrust::host_vector<K>::iterator HostKeyIterator;
+ typedef typename thrust::host_vector<V>::iterator HostValIterator;
+ typedef typename thrust::device_vector<K>::iterator DeviceKeyIterator;
+ typedef typename thrust::device_vector<V>::iterator DeviceValIterator;
+
+ typedef typename thrust::pair<HostKeyIterator, HostValIterator> HostIteratorPair;
+ typedef typename thrust::pair<DeviceKeyIterator,DeviceValIterator> DeviceIteratorPair;
+
+ HostIteratorPair h_last = thrust::reduce_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_keys_output.begin(), h_vals_output.begin());
+ DeviceIteratorPair d_last = thrust::reduce_by_key(d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys_output.begin(), d_vals_output.begin());
+
+ ASSERT_EQ(h_last.first - h_keys_output.begin(), d_last.first - d_keys_output.begin());
+ ASSERT_EQ(h_last.second - h_vals_output.begin(), d_last.second - d_vals_output.begin());
+
+ size_t N = h_last.first - h_keys_output.begin();
+
+ h_keys_output.resize(N);
+ h_vals_output.resize(N);
+ d_keys_output.resize(N);
+ d_vals_output.resize(N);
+
+ ASSERT_EQ(h_keys_output, d_keys_output);
+ ASSERT_EQ(h_vals_output, d_vals_output);
+ }
+};
+
+TYPED_TEST(ReduceByKeysTests, TestReduceByKeyToDiscardIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using V = typename Vector::value_type; // value type
+ typedef unsigned int K; // key type
+
+ const std::vector<size_t> sizes = get_sizes();
+ for(auto size : sizes)
+ {
+ thrust::host_vector<K> h_keys = get_random_data<K>(size,
+ std::numeric_limits<K>::min(),
+ std::numeric_limits<K>::max());
+ thrust::host_vector<V> h_vals = get_random_data<V>(size,
+ std::numeric_limits<V>::min(),
+ std::numeric_limits<V>::max());
+ thrust::device_vector<K> d_keys = h_keys;
+ thrust::device_vector<V> d_vals = h_vals;
+
+ thrust::host_vector<K> h_keys_output(size);
+ thrust::host_vector<V> h_vals_output(size);
+ thrust::device_vector<K> d_keys_output(size);
+ thrust::device_vector<V> d_vals_output(size);
+
+ thrust::host_vector<K> unique_keys = h_keys;
+ unique_keys.erase(thrust::unique(unique_keys.begin(), unique_keys.end()), unique_keys.end());
+
+ // discard key output
+ size_t h_size =
+ thrust::reduce_by_key(h_keys.begin(), h_keys.end(),
+ h_vals.begin(),
+ thrust::make_discard_iterator(),
+ h_vals_output.begin()).second - h_vals_output.begin();
+
+ size_t d_size =
+ thrust::reduce_by_key(d_keys.begin(), d_keys.end(),
+ d_vals.begin(),
+ thrust::make_discard_iterator(),
+ d_vals_output.begin()).second - d_vals_output.begin();
+
+ h_vals_output.resize(h_size);
+ d_vals_output.resize(d_size);
+
+ ASSERT_EQ(h_vals_output.size(), unique_keys.size());
+ ASSERT_EQ(d_vals_output.size(), unique_keys.size());
+ ASSERT_EQ(d_vals_output.size(), h_vals_output.size());
+ }
+};
+
+template<typename InputIterator1,
+typename InputIterator2,
+typename OutputIterator1,
+typename OutputIterator2>
+thrust::pair<OutputIterator1,OutputIterator2>
+reduce_by_key(my_system &system,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ OutputIterator1 keys_output,
+ OutputIterator2 values_output)
+{
+ system.validate_dispatch();
+ return thrust::make_pair(keys_output, values_output);
+}
+
+TYPED_TEST(ReduceByKeysTests, TestReduceByKeyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::reduce_by_key(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename InputIterator1,
+typename InputIterator2,
+typename OutputIterator1,
+typename OutputIterator2>
+thrust::pair<OutputIterator1,OutputIterator2>
+reduce_by_key(my_tag,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ OutputIterator1 keys_output,
+ OutputIterator2 values_output)
+{
+ *keys_output = 13;
+ return thrust::make_pair(keys_output, values_output);
+}
+
+TYPED_TEST(ReduceByKeysTests, TestReduceByKeyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::reduce_by_key(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

Add tests to test_reduce_by_key