CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
660 lines
23 KiB
C++
660 lines
23 KiB
C++
/*
|
|
* 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/functional.h>
|
|
#include <thrust/iterator/discard_iterator.h>
|
|
#include <thrust/iterator/retag.h>
|
|
#include <thrust/unique.h>
|
|
|
|
#include "test_param_fixtures.hpp"
|
|
#include "test_real_assertions.hpp"
|
|
#include "test_utils.hpp"
|
|
|
|
using IntegralTypes = ::testing::Types<
|
|
Params<char>,
|
|
Params<signed char>,
|
|
Params<unsigned char>,
|
|
Params<short>,
|
|
Params<unsigned short>,
|
|
Params<int>,
|
|
Params<unsigned int>,
|
|
Params<long>,
|
|
Params<unsigned long>,
|
|
Params<long long>,
|
|
Params<unsigned long long>>;
|
|
|
|
TESTS_DEFINE(UniqueByKeyTests, FullTestsParams);
|
|
|
|
TESTS_DEFINE(UniqueByKeyIntegralTests, IntegralTypes);
|
|
|
|
template <typename ValueT>
|
|
struct index_to_value_t
|
|
{
|
|
template <typename IndexT>
|
|
THRUST_HOST_DEVICE THRUST_FORCEINLINE ValueT operator()(IndexT index)
|
|
{
|
|
if (static_cast<std::uint64_t>(index) == 4300000000ULL)
|
|
{
|
|
return static_cast<ValueT>(1);
|
|
}
|
|
else
|
|
{
|
|
return static_cast<ValueT>(0);
|
|
}
|
|
}
|
|
};
|
|
|
|
template <typename ForwardIterator1, typename ForwardIterator2>
|
|
thrust::pair<ForwardIterator1, ForwardIterator2>
|
|
unique_by_key(my_system& system, ForwardIterator1 keys_first, ForwardIterator1, ForwardIterator2 values_first)
|
|
{
|
|
system.validate_dispatch();
|
|
return thrust::make_pair(keys_first, values_first);
|
|
}
|
|
|
|
TEST(UniqueByKeyTests, TestUniqueByKeyDispatchExplicit)
|
|
{
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
thrust::device_vector<int> vec(1);
|
|
|
|
my_system sys(0);
|
|
thrust::unique_by_key(sys, vec.begin(), vec.begin(), vec.begin());
|
|
|
|
ASSERT_EQ(true, sys.is_valid());
|
|
}
|
|
|
|
template <typename ForwardIterator1, typename ForwardIterator2>
|
|
thrust::pair<ForwardIterator1, ForwardIterator2>
|
|
unique_by_key(my_tag, ForwardIterator1 keys_first, ForwardIterator1, ForwardIterator2 values_first)
|
|
{
|
|
*keys_first = 13;
|
|
return thrust::make_pair(keys_first, values_first);
|
|
}
|
|
|
|
TEST(UniqueByKeyTests, TestUniqueByKeyDispatchImplicit)
|
|
{
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
thrust::device_vector<int> vec(1);
|
|
|
|
thrust::unique_by_key(
|
|
thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()));
|
|
|
|
ASSERT_EQ(13, vec.front());
|
|
}
|
|
|
|
template <typename InputIterator1, typename InputIterator2, typename OutputIterator1, typename OutputIterator2>
|
|
thrust::pair<OutputIterator1, OutputIterator2> unique_by_key_copy(
|
|
my_system& system,
|
|
InputIterator1,
|
|
InputIterator1,
|
|
InputIterator2,
|
|
OutputIterator1 keys_output,
|
|
OutputIterator2 values_output)
|
|
{
|
|
system.validate_dispatch();
|
|
return thrust::make_pair(keys_output, values_output);
|
|
}
|
|
|
|
TEST(UniqueByKeyTests, TestUniqueByKeyCopyDispatchExplicit)
|
|
{
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
thrust::device_vector<int> vec(1);
|
|
|
|
my_system sys(0);
|
|
thrust::unique_by_key_copy(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> unique_by_key_copy(
|
|
my_tag, InputIterator1, InputIterator1, InputIterator2, OutputIterator1 keys_output, OutputIterator2 values_output)
|
|
{
|
|
*keys_output = 13;
|
|
return thrust::make_pair(keys_output, values_output);
|
|
}
|
|
|
|
TEST(UniqueByKeyTests, TestUniqueByKeyCopyDispatchImplicit)
|
|
{
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
thrust::device_vector<int> vec(1);
|
|
|
|
thrust::unique_by_key_copy(
|
|
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());
|
|
}
|
|
|
|
template <typename T>
|
|
struct is_equal_div_10_unique
|
|
{
|
|
THRUST_HOST_DEVICE bool operator()(const T x, const T& y) const
|
|
{
|
|
return ((int) x / 10) == ((int) y / 10);
|
|
}
|
|
};
|
|
|
|
template <typename Vector>
|
|
void initialize_keys(Vector& keys)
|
|
{
|
|
keys.resize(9);
|
|
keys = {11, 11, 21, 20, 21, 21, 21, 37, 37};
|
|
}
|
|
|
|
template <typename Vector>
|
|
void initialize_values(Vector& values)
|
|
{
|
|
values.resize(9);
|
|
values = {0, 1, 2, 3, 4, 5, 6, 7, 8};
|
|
}
|
|
|
|
TYPED_TEST(UniqueByKeyTests, TestUniqueByKeySimple)
|
|
{
|
|
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 keys;
|
|
Vector values;
|
|
|
|
typename thrust::pair<typename Vector::iterator, typename Vector::iterator> new_last;
|
|
|
|
// basic test
|
|
initialize_keys(keys);
|
|
initialize_values(values);
|
|
|
|
new_last = thrust::unique_by_key(keys.begin(), keys.end(), values.begin());
|
|
|
|
ASSERT_EQ(new_last.first - keys.begin(), 5);
|
|
ASSERT_EQ(new_last.second - values.begin(), 5);
|
|
keys.resize(5);
|
|
values.resize(5);
|
|
Vector keys_ref{11, 21, 20, 21, 37};
|
|
ASSERT_EQ(keys, keys_ref);
|
|
|
|
Vector values_ref{0, 2, 3, 4, 7};
|
|
ASSERT_EQ(values, values_ref);
|
|
|
|
// test BinaryPredicate
|
|
initialize_keys(keys);
|
|
initialize_values(values);
|
|
|
|
new_last = thrust::unique_by_key(keys.begin(), keys.end(), values.begin(), is_equal_div_10_unique<T>());
|
|
|
|
ASSERT_EQ(new_last.first - keys.begin(), 3);
|
|
ASSERT_EQ(new_last.second - values.begin(), 3);
|
|
keys_ref.resize(3);
|
|
keys.resize(3);
|
|
keys_ref = {11, 21, 37};
|
|
ASSERT_EQ(keys, keys_ref);
|
|
|
|
values.resize(3);
|
|
values_ref.resize(3);
|
|
values_ref = {0, 2, 7};
|
|
ASSERT_EQ(values, values_ref);
|
|
}
|
|
|
|
TYPED_TEST(UniqueByKeyTests, TestUniqueCopyByKeySimple)
|
|
{
|
|
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 keys;
|
|
Vector values;
|
|
|
|
typename thrust::pair<typename Vector::iterator, typename Vector::iterator> new_last;
|
|
|
|
// basic test
|
|
initialize_keys(keys);
|
|
initialize_values(values);
|
|
|
|
Vector output_keys(keys.size());
|
|
Vector output_values(values.size());
|
|
|
|
new_last =
|
|
thrust::unique_by_key_copy(keys.begin(), keys.end(), values.begin(), output_keys.begin(), output_values.begin());
|
|
|
|
ASSERT_EQ(new_last.first - output_keys.begin(), 5);
|
|
ASSERT_EQ(new_last.second - output_values.begin(), 5);
|
|
output_keys.resize(5);
|
|
output_values.resize(5);
|
|
Vector keys_ref{11, 21, 20, 21, 37};
|
|
ASSERT_EQ(output_keys, keys_ref);
|
|
|
|
Vector values_ref{0, 2, 3, 4, 7};
|
|
ASSERT_EQ(output_values, values_ref);
|
|
|
|
// test BinaryPredicate
|
|
initialize_keys(keys);
|
|
initialize_values(values);
|
|
|
|
new_last = thrust::unique_by_key_copy(
|
|
keys.begin(), keys.end(), values.begin(), output_keys.begin(), output_values.begin(), is_equal_div_10_unique<T>());
|
|
|
|
ASSERT_EQ(new_last.first - output_keys.begin(), 3);
|
|
ASSERT_EQ(new_last.second - output_values.begin(), 3);
|
|
output_keys.resize(3);
|
|
output_values.resize(3);
|
|
keys_ref = {11, 21, 37};
|
|
ASSERT_EQ(output_keys, keys_ref);
|
|
|
|
values_ref.resize(3);
|
|
values_ref = {0, 2, 7};
|
|
ASSERT_EQ(output_values, values_ref);
|
|
}
|
|
|
|
TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueByKey)
|
|
{
|
|
using K = typename TestFixture::input_type;
|
|
using V = unsigned int; // ValueType
|
|
|
|
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<K> h_keys =
|
|
get_random_data<K>(size, get_default_limits<K>::min(), get_default_limits<K>::max(), seed);
|
|
thrust::host_vector<V> h_vals = get_random_data<V>(
|
|
size, get_default_limits<V>::min(), get_default_limits<V>::max(), seed + seed_value_addition);
|
|
thrust::device_vector<K> d_keys = h_keys;
|
|
thrust::device_vector<V> d_vals = h_vals;
|
|
|
|
using HostKeyIterator = typename thrust::host_vector<K>::iterator;
|
|
using HostValIterator = typename thrust::host_vector<V>::iterator;
|
|
using DeviceKeyIterator = typename thrust::device_vector<K>::iterator;
|
|
using DeviceValIterator = typename thrust::device_vector<V>::iterator;
|
|
|
|
using HostIteratorPair = typename thrust::pair<HostKeyIterator, HostValIterator>;
|
|
using DeviceIteratorPair = typename thrust::pair<DeviceKeyIterator, DeviceValIterator>;
|
|
|
|
HostIteratorPair h_last = thrust::unique_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
|
|
DeviceIteratorPair d_last = thrust::unique_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());
|
|
|
|
ASSERT_EQ(h_last.first - h_keys.begin(), d_last.first - d_keys.begin());
|
|
ASSERT_EQ(h_last.second - h_vals.begin(), d_last.second - d_vals.begin());
|
|
|
|
size_t N = h_last.first - h_keys.begin();
|
|
|
|
h_keys.resize(N);
|
|
h_vals.resize(N);
|
|
d_keys.resize(N);
|
|
d_vals.resize(N);
|
|
|
|
ASSERT_EQ(h_keys, d_keys);
|
|
ASSERT_EQ(h_vals, d_vals);
|
|
}
|
|
}
|
|
}
|
|
|
|
TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueCopyByKey)
|
|
{
|
|
using K = typename TestFixture::input_type;
|
|
using V = unsigned int; // ValueType
|
|
|
|
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<K> h_keys =
|
|
get_random_data<K>(size, get_default_limits<K>::min(), get_default_limits<K>::max(), seed);
|
|
thrust::host_vector<V> h_vals = get_random_data<V>(
|
|
size, get_default_limits<V>::min(), get_default_limits<V>::max(), seed + seed_value_addition);
|
|
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);
|
|
|
|
using HostKeyIterator = typename thrust::host_vector<K>::iterator;
|
|
using HostValIterator = typename thrust::host_vector<V>::iterator;
|
|
using DeviceKeyIterator = typename thrust::device_vector<K>::iterator;
|
|
using DeviceValIterator = typename thrust::device_vector<V>::iterator;
|
|
|
|
using HostIteratorPair = typename thrust::pair<HostKeyIterator, HostValIterator>;
|
|
using DeviceIteratorPair = typename thrust::pair<DeviceKeyIterator, DeviceValIterator>;
|
|
|
|
HostIteratorPair h_last = thrust::unique_by_key_copy(
|
|
h_keys.begin(), h_keys.end(), h_vals.begin(), h_keys_output.begin(), h_vals_output.begin());
|
|
DeviceIteratorPair d_last = thrust::unique_by_key_copy(
|
|
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(UniqueByKeyIntegralTests, TestUniqueCopyByKeyToDiscardIterator)
|
|
{
|
|
using K = typename TestFixture::input_type;
|
|
using V = unsigned int; // ValueType
|
|
|
|
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<K> h_keys =
|
|
get_random_data<K>(size, get_default_limits<K>::min(), get_default_limits<K>::max(), seed);
|
|
thrust::host_vector<V> h_vals = get_random_data<V>(
|
|
size, get_default_limits<V>::min(), get_default_limits<V>::max(), seed + seed_value_addition);
|
|
thrust::device_vector<K> d_keys = h_keys;
|
|
thrust::device_vector<V> d_vals = h_vals;
|
|
|
|
thrust::host_vector<V> h_vals_output(size);
|
|
thrust::device_vector<V> d_vals_output(size);
|
|
|
|
thrust::host_vector<K> h_keys_output(size);
|
|
thrust::device_vector<K> d_keys_output(size);
|
|
|
|
thrust::host_vector<K> h_unique_keys = h_keys;
|
|
h_unique_keys.erase(thrust::unique(h_unique_keys.begin(), h_unique_keys.end()), h_unique_keys.end());
|
|
|
|
size_t num_unique_keys = h_unique_keys.size();
|
|
|
|
// mask both outputs
|
|
thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> h_result1 = thrust::unique_by_key_copy(
|
|
h_keys.begin(), h_keys.end(), h_vals.begin(), thrust::make_discard_iterator(), thrust::make_discard_iterator());
|
|
|
|
thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> d_result1 = thrust::unique_by_key_copy(
|
|
d_keys.begin(), d_keys.end(), d_vals.begin(), thrust::make_discard_iterator(), thrust::make_discard_iterator());
|
|
|
|
thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> reference1 = thrust::make_pair(
|
|
thrust::make_discard_iterator(num_unique_keys), thrust::make_discard_iterator(num_unique_keys));
|
|
|
|
ASSERT_EQ_QUIET(reference1, h_result1);
|
|
ASSERT_EQ_QUIET(reference1, d_result1);
|
|
|
|
// mask values output
|
|
thrust::pair<typename thrust::host_vector<K>::iterator, thrust::discard_iterator<>> h_result2 =
|
|
thrust::unique_by_key_copy(
|
|
h_keys.begin(), h_keys.end(), h_vals.begin(), h_keys_output.begin(), thrust::make_discard_iterator());
|
|
|
|
thrust::pair<typename thrust::device_vector<K>::iterator, thrust::discard_iterator<>> d_result2 =
|
|
thrust::unique_by_key_copy(
|
|
d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys_output.begin(), thrust::make_discard_iterator());
|
|
|
|
thrust::pair<typename thrust::host_vector<K>::iterator, thrust::discard_iterator<>> h_reference2 =
|
|
thrust::make_pair(h_keys_output.begin() + num_unique_keys, thrust::make_discard_iterator(num_unique_keys));
|
|
|
|
thrust::pair<typename thrust::device_vector<K>::iterator, thrust::discard_iterator<>> d_reference2 =
|
|
thrust::make_pair(d_keys_output.begin() + num_unique_keys, thrust::make_discard_iterator(num_unique_keys));
|
|
|
|
ASSERT_EQ(h_keys_output, d_keys_output);
|
|
ASSERT_EQ_QUIET(h_reference2, h_result2);
|
|
ASSERT_EQ_QUIET(d_reference2, d_result2);
|
|
|
|
// mask keys output
|
|
thrust::pair<thrust::discard_iterator<>, typename thrust::host_vector<V>::iterator> h_result3 =
|
|
thrust::unique_by_key_copy(
|
|
h_keys.begin(), h_keys.end(), h_vals.begin(), thrust::make_discard_iterator(), h_vals_output.begin());
|
|
|
|
thrust::pair<thrust::discard_iterator<>, typename thrust::device_vector<V>::iterator> d_result3 =
|
|
thrust::unique_by_key_copy(
|
|
d_keys.begin(), d_keys.end(), d_vals.begin(), thrust::make_discard_iterator(), d_vals_output.begin());
|
|
|
|
thrust::pair<thrust::discard_iterator<>, typename thrust::host_vector<V>::iterator> h_reference3 =
|
|
thrust::make_pair(thrust::make_discard_iterator(num_unique_keys), h_vals_output.begin() + num_unique_keys);
|
|
|
|
thrust::pair<thrust::discard_iterator<>, typename thrust::device_vector<V>::iterator> d_reference3 =
|
|
thrust::make_pair(thrust::make_discard_iterator(num_unique_keys), d_vals_output.begin() + num_unique_keys);
|
|
|
|
ASSERT_EQ(h_vals_output, d_vals_output);
|
|
ASSERT_EQ_QUIET(h_reference3, h_result3);
|
|
ASSERT_EQ_QUIET(d_reference3, d_result3);
|
|
}
|
|
}
|
|
}
|
|
|
|
__global__ THRUST_HIP_LAUNCH_BOUNDS_DEFAULT void
|
|
UniqueByKeyKernel(int const N, int* in_array, int* in_keys, int* out_size)
|
|
{
|
|
if (threadIdx.x == 0)
|
|
{
|
|
thrust::device_ptr<int> in_begin(in_array);
|
|
thrust::device_ptr<int> in_end(in_array + N);
|
|
|
|
thrust::device_ptr<int> keys_begin(in_keys);
|
|
thrust::device_ptr<int> keys_end(in_keys + N);
|
|
|
|
thrust::pair<thrust::device_vector<int>::iterator, thrust::device_vector<int>::iterator> d_result =
|
|
thrust::unique_by_key(thrust::hip::par, keys_begin, keys_end, in_begin);
|
|
out_size[0] = d_result.second - thrust::device_vector<int>::iterator(in_begin);
|
|
out_size[1] = d_result.first - thrust::device_vector<int>::iterator(keys_begin);
|
|
}
|
|
}
|
|
|
|
TEST(UniqueByKeyIntegralTests, TestUniqueDevice)
|
|
{
|
|
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, 15, seed);
|
|
thrust::host_vector<int> h_keys = get_random_data<int>(size, 0, 15, seed);
|
|
thrust::device_vector<int> d_data = h_data;
|
|
thrust::device_vector<int> d_keys = h_keys;
|
|
|
|
thrust::device_vector<int> d_output_size(2, 0);
|
|
|
|
thrust::pair<thrust::host_vector<int>::iterator, thrust::host_vector<int>::iterator> h_new_last;
|
|
// thrust::pair<thrust::device_vector<int>::iterator,thrust::device_vector<int>::iterator> d_new_last;
|
|
|
|
h_new_last = thrust::unique_by_key(h_keys.begin(), h_keys.end(), h_data.begin());
|
|
auto h_values_last = h_new_last.second;
|
|
auto h_values_keys = h_new_last.first;
|
|
|
|
hipLaunchKernelGGL(
|
|
UniqueByKeyKernel,
|
|
dim3(1, 1, 1),
|
|
dim3(128, 1, 1),
|
|
0,
|
|
0,
|
|
size,
|
|
thrust::raw_pointer_cast(&d_data[0]),
|
|
thrust::raw_pointer_cast(&d_keys[0]),
|
|
thrust::raw_pointer_cast(&d_output_size[0]));
|
|
HIP_CHECK(hipGetLastError());
|
|
|
|
ASSERT_EQ(h_values_last - h_data.begin(), d_output_size[0]);
|
|
ASSERT_EQ(h_values_keys - h_keys.begin(), d_output_size[1]);
|
|
|
|
h_data.resize(h_values_last - h_data.begin());
|
|
h_keys.resize(h_values_last - h_data.begin());
|
|
|
|
d_data.resize(d_output_size[0]);
|
|
d_keys.resize(d_output_size[1]);
|
|
|
|
ASSERT_EQ(h_data, d_data);
|
|
ASSERT_EQ(h_keys, d_keys);
|
|
}
|
|
}
|
|
}
|
|
|
|
// OpenMP has issues with these tests, NVIDIA/cccl#1715
|
|
#if THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_OMP
|
|
|
|
# ifndef THRUST_FORCE_32_BIT_OFFSET_TYPE
|
|
|
|
TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueCopyByKeyLargeInput)
|
|
{
|
|
#ifdef ADDRESS_SANITIZER_BUILD
|
|
GTEST_SKIP() << "Skipping large size test for address sanitizer build.";
|
|
#endif
|
|
|
|
using K = typename TestFixture::input_type;
|
|
using type = K;
|
|
using index_type = std::int64_t;
|
|
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
const std::size_t num_items = 4400000000ULL;
|
|
thrust::host_vector<type> reference_keys{static_cast<type>(0), static_cast<type>(1), static_cast<type>(0)};
|
|
thrust::host_vector<index_type> reference_values{0, 4300000000ULL, 4300000001ULL};
|
|
|
|
auto keys_in = thrust::make_transform_iterator(thrust::make_counting_iterator(0ULL), index_to_value_t<type>{});
|
|
auto values_in = thrust::make_counting_iterator(0ULL);
|
|
thrust::device_vector<type> keys_out(reference_keys.size());
|
|
thrust::device_vector<index_type> values_out(reference_values.size());
|
|
|
|
// Run test
|
|
const auto selected_aut_end =
|
|
thrust::unique_by_key_copy(keys_in, keys_in + num_items, values_in, keys_out.begin(), values_out.begin());
|
|
|
|
// Ensure that we created the correct output
|
|
auto const num_selected_out = thrust::distance(keys_out.begin(), selected_aut_end.first);
|
|
ASSERT_EQ(reference_keys.size(), static_cast<std::size_t>(num_selected_out));
|
|
ASSERT_EQ(num_selected_out, thrust::distance(values_out.begin(), selected_aut_end.second));
|
|
keys_out.resize(num_selected_out);
|
|
values_out.resize(num_selected_out);
|
|
ASSERT_EQ(reference_keys, keys_out);
|
|
ASSERT_EQ(reference_values, values_out);
|
|
}
|
|
|
|
TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueCopyByKeyLargeOutCount)
|
|
{
|
|
#ifdef ADDRESS_SANITIZER_BUILD
|
|
GTEST_SKIP() << "Skipping large size test for address sanitizer build.";
|
|
#endif
|
|
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
constexpr std::size_t num_items = 4400000000ULL;
|
|
|
|
auto keys_in = thrust::make_counting_iterator(0ULL);
|
|
auto values_in = thrust::make_counting_iterator(0ULL);
|
|
|
|
// Run test
|
|
auto keys_out = thrust::make_discard_iterator();
|
|
auto values_out = thrust::make_discard_iterator();
|
|
const auto selected_aut_end =
|
|
thrust::unique_by_key_copy(thrust::device, keys_in, keys_in + num_items, values_in, keys_out, values_out);
|
|
|
|
// Ensure that we created the correct output
|
|
auto const num_selected_out = thrust::distance(keys_out, selected_aut_end.first);
|
|
ASSERT_EQ(num_items, static_cast<std::size_t>(num_selected_out));
|
|
ASSERT_EQ(num_selected_out, thrust::distance(values_out, selected_aut_end.second));
|
|
}
|
|
|
|
# endif // THRUST_FORCE_32_BIT_OFFSET_TYPE
|
|
|
|
#endif // non-OpenMP backend
|
|
|
|
// This test fails only on GCC 6
|
|
#if !defined(__GNUC__) || __GNUC__ != 6
|
|
|
|
// Based on GitHub issue: https://github.com/NVIDIA/cccl/issues/1956
|
|
namespace
|
|
{
|
|
struct CompareFirst
|
|
{
|
|
template <typename T>
|
|
THRUST_HOST_DEVICE bool operator()(T const& lhs, T const& rhs) const
|
|
{
|
|
return lhs.first == rhs.first;
|
|
}
|
|
};
|
|
struct Entry
|
|
{
|
|
std::int32_t a;
|
|
float b;
|
|
};
|
|
} // namespace
|
|
|
|
TEST(UniqueWithoutEqualityOperatorTests, TestKeysWithoutEqualityOperator)
|
|
{
|
|
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
|
|
|
|
using Key = thrust::pair<std::int32_t, Entry>;
|
|
|
|
const auto k1 = Key{1, {}};
|
|
const auto k2 = Key{2, {}};
|
|
const thrust::device_vector<Key> keys{k1, k1, k1, k2, k2};
|
|
thrust::device_vector<Entry> data{{0, 0}, {1, 1}, {2, 2}, {3, 3}, {4, 4}};
|
|
|
|
thrust::device_vector<Key> unique_keys(5);
|
|
thrust::device_vector<Entry> unique_data(5);
|
|
|
|
const auto result = thrust::unique_by_key_copy(
|
|
thrust::device, keys.cbegin(), keys.cend(), data.begin(), unique_keys.begin(), unique_data.begin(), CompareFirst{});
|
|
|
|
unique_keys.erase(result.first, unique_keys.end());
|
|
unique_data.erase(result.second, unique_data.end());
|
|
|
|
auto unique_keys_h = thrust::host_vector<Key>(unique_keys);
|
|
auto unique_data_h = thrust::host_vector<Entry>(unique_data);
|
|
|
|
ASSERT_EQ(unique_keys_h[0].first, k1.first);
|
|
ASSERT_EQ(unique_keys_h[0].second.a, k1.second.a);
|
|
ASSERT_EQ(unique_keys_h[0].second.b, k1.second.b);
|
|
ASSERT_EQ(unique_keys_h[1].first, k2.first);
|
|
ASSERT_EQ(unique_keys_h[1].second.a, k2.second.a);
|
|
ASSERT_EQ(unique_keys_h[1].second.b, k2.second.b);
|
|
|
|
ASSERT_EQ(unique_data_h[0].a, 0);
|
|
ASSERT_EQ(unique_data_h[0].b, 0);
|
|
ASSERT_EQ(unique_data_h[1].a, 3);
|
|
ASSERT_EQ(unique_data_h[1].b, 3);
|
|
}
|
|
#endif // !defined(__GNUC__) || __GNUC__ != 6
|