1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
|
/*
* 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/functional.h>
#include <thrust/iterator/retag.h>
#include <thrust/partition.h>
#include "test_header.hpp"
TESTS_DEFINE(PartitionPointVectorTests, VectorSignedIntegerTestsParams);
template <typename T>
struct is_even
{
__host__ __device__ bool operator()(T x) const
{
return ((int)x % 2) == 0;
}
};
TYPED_TEST(PartitionPointVectorTests, TestPartitionPointSimple)
{
using Vector = typename TestFixture::input_type;
using T = typename Vector::value_type;
using Iterator = typename Vector::iterator;
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
Vector v(4);
v[0] = 1;
v[1] = 1;
v[2] = 1;
v[3] = 0;
Iterator first = v.begin();
Iterator last = v.begin() + 4;
Iterator ref = first + 3;
ASSERT_EQ_QUIET(ref, thrust::partition_point(first, last, thrust::identity<T>()));
last = v.begin() + 3;
ref = last;
ASSERT_EQ_QUIET(ref, thrust::partition_point(first, last, thrust::identity<T>()));
}
TYPED_TEST(PartitionPointVectorTests, TestPartitionPoint)
{
using Vector = typename TestFixture::input_type;
using T = typename Vector::value_type;
using Iterator = typename Vector::iterator;
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
const size_t n = (1 << 16) + 13;
for(auto seed : get_seeds())
{
SCOPED_TRACE(testing::Message() << "with seed= " << seed);
Vector v = get_random_data<T>(
n, std::numeric_limits<T>::min(), std::numeric_limits<T>::max(), seed);
Iterator ref = thrust::stable_partition(v.begin(), v.end(), is_even<T>());
ASSERT_EQ(ref - v.begin(),
thrust::partition_point(v.begin(), v.end(), is_even<T>()) - v.begin());
}
}
template <typename ForwardIterator, typename Predicate>
__host__ __device__ ForwardIterator
partition_point(my_system& system, ForwardIterator first, ForwardIterator, Predicate)
{
system.validate_dispatch();
return first;
}
TEST(PartitionPointTests, TestPartitionPointDispatchExplicit)
{
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
thrust::device_vector<int> vec(1);
my_system sys(0);
thrust::partition_point(sys, vec.begin(), vec.begin(), 0);
ASSERT_EQ(true, sys.is_valid());
}
template <typename ForwardIterator, typename Predicate>
__host__ __device__ ForwardIterator
partition_point(my_tag, ForwardIterator first, ForwardIterator, Predicate)
{
*first = 13;
return first;
}
TEST(PartitionPointTests, TestPartitionPointDispatchImplicit)
{
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
thrust::device_vector<int> vec(1);
thrust::partition_point(
thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.begin()), 0);
ASSERT_EQ(13, vec.front());
}
|
