TheRock/rocm-libraries/projects/rocthrust/test/test_remove.cpp

/*
 *  Copyright 2008-2013 NVIDIA Corporation
 *  Modifications Copyright© 2019-2025 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/count.h>
#include <thrust/functional.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/retag.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/remove.h>

#include <stdexcept>

#include "test_param_fixtures.hpp"
#include "test_real_assertions.hpp"
#include "test_utils.hpp"

TESTS_DEFINE(RemoveTests, FullTestsParams);

TESTS_DEFINE(RemoveVariableTests, NumericalTestsParams);

template <typename T>
struct is_even
{
  THRUST_HOST_DEVICE bool operator()(T x)
  {
    return (static_cast<unsigned int>(x) & 1) == 0;
  }
};

template <typename T>
struct is_true
{
  THRUST_HOST_DEVICE bool operator()(T x)
  {
    return x ? true : false;
  }
};

TYPED_TEST(RemoveTests, TestRemoveSimple)
{
  using Vector = typename TestFixture::input_type;
  using T      = typename Vector::value_type;

  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  Vector data{1, 2, 1, 3, 2};

  typename Vector::iterator end = thrust::remove(data.begin(), data.end(), (T) 2);

  ASSERT_EQ(end - data.begin(), 3);
  data.resize(end - data.begin());

  Vector ref{1, 1, 3};
  ASSERT_EQ(data, ref);
}

template <typename ForwardIterator, typename T>
ForwardIterator remove(my_system& system, ForwardIterator first, ForwardIterator, const T&)
{
  system.validate_dispatch();
  return first;
}

TEST(RemoveTests, TestRemoveDispatchExplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  my_system sys(0);
  thrust::remove(sys, vec.begin(), vec.end(), 0);

  ASSERT_EQ(true, sys.is_valid());
}

template <typename ForwardIterator, typename T>
ForwardIterator remove(my_tag, ForwardIterator first, ForwardIterator, const T&)
{
  *first = 13;
  return first;
}

TEST(RemoveTests, TestRemoveDispatchImplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
  thrust::device_vector<int> vec(1);

  thrust::remove(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), 0);

  ASSERT_EQ(13, vec.front());
}

TYPED_TEST(RemoveTests, TestRemoveCopySimple)
{
  using Vector = typename TestFixture::input_type;
  using T      = typename Vector::value_type;

  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  Vector data{1, 2, 1, 3, 2};

  Vector result(5);

  typename Vector::iterator end = thrust::remove_copy(data.begin(), data.end(), result.begin(), (T) 2);

  ASSERT_EQ(end - result.begin(), 3);
  result.resize(end - result.begin());

  Vector ref{1, 1, 3};
  ASSERT_EQ(result, ref);
}

template <typename InputIterator, typename OutputIterator, typename T>
OutputIterator remove_copy(my_system& system, InputIterator, InputIterator, OutputIterator result, const T&)
{
  system.validate_dispatch();
  return result;
}

TEST(RemoveTests, TestRemoveCopyDispatchExplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  my_system sys(0);
  thrust::remove_copy(sys, vec.begin(), vec.begin(), vec.begin(), 0);

  ASSERT_EQ(true, sys.is_valid());
}

template <typename InputIterator, typename OutputIterator, typename T>
OutputIterator remove_copy(my_tag, InputIterator, InputIterator, OutputIterator result, const T&)
{
  *result = 13;
  return result;
}

TEST(RemoveTests, TestRemoveCopyDispatchImplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  thrust::remove_copy(
    thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), 0);

  ASSERT_EQ(13, vec.front());
}

TYPED_TEST(RemoveTests, TestRemoveIfSimple)
{
  using Vector = typename TestFixture::input_type;
  using T      = typename Vector::value_type;

  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  Vector data{1, 2, 1, 3, 2};

  typename Vector::iterator end = thrust::remove_if(data.begin(), data.end(), is_even<T>());

  ASSERT_EQ(end - data.begin(), 3);
  data.resize(end - data.begin());

  Vector ref{1, 1, 3};
  ASSERT_EQ(data, ref);
}

template <typename ForwardIterator, typename Predicate>
ForwardIterator remove_if(my_system& system, ForwardIterator first, ForwardIterator, Predicate)
{
  system.validate_dispatch();
  return first;
}

TEST(RemoveTests, TestRemoveIfDispatchExplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  my_system sys(0);
  thrust::remove_if(sys, vec.begin(), vec.end(), 0);

  ASSERT_EQ(true, sys.is_valid());
}

template <typename ForwardIterator, typename Predicate>
ForwardIterator remove_if(my_tag, ForwardIterator first, ForwardIterator, Predicate)
{
  *first = 13;
  return first;
}

TEST(RemoveTests, TestRemoveIfDispatchImplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  thrust::remove_if(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), 0);

  ASSERT_EQ(13, vec.front());
}

TYPED_TEST(RemoveTests, TestRemoveIfStencilSimple)
{
  using Vector = typename TestFixture::input_type;

  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  Vector data{1, 2, 1, 3, 2};
  Vector stencil{0, 1, 0, 0, 1};

  typename Vector::iterator end = thrust::remove_if(data.begin(), data.end(), stencil.begin(), ::internal::identity{});

  ASSERT_EQ(end - data.begin(), 3);
  data.resize(end - data.begin());

  Vector ref{1, 1, 3};
  ASSERT_EQ(data, ref);
}

template <typename ForwardIterator, typename InputIterator, typename Predicate>
ForwardIterator remove_if(my_system& system, ForwardIterator first, ForwardIterator, InputIterator, Predicate)
{
  system.validate_dispatch();
  return first;
}

TEST(RemoveTests, TestRemoveIfStencilDispatchExplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  my_system sys(0);
  thrust::remove_if(sys, vec.begin(), vec.begin(), vec.begin(), 0);

  ASSERT_EQ(true, sys.is_valid());
}

template <typename ForwardIterator, typename InputIterator, typename Predicate>
ForwardIterator remove_if(my_tag, ForwardIterator first, ForwardIterator, InputIterator, Predicate)
{
  *first = 13;
  return first;
}

TEST(RemoveTests, TestRemoveIfStencilDispatchImplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  thrust::remove_if(
    thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), 0);

  ASSERT_EQ(13, vec.front());
}

TYPED_TEST(RemoveTests, TestRemoveCopyIfSimple)
{
  using Vector = typename TestFixture::input_type;
  using T      = typename Vector::value_type;

  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  Vector data{1, 2, 1, 3, 2};

  Vector result(5);

  typename Vector::iterator end = thrust::remove_copy_if(data.begin(), data.end(), result.begin(), is_even<T>());

  ASSERT_EQ(end - result.begin(), 3);
  result.resize(end - result.begin());

  Vector ref{1, 1, 3};
  ASSERT_EQ(result, ref);
}

template <typename InputIterator, typename OutputIterator, typename Predicate>
InputIterator remove_copy_if(my_system& system, InputIterator first, InputIterator, OutputIterator, Predicate)
{
  system.validate_dispatch();
  return first;
}

TEST(RemoveTests, TestRemoveCopyIfDispatchExplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  my_system sys(0);
  thrust::remove_copy_if(sys, vec.begin(), vec.begin(), vec.begin(), 0);

  ASSERT_EQ(true, sys.is_valid());
}

template <typename InputIterator, typename OutputIterator, typename Predicate>
InputIterator remove_copy_if(my_tag, InputIterator first, InputIterator, OutputIterator, Predicate)
{
  *first = 13;
  return first;
}

TEST(RemoveTests, TestRemoveCopyIfDispatchImplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  thrust::remove_copy_if(
    thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), 0);

  ASSERT_EQ(13, vec.front());
}

TYPED_TEST(RemoveTests, TestRemoveCopyIfStencilSimple)
{
  using Vector = typename TestFixture::input_type;

  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  Vector data{1, 2, 1, 3, 2};
  Vector stencil{0, 1, 0, 0, 1};

  Vector result(5);

  typename Vector::iterator end =
    thrust::remove_copy_if(data.begin(), data.end(), stencil.begin(), result.begin(), ::internal::identity{});

  ASSERT_EQ(end - result.begin(), 3);
  result.resize(end - result.begin());

  Vector ref{1, 1, 3};
  ASSERT_EQ(result, ref);
}

template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename Predicate>
OutputIterator
remove_copy_if(my_system& system, InputIterator1, InputIterator1, InputIterator2, OutputIterator result, Predicate)
{
  system.validate_dispatch();
  return result;
}

TEST(RemoveTests, TestRemoveCopyIfStencilDispatchExplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  my_system sys(0);
  thrust::remove_copy_if(sys, vec.begin(), vec.begin(), vec.begin(), vec.begin(), 0);

  ASSERT_EQ(true, sys.is_valid());
}

template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename Predicate>
OutputIterator remove_copy_if(my_tag, InputIterator1, InputIterator1, InputIterator2, OutputIterator result, Predicate)
{
  *result = 13;
  return result;
}

TEST(RemoveTests, TestRemoveCopyIfStencilDispatchImplicit)
{
  SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());

  thrust::device_vector<int> vec(1);

  thrust::remove_copy_if(
    thrust::retag<my_tag>(vec.begin()),
    thrust::retag<my_tag>(vec.begin()),
    thrust::retag<my_tag>(vec.begin()),
    thrust::retag<my_tag>(vec.begin()),
    0);

  ASSERT_EQ(13, vec.front());
}

TYPED_TEST(RemoveVariableTests, TestRemove)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      size_t h_size = thrust::remove(h_data.begin(), h_data.end(), T(0)) - h_data.begin();
      size_t d_size = thrust::remove(d_data.begin(), d_data.end(), T(0)) - d_data.begin();

      ASSERT_EQ(h_size, d_size);

      h_data.resize(h_size);
      d_data.resize(d_size);

      ASSERT_EQ(h_data, d_data);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveIf)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), is_true<T>()) - h_data.begin();
      size_t d_size = thrust::remove_if(d_data.begin(), d_data.end(), is_true<T>()) - d_data.begin();

      ASSERT_EQ(h_size, d_size);

      h_data.resize(h_size);
      d_data.resize(d_size);

      ASSERT_EQ(h_data, d_data);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveIfStencil)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      thrust::host_vector<bool> h_stencil = get_random_data<bool>(
        size, std::numeric_limits<bool>::min(), std::numeric_limits<bool>::max(), seed + seed_value_addition);
      thrust::device_vector<bool> d_stencil = h_stencil;

      size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), h_stencil.begin(), is_true<T>()) - h_data.begin();
      size_t d_size = thrust::remove_if(d_data.begin(), d_data.end(), d_stencil.begin(), is_true<T>()) - d_data.begin();

      ASSERT_EQ(h_size, d_size);

      h_data.resize(h_size);
      d_data.resize(d_size);

      ASSERT_EQ(h_data, d_data);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopy)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      thrust::host_vector<T> h_result(size);
      thrust::device_vector<T> d_result(size);

      size_t h_size = thrust::remove_copy(h_data.begin(), h_data.end(), h_result.begin(), T(0)) - h_result.begin();
      size_t d_size = thrust::remove_copy(d_data.begin(), d_data.end(), d_result.begin(), T(0)) - d_result.begin();

      ASSERT_EQ(h_size, d_size);

      h_result.resize(h_size);
      d_result.resize(d_size);

      ASSERT_EQ(h_result, d_result);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopyToDiscardIterator)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      size_t num_zeros    = thrust::count(h_data.begin(), h_data.end(), T(0));
      size_t num_nonzeros = h_data.size() - num_zeros;

      thrust::discard_iterator<> h_result =
        thrust::remove_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), T(0));

      thrust::discard_iterator<> d_result =
        thrust::remove_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), T(0));

      thrust::discard_iterator<> reference(num_nonzeros);

      ASSERT_EQ_QUIET(reference, h_result);
      ASSERT_EQ_QUIET(reference, d_result);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopyToDiscardIteratorZipped)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      thrust::host_vector<T> h_output(size);
      thrust::device_vector<T> d_output(size);

      size_t num_zeros    = thrust::count(h_data.begin(), h_data.end(), T(0));
      size_t num_nonzeros = h_data.size() - num_zeros;

      using Tuple1 = thrust::tuple<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>>;
      using Tuple2 = thrust::tuple<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>>;

      using ZipIterator1 = thrust::zip_iterator<Tuple1>;
      using ZipIterator2 = thrust::zip_iterator<Tuple2>;

      ZipIterator1 h_result = thrust::remove_copy(
        thrust::make_zip_iterator(thrust::make_tuple(h_data.begin(), h_data.begin())),
        thrust::make_zip_iterator(thrust::make_tuple(h_data.end(), h_data.end())),
        thrust::make_zip_iterator(thrust::make_tuple(h_output.begin(), thrust::make_discard_iterator())),
        thrust::make_tuple(T(0), T(0)));

      ZipIterator2 d_result = thrust::remove_copy(
        thrust::make_zip_iterator(thrust::make_tuple(d_data.begin(), d_data.begin())),
        thrust::make_zip_iterator(thrust::make_tuple(d_data.end(), d_data.end())),
        thrust::make_zip_iterator(thrust::make_tuple(d_output.begin(), thrust::make_discard_iterator())),
        thrust::make_tuple(T(0), T(0)));

      thrust::discard_iterator<> reference(num_nonzeros);

      ASSERT_EQ(h_output, d_output);
      ASSERT_EQ_QUIET(reference, thrust::get<1>(h_result.get_iterator_tuple()));
      ASSERT_EQ_QUIET(reference, thrust::get<1>(d_result.get_iterator_tuple()));
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopyIf)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      thrust::host_vector<T> h_result(size);
      thrust::device_vector<T> d_result(size);

      size_t h_size =
        thrust::remove_copy_if(h_data.begin(), h_data.end(), h_result.begin(), is_true<T>()) - h_result.begin();
      size_t d_size =
        thrust::remove_copy_if(d_data.begin(), d_data.end(), d_result.begin(), is_true<T>()) - d_result.begin();

      ASSERT_EQ(h_size, d_size);

      h_result.resize(h_size);
      d_result.resize(d_size);

      ASSERT_EQ(h_result, d_result);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopyIfToDiscardIterator)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      size_t num_false = thrust::count_if(h_data.begin(), h_data.end(), thrust::not_fn(is_true<T>()));

      thrust::discard_iterator<> h_result =
        thrust::remove_copy_if(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), is_true<T>());

      thrust::discard_iterator<> d_result =
        thrust::remove_copy_if(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), is_true<T>());

      thrust::discard_iterator<> reference(num_false);

      ASSERT_EQ_QUIET(reference, h_result);
      ASSERT_EQ_QUIET(reference, d_result);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopyIfStencil)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      thrust::host_vector<bool> h_stencil =
        get_random_data<bool>(size, get_default_limits<bool>::min(), get_default_limits<bool>::max(), seed);
      thrust::device_vector<bool> d_stencil = h_stencil;

      thrust::host_vector<T> h_result(size);
      thrust::device_vector<T> d_result(size);

      size_t h_size =
        thrust::remove_copy_if(h_data.begin(), h_data.end(), h_stencil.begin(), h_result.begin(), is_true<T>())
        - h_result.begin();
      size_t d_size =
        thrust::remove_copy_if(d_data.begin(), d_data.end(), d_stencil.begin(), d_result.begin(), is_true<T>())
        - d_result.begin();

      ASSERT_EQ(h_size, d_size);

      h_result.resize(h_size);
      d_result.resize(d_size);

      ASSERT_EQ(h_result, d_result);
    }
  }
}

TYPED_TEST(RemoveVariableTests, TestRemoveCopyIfStencilToDiscardIterator)
{
  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_data =
        get_random_data<T>(size, get_default_limits<T>::min(), get_default_limits<T>::max(), seed);
      thrust::device_vector<T> d_data = h_data;

      thrust::host_vector<bool> h_stencil = get_random_data<bool>(
        size, get_default_limits<bool>::min(), get_default_limits<bool>::max(), seed + seed_value_addition);
      thrust::device_vector<bool> d_stencil = h_stencil;

      size_t num_false = thrust::count_if(h_stencil.begin(), h_stencil.end(), thrust::not_fn(is_true<T>()));

      thrust::discard_iterator<> h_result = thrust::remove_copy_if(
        h_data.begin(), h_data.end(), h_stencil.begin(), thrust::make_discard_iterator(), is_true<T>());

      thrust::discard_iterator<> d_result = thrust::remove_copy_if(
        d_data.begin(), d_data.end(), d_stencil.begin(), thrust::make_discard_iterator(), is_true<T>());

      thrust::discard_iterator<> reference(num_false);

      ASSERT_EQ_QUIET(reference, h_result);
      ASSERT_EQ_QUIET(reference, d_result);
    }
  }
}

__global__ THRUST_HIP_LAUNCH_BOUNDS_DEFAULT void RemoveKernel(int const N, int* array, int remove_value)
{
  if (threadIdx.x == 0)
  {
    thrust::device_ptr<int> begin(array);
    thrust::device_ptr<int> end(array + N);
    thrust::remove(thrust::hip::par, begin, end, remove_value) - begin;
  }
}

TEST(RemoveTests, TestRemoveDevice)
{
  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<int> h_data   = get_random_data<int>(size, 0, size, seed);
      thrust::device_vector<int> d_data = h_data;

      int remove_value = get_random_data<int>(1, 0, size, seed)[0];
      SCOPED_TRACE(testing::Message() << "with remove_value= " << remove_value);

      thrust::remove(h_data.begin(), h_data.end(), remove_value);
      hipLaunchKernelGGL(
        RemoveKernel, dim3(1, 1, 1), dim3(128, 1, 1), 0, 0, size, thrust::raw_pointer_cast(&d_data[0]), remove_value);
      HIP_CHECK(hipGetLastError());

      ASSERT_EQ(h_data, d_data);
    }
  }
}