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
127
128
129
130
|
#include <thrust/execution_policy.h>
#include <thrust/extrema.h>
#include <unittest/unittest.h>
#ifdef THRUST_TEST_DEVICE_SIDE
template <typename ExecutionPolicy, typename Iterator, typename Iterator2>
__global__ void max_element_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Iterator2 result)
{
*result = thrust::max_element(exec, first, last);
}
template <typename ExecutionPolicy, typename Iterator, typename BinaryPredicate, typename Iterator2>
__global__ void
max_element_kernel(ExecutionPolicy exec, Iterator first, Iterator last, BinaryPredicate pred, Iterator2 result)
{
*result = thrust::max_element(exec, first, last, pred);
}
template <typename ExecutionPolicy>
void TestMaxElementDevice(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;
typedef typename thrust::device_vector<int>::iterator iter_type;
thrust::device_vector<iter_type> d_result(1);
typename thrust::host_vector<int>::iterator h_max = thrust::max_element(h_data.begin(), h_data.end());
max_element_kernel<<<1, 1>>>(exec, d_data.begin(), d_data.end(), d_result.begin());
{
cudaError_t const err = cudaDeviceSynchronize();
ASSERT_EQUAL(cudaSuccess, err);
}
ASSERT_EQUAL(h_max - h_data.begin(), (iter_type) d_result[0] - d_data.begin());
typename thrust::host_vector<int>::iterator h_min =
thrust::max_element(h_data.begin(), h_data.end(), thrust::greater<int>());
max_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);
}
ASSERT_EQUAL(h_min - h_data.begin(), (iter_type) d_result[0] - d_data.begin());
}
void TestMaxElementDeviceSeq()
{
TestMaxElementDevice(thrust::seq);
}
DECLARE_UNITTEST(TestMaxElementDeviceSeq);
void TestMaxElementDeviceDevice()
{
TestMaxElementDevice(thrust::device);
}
DECLARE_UNITTEST(TestMaxElementDeviceDevice);
void TestMaxElementDeviceNoSync()
{
TestMaxElementDevice(thrust::cuda::par_nosync);
}
DECLARE_UNITTEST(TestMaxElementDeviceNoSync);
#endif
template <typename ExecutionPolicy>
void TestMaxElementCudaStreams(ExecutionPolicy policy)
{
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;
cudaStream_t s;
cudaStreamCreate(&s);
auto streampolicy = policy.on(s);
ASSERT_EQUAL(*thrust::max_element(streampolicy, data.begin(), data.end()), 5);
ASSERT_EQUAL(thrust::max_element(streampolicy, data.begin(), data.end()) - data.begin(), 1);
ASSERT_EQUAL(*thrust::max_element(streampolicy, data.begin(), data.end(), thrust::greater<T>()), 1);
ASSERT_EQUAL(thrust::max_element(streampolicy, data.begin(), data.end(), thrust::greater<T>()) - data.begin(), 2);
cudaStreamDestroy(s);
}
void TestMaxElementCudaStreamsSync()
{
TestMaxElementCudaStreams(thrust::cuda::par);
}
DECLARE_UNITTEST(TestMaxElementCudaStreamsSync);
void TestMaxElementCudaStreamsNoSync()
{
TestMaxElementCudaStreams(thrust::cuda::par_nosync);
}
DECLARE_UNITTEST(TestMaxElementCudaStreamsNoSync);
void TestMaxElementDevicePointer()
{
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::max_element(thrust::device, raw_ptr, raw_ptr + n) - raw_ptr, 1);
ASSERT_EQUAL(thrust::max_element(thrust::device, raw_ptr, raw_ptr + n, thrust::greater<T>()) - raw_ptr, 2);
}
DECLARE_UNITTEST(TestMaxElementDevicePointer);
|
