text: /*
   *  Copyright 2008-2013 NVIDIA Corporation
   *  Modifications Copyright© 2019 Advanced Micro Devices, Inc. All rights reserved.
   *
   *  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.
   */
  
  #include <thrust/device_vector.h>
  #include <thrust/functional.h>
  #include <thrust/host_vector.h>
  #include <thrust/iterator/discard_iterator.h>
  #include <thrust/iterator/retag.h>
  #include <thrust/merge.h>
  #include <thrust/sort.h>
  #include <thrust/unique.h>
  
  #include "test_header.hpp"
  
  TESTS_DEFINE(MergeByKeyTests, FullTestsParams);
  TESTS_DEFINE(PrimitiveMergeByKeyTests, NumericalTestsParams);
  
  TEST(MergeByKeyTests, UsingHip)
  {
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      ASSERT_EQ(THRUST_DEVICE_SYSTEM, THRUST_DEVICE_SYSTEM_HIP);
  }
  
  TYPED_TEST(MergeByKeyTests, MergeByKeySimple)
  {
      using Vector   = typename TestFixture::input_type;
      using Iterator = typename Vector::iterator;
  
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      Vector a_key(3), a_val(3), b_key(4), b_val(4);
  
      a_key[0] = 0;
      a_key[1] = 2;
      a_key[2] = 4;
      a_val[0] = 13;
      a_val[1] = 7;
      a_val[2] = 42;
  
      b_key[0] = 0;
      b_key[1] = 3;
      b_key[2] = 3;
      b_key[3] = 4;
      b_val[0] = 42;
      b_val[1] = 42;
      b_val[2] = 7;
      b_val[3] = 13;
  
      Vector ref_key(7), ref_val(7);
      ref_key[0] = 0;
      ref_val[0] = 13;
      ref_key[1] = 0;
      ref_val[1] = 42;
      ref_key[2] = 2;
      ref_val[2] = 7;
      ref_key[3] = 3;
      ref_val[3] = 42;
      ref_key[4] = 3;
      ref_val[4] = 7;
      ref_key[5] = 4;
      ref_val[5] = 42;
      ref_key[6] = 4;
      ref_val[6] = 13;
  
      Vector result_key(7), result_val(7);
  
      thrust::pair<Iterator, Iterator> ends = thrust::merge_by_key(a_key.begin(),
                                                                   a_key.end(),
                                                                   b_key.begin(),
                                                                   b_key.end(),
                                                                   a_val.begin(),
                                                                   b_val.begin(),
                                                                   result_key.begin(),
                                                                   result_val.begin());
  
      EXPECT_EQ(result_key.end(), ends.first);
      EXPECT_EQ(result_val.end(), ends.second);
      ASSERT_EQ(ref_key, result_key);
      ASSERT_EQ(ref_val, result_val);
  }
  
  template <typename InputIterator1,
            typename InputIterator2,
            typename InputIterator3,
            typename InputIterator4,
            typename OutputIterator1,
            typename OutputIterator2>
  thrust::pair<OutputIterator1, OutputIterator2> merge_by_key(my_system& system,
                                                              InputIterator1,
                                                              InputIterator1,
                                                              InputIterator2,
                                                              InputIterator2,
                                                              InputIterator3,
                                                              InputIterator4,
                                                              OutputIterator1 keys_result,
                                                              OutputIterator2 values_result)
  {
      system.validate_dispatch();
      return thrust::make_pair(keys_result, values_result);
  }
  
  TEST(MergeByKeyTests, MergeByKeyDispatchExplicit)
  {
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      thrust::device_vector<int> vec(1);
  
      my_system sys(0);
      thrust::merge_by_key(sys,
                           vec.begin(),
                           vec.begin(),
                           vec.begin(),
                           vec.begin(),
                           vec.begin(),
                           vec.begin(),
                           vec.begin(),
                           vec.begin());
  
      ASSERT_EQ(true, sys.is_valid());
  }
  
  template <typename InputIterator1,
            typename InputIterator2,
            typename InputIterator3,
            typename InputIterator4,
            typename OutputIterator1,
            typename OutputIterator2>
  thrust::pair<OutputIterator1, OutputIterator2> merge_by_key(my_tag,
                                                              InputIterator1,
                                                              InputIterator1,
                                                              InputIterator2,
                                                              InputIterator2,
                                                              InputIterator3,
                                                              InputIterator4,
                                                              OutputIterator1 keys_result,
                                                              OutputIterator2 values_result)
  {
      *keys_result = 13;
      return thrust::make_pair(keys_result, values_result);
  }
  
  TEST(MergeByKeyTests, MergeByKeyDispatchImplicit)
  {
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      thrust::device_vector<int> vec(1);
  
      thrust::merge_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()),
                           thrust::retag<my_tag>(vec.begin()),
                           thrust::retag<my_tag>(vec.begin()),
                           thrust::retag<my_tag>(vec.begin()));
  
      ASSERT_EQ(13, vec.front());
  }
  
  TYPED_TEST(PrimitiveMergeByKeyTests, TestMergeByKeyWithRandomData)
  {
      using T = typename TestFixture::input_type;
  
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      for(auto size : get_sizes())
      {
          SCOPED_TRACE(testing::Message() << "with size= " << size);
  
          for(auto seed : get_seeds())
          {
              SCOPED_TRACE(testing::Message() << "with seed= " << seed);
  
              thrust::host_vector<T> random_keys = get_random_data<unsigned short int>(
                  size,
                  0,
                  255,
                  seed
              );
              thrust::host_vector<T> random_vals = get_random_data<unsigned short int>(
                  size,
                  0,
                  255,
                  seed + seed_value_addition
              );
  
              size_t denominators[]   = {1, 2, 3, 4, 5, 6, 7, 8, 9};
              size_t num_denominators = sizeof(denominators) / sizeof(size_t);
  
              for(size_t i = 0; i < num_denominators; ++i)
              {
                  size_t size_a = size / denominators[i];
  
                  thrust::host_vector<T> h_a_keys(random_keys.begin(), random_keys.begin() + size_a);
                  thrust::host_vector<T> h_b_keys(random_keys.begin() + size_a, random_keys.end());
  
                  thrust::host_vector<T> h_a_vals(random_vals.begin(), random_vals.begin() + size_a);
                  thrust::host_vector<T> h_b_vals(random_vals.begin() + size_a, random_vals.end());
  
                  thrust::stable_sort(h_a_keys.begin(), h_a_keys.end());
                  thrust::stable_sort(h_b_keys.begin(), h_b_keys.end());
  
                  thrust::device_vector<T> d_a_keys = h_a_keys;
                  thrust::device_vector<T> d_b_keys = h_b_keys;
  
                  thrust::device_vector<T> d_a_vals = h_a_vals;
                  thrust::device_vector<T> d_b_vals = h_b_vals;
  
                  thrust::host_vector<T> h_result_keys(size);
                  thrust::host_vector<T> h_result_vals(size);
  
                  thrust::device_vector<T> d_result_keys(size);
                  thrust::device_vector<T> d_result_vals(size);
  
                  thrust::pair<typename thrust::host_vector<T>::iterator,
                               typename thrust::host_vector<T>::iterator>
                      h_end;
  
                  thrust::pair<typename thrust::device_vector<T>::iterator,
                               typename thrust::device_vector<T>::iterator>
                      d_end;
  
                  h_end = thrust::merge_by_key(h_a_keys.begin(),
                                               h_a_keys.end(),
                                               h_b_keys.begin(),
                                               h_b_keys.end(),
                                               h_a_vals.begin(),
                                               h_b_vals.begin(),
                                               h_result_keys.begin(),
                                               h_result_vals.begin());
                  h_result_keys.erase(h_end.first, h_result_keys.end());
                  h_result_vals.erase(h_end.second, h_result_vals.end());
  
                  d_end = thrust::merge_by_key(d_a_keys.begin(),
                                               d_a_keys.end(),
                                               d_b_keys.begin(),
                                               d_b_keys.end(),
                                               d_a_vals.begin(),
                                               d_b_vals.begin(),
                                               d_result_keys.begin(),
                                               d_result_vals.begin());
                  d_result_keys.erase(d_end.first, d_result_keys.end());
                  d_result_vals.erase(d_end.second, d_result_vals.end());
  
                  thrust::host_vector<T> d_result_keys_h = d_result_keys;
                  thrust::host_vector<T> d_result_vals_h = d_result_vals;
  
                  ASSERT_EQ(h_result_keys, d_result_keys_h);
                  ASSERT_EQ(h_result_vals, d_result_vals_h);
              }
          }
      }
  }
  
  TYPED_TEST(PrimitiveMergeByKeyTests, MergeByKeyToDiscardIterator)
  {
      using T = typename TestFixture::input_type;
  
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      for(auto size : get_sizes())
      {
          SCOPED_TRACE(testing::Message() << "with size= " << size);
  
          for(auto seed : get_seeds())
          {
              SCOPED_TRACE(testing::Message() << "with seed= " << seed);
  
              thrust::host_vector<T> h_a_keys = get_random_data<T>(
                  size,
                  std::numeric_limits<T>::min(),
                  std::numeric_limits<T>::max(),
                  seed
              );
              thrust::host_vector<T> h_b_keys = get_random_data<T>(
                  size,
                  std::numeric_limits<T>::min(),
                  std::numeric_limits<T>::max(),
                  seed + seed_value_addition
              );
  
              thrust::host_vector<T> h_a_vals = get_random_data<T>(
                  size,
                  std::numeric_limits<T>::min(),
                  std::numeric_limits<T>::max(),
                  seed + 2 * seed_value_addition
              );
              thrust::host_vector<T> h_b_vals = get_random_data<T>(
                  size,
                  std::numeric_limits<T>::min(),
                  std::numeric_limits<T>::max(),
                  seed + 3 * seed_value_addition
              );
  
              thrust::stable_sort(h_a_keys.begin(), h_a_keys.end());
              thrust::stable_sort(h_b_keys.begin(), h_b_keys.end());
  
              thrust::device_vector<T> d_a_keys = h_a_keys;
              thrust::device_vector<T> d_b_keys = h_b_keys;
  
              thrust::device_vector<T> d_a_vals = h_a_vals;
              thrust::device_vector<T> d_b_vals = h_b_vals;
  
              typedef thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<>>
                  discard_pair;
  
              discard_pair h_result = thrust::merge_by_key(h_a_keys.begin(),
                                                           h_a_keys.end(),
                                                           h_b_keys.begin(),
                                                           h_b_keys.end(),
                                                           h_a_vals.begin(),
                                                           h_b_vals.begin(),
                                                           thrust::make_discard_iterator(),
                                                           thrust::make_discard_iterator());
  
              discard_pair d_result = thrust::merge_by_key(d_a_keys.begin(),
                                                           d_a_keys.end(),
                                                           d_b_keys.begin(),
                                                           d_b_keys.end(),
                                                           d_a_vals.begin(),
                                                           d_b_vals.begin(),
                                                           thrust::make_discard_iterator(),
                                                           thrust::make_discard_iterator());
  
              thrust::discard_iterator<> reference(2 * size);
  
              ASSERT_EQ(reference, h_result.first);
              ASSERT_EQ(reference, h_result.second);
              ASSERT_EQ(reference, d_result.first);
              ASSERT_EQ(reference, d_result.second);
          }
      }
  }
  
  TYPED_TEST(PrimitiveMergeByKeyTests, MergeByKeyDescending)
  {
      using T = typename TestFixture::input_type;
  
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      for(auto size : get_sizes())
      {
          SCOPED_TRACE(testing::Message() << "with size= " << size);
  
          for(auto seed : get_seeds())
          {
              SCOPED_TRACE(testing::Message() << "with seed= " << seed);
  
              thrust::host_vector<T> random_keys = get_random_data<unsigned short int>(
                  size,
                  0,
                  255,
                  seed
              );
              thrust::host_vector<T> random_vals = get_random_data<unsigned short int>(
                  size,
                  0,
                  255,
                  seed + seed_value_addition
              );
  
              size_t denominators[]   = {1, 2, 3, 4, 5, 6, 7, 8, 9};
              size_t num_denominators = sizeof(denominators) / sizeof(size_t);
  
              for(size_t i = 0; i < num_denominators; ++i)
              {
                  size_t size_a = size / denominators[i];
  
                  thrust::host_vector<T> h_a_keys(random_keys.begin(), random_keys.begin() + size_a);
                  thrust::host_vector<T> h_b_keys(random_keys.begin() + size_a, random_keys.end());
  
                  thrust::host_vector<T> h_a_vals(random_vals.begin(), random_vals.begin() + size_a);
                  thrust::host_vector<T> h_b_vals(random_vals.begin() + size_a, random_vals.end());
  
                  thrust::stable_sort(h_a_keys.begin(), h_a_keys.end(), thrust::greater<T>());
                  thrust::stable_sort(h_b_keys.begin(), h_b_keys.end(), thrust::greater<T>());
  
                  thrust::device_vector<T> d_a_keys = h_a_keys;
                  thrust::device_vector<T> d_b_keys = h_b_keys;
  
                  thrust::device_vector<T> d_a_vals = h_a_vals;
                  thrust::device_vector<T> d_b_vals = h_b_vals;
  
                  thrust::host_vector<T> h_result_keys(size);
                  thrust::host_vector<T> h_result_vals(size);
  
                  thrust::device_vector<T> d_result_keys(size);
                  thrust::device_vector<T> d_result_vals(size);
  
                  thrust::pair<typename thrust::host_vector<T>::iterator,
                               typename thrust::host_vector<T>::iterator>
                      h_end;
  
                  thrust::pair<typename thrust::device_vector<T>::iterator,
                               typename thrust::device_vector<T>::iterator>
                      d_end;
  
                  h_end = thrust::merge_by_key(h_a_keys.begin(),
                                               h_a_keys.end(),
                                               h_b_keys.begin(),
                                               h_b_keys.end(),
                                               h_a_vals.begin(),
                                               h_b_vals.begin(),
                                               h_result_keys.begin(),
                                               h_result_vals.begin(),
                                               thrust::greater<T>());
                  h_result_keys.erase(h_end.first, h_result_keys.end());
                  h_result_vals.erase(h_end.second, h_result_vals.end());
  
                  d_end = thrust::merge_by_key(d_a_keys.begin(),
                                               d_a_keys.end(),
                                               d_b_keys.begin(),
                                               d_b_keys.end(),
                                               d_a_vals.begin(),
                                               d_b_vals.begin(),
                                               d_result_keys.begin(),
                                               d_result_vals.begin(),
                                               thrust::greater<T>());
                  d_result_keys.erase(d_end.first, d_result_keys.end());
                  d_result_vals.erase(d_end.second, d_result_vals.end());
  
                  ASSERT_EQ(h_result_keys, d_result_keys);
                  ASSERT_EQ(h_result_vals, d_result_vals);
              }
          }
      }
  }
  
  
  
  template<class T>
  __global__
  THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
  void MergeByKeyKernel(int const N, T* keys_A, T* keys_B, T * values_A, T* values_B, T* keys_result, T* values_result)
  {
      if(threadIdx.x == 0)
      {
          thrust::device_ptr<int> keyA_begin(keys_A);
          thrust::device_ptr<int> keyA_end(keys_A + N);
          thrust::device_ptr<int> keyB_begin(keys_B);
          thrust::device_ptr<int> keyB_end(keys_B + N);
          thrust::device_ptr<int> valuesA_begin(values_A);
          thrust::device_ptr<int> valuesB_begin(values_B);
          thrust::device_ptr<int> keys_result_begin(keys_result);
          thrust::device_ptr<int> values_result_begin(values_result);
  
  
          thrust::merge_by_key(thrust::hip::par, keyA_begin,           keyA_end,
                                                 keyB_begin,           keyB_end,
                                                 valuesA_begin,        valuesB_begin,
                                                 keys_result_begin,
                                                 values_result_begin);
  
      }
  }
  TEST(MergeByKeyTests, TestMergeByKeyDevice)
  {
      using T = int;
  
      SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  
      for(auto size : get_sizes())
      {
          SCOPED_TRACE(testing::Message() << "with size= " << size);
  
          for(auto seed : get_seeds())
          {
              SCOPED_TRACE(testing::Message() << "with seed= " << seed);
  
              thrust::host_vector<T> h_keys_a = get_random_data<T>(
                  size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed);
              thrust::host_vector<T> h_keys_b = get_random_data<T>(
                  size,
                  std::numeric_limits<T>::min(),
                  std::numeric_limits<T>::max(),
                  seed + seed_value_addition
              );
  
              thrust::host_vector<T> h_values_a = get_random_data<T>(
                  size, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed);
              thrust::host_vector<T> h_values_b = get_random_data<T>(
                  size,
                  std::numeric_limits<T>::min(),
                  std::numeric_limits<T>::max(),
                  seed + seed_value_addition
              );
  
  
              thrust::stable_sort(h_keys_a.begin(), h_keys_a.end(), thrust::greater<T>());
              thrust::stable_sort(h_keys_b.begin(), h_keys_b.end(), thrust::greater<T>());
  
              thrust::device_vector<T> d_keys_a = h_keys_a;
              thrust::device_vector<T> d_keys_b = h_keys_b;
  
              thrust::device_vector<T> d_values_a = h_values_a;
              thrust::device_vector<T> d_values_b = h_values_b;
  
              thrust::host_vector<T>   h_keys_result(h_keys_a.size() + h_keys_b.size());
              thrust::device_vector<T> d_keys_result(d_keys_a.size() + d_keys_b.size());
  
              thrust::host_vector<T>   h_values_result(h_values_a.size() + h_values_b.size());
              thrust::device_vector<T> d_values_result(d_values_a.size() + d_values_b.size());
  
              typename thrust::host_vector<T>::iterator   h_end;
              typename thrust::device_vector<T>::iterator d_end;
  
              thrust::merge_by_key(h_keys_a.begin(),h_keys_a.end(),
                                           h_keys_b.begin(),h_keys_b.end(),
                                           h_values_a.begin(),h_values_b.begin(),
                                           h_keys_result.begin(),
                                           h_values_result.begin());
  
                hipLaunchKernelGGL(MergeByKeyKernel,
                                   dim3(1, 1, 1),
                                   dim3(128, 1, 1),
                                   0,
                                   0,
                                   size,
                                   thrust::raw_pointer_cast(&d_keys_a[0]),
                                   thrust::raw_pointer_cast(&d_keys_b[0]),
                                   thrust::raw_pointer_cast(&d_values_a[0]),
                                   thrust::raw_pointer_cast(&d_values_b[0]),
                                   thrust::raw_pointer_cast(&d_keys_result[0]),
                                   thrust::raw_pointer_cast(&d_values_result[0]));
  
              ASSERT_EQ(h_keys_result, d_keys_result);
              ASSERT_EQ(h_values_result, d_values_result);
  
          }
      }
  }
  