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
123
124
125
126
|
#include <thrust/extrema.h>
#include <unittest/unittest.h>
#ifdef THRUST_TEST_DEVICE_SIDE
template <typename ExecutionPolicy, typename Iterator1, typename Iterator2>
__global__ void minmax_element_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result)
{
*result = thrust::minmax_element(exec, first, last);
}
template <typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename BinaryPredicate>
__global__ void
minmax_element_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, BinaryPredicate pred, Iterator2 result)
{
*result = thrust::minmax_element(exec, first, last, pred);
}
template <typename ExecutionPolicy>
void TestMinMaxElementDevice(ExecutionPolicy exec)
{
size_t n = 1000;
thrust::host_vector<int> h_data = unittest::random_samples<int>(n);
thrust::device_vector<int> d_data = h_data;
typename thrust::host_vector<int>::iterator h_min;
typename thrust::host_vector<int>::iterator h_max;
typename thrust::device_vector<int>::iterator d_min;
typename thrust::device_vector<int>::iterator d_max;
typedef thrust::pair<typename thrust::device_vector<int>::iterator, typename thrust::device_vector<int>::iterator>
pair_type;
thrust::device_vector<pair_type> d_result(1);
h_min = thrust::minmax_element(h_data.begin(), h_data.end()).first;
h_max = thrust::minmax_element(h_data.begin(), h_data.end()).second;
d_min = thrust::minmax_element(d_data.begin(), d_data.end()).first;
d_max = thrust::minmax_element(d_data.begin(), d_data.end()).second;
minmax_element_kernel<<<1, 1>>>(exec, d_data.begin(), d_data.end(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}
d_min = ((pair_type) d_result[0]).first;
d_max = ((pair_type) d_result[0]).second;
ASSERT_EQUAL(h_min - h_data.begin(), d_min - d_data.begin());
ASSERT_EQUAL(h_max - h_data.begin(), d_max - d_data.begin());
h_max = thrust::minmax_element(h_data.begin(), h_data.end(), thrust::greater<int>()).first;
h_min = thrust::minmax_element(h_data.begin(), h_data.end(), thrust::greater<int>()).second;
minmax_element_kernel<<<1, 1>>>(exec, d_data.begin(), d_data.end(), thrust::greater<int>(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}
d_max = ((pair_type) d_result[0]).first;
d_min = ((pair_type) d_result[0]).second;
ASSERT_EQUAL(h_min - h_data.begin(), d_min - d_data.begin());
ASSERT_EQUAL(h_max - h_data.begin(), d_max - d_data.begin());
}
void TestMinMaxElementDeviceSeq()
{
TestMinMaxElementDevice(thrust::seq);
}
DECLARE_UNITTEST(TestMinMaxElementDeviceSeq);
void TestMinMaxElementDeviceDevice()
{
TestMinMaxElementDevice(thrust::device);
}
DECLARE_UNITTEST(TestMinMaxElementDeviceDevice);
#endif
void TestMinMaxElementCudaStreams()
{
typedef thrust::device_vector<int> Vector;
Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;
cudaStream_t s;
cudaStreamCreate(&s);
ASSERT_EQUAL(*thrust::minmax_element(thrust::cuda::par.on(s), data.begin(), data.end()).first, 1);
ASSERT_EQUAL(*thrust::minmax_element(thrust::cuda::par.on(s), data.begin(), data.end()).second, 5);
ASSERT_EQUAL(thrust::minmax_element(thrust::cuda::par.on(s), data.begin(), data.end()).first - data.begin(), 2);
ASSERT_EQUAL(thrust::minmax_element(thrust::cuda::par.on(s), data.begin(), data.end()).second - data.begin(), 1);
cudaStreamDestroy(s);
}
DECLARE_UNITTEST(TestMinMaxElementCudaStreams);
void TestMinMaxElementDevicePointer()
{
typedef thrust::device_vector<int> Vector;
typedef Vector::value_type T;
Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;
T* raw_ptr = thrust::raw_pointer_cast(data.data());
size_t n = data.size();
ASSERT_EQUAL(thrust::minmax_element(thrust::device, raw_ptr, raw_ptr + n).first - raw_ptr, 2);
ASSERT_EQUAL(thrust::minmax_element(thrust::device, raw_ptr, raw_ptr + n).second - raw_ptr, 1);
}
DECLARE_UNITTEST(TestMinMaxElementDevicePointer);
|
