Skip to content

Commit 01bbe09

Browse files
griwesgriwes
committed
The great Thrust index type fix: do not test adj_diff for OMP and TBB.
1 parent 086613d commit 01bbe09

File tree

4 files changed

+115
-59
lines changed

4 files changed

+115
-59
lines changed

‎CMakeLists.txt‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -392,6 +392,9 @@ list(APPEND THRUST_TEST_GLOBS testing/*.cpp)
392392

393393
if ("CUDA" STREQUAL "${THRUST_DEVICE_SYSTEM}")
394394
list(APPEND THRUST_TEST_GLOBS testing/cuda/*.cu)
395+
elseif ("CPP" STREQUAL "${THRUST_DEVICE_SYSTEM}")
396+
list(APPEND THRUST_TEST_GLOBS testing/cpp/*.cu)
397+
list(APPEND THRUST_TEST_GLOBS testing/cpp/*.cpp)
395398
elseif ("OMP" STREQUAL "${THRUST_DEVICE_SYSTEM}")
396399
list(APPEND THRUST_TEST_GLOBS testing/omp/*.cu)
397400
list(APPEND THRUST_TEST_GLOBS testing/omp/*.cpp)

‎testing/adjacent_difference.cu‎

Lines changed: 0 additions & 49 deletions
Original file line numberDiff line numberDiff line change
@@ -161,52 +161,3 @@ void TestAdjacentDifferenceDispatchImplicit()
161161
ASSERT_EQUAL(13, d_input.front());
162162
}
163163
DECLARE_UNITTEST(TestAdjacentDifferenceDispatchImplicit);
164-
165-
struct detect_wrong_difference
166-
{
167-
bool * flag;
168-
169-
__host__ __device__ detect_wrong_difference operator++() const { return *this; }
170-
__host__ __device__ detect_wrong_difference operator*() const { return *this; }
171-
template<typename Difference>
172-
__host__ __device__ detect_wrong_difference operator+(Difference) const { return *this; }
173-
template<typename Index>
174-
__host__ __device__ detect_wrong_difference operator[](Index) const { return *this; }
175-
176-
__device__
177-
void operator=(long long difference) const
178-
{
179-
if (difference != 1)
180-
{
181-
*flag = false;
182-
}
183-
}
184-
};
185-
186-
void TestAdjacentDifferenceWithBigIndexesHelper(int magnitude)
187-
{
188-
thrust::counting_iterator<long long> begin(1);
189-
thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
190-
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
191-
192-
thrust::device_ptr<bool> all_differences_correct = thrust::device_malloc<bool>(1);
193-
*all_differences_correct = true;
194-
195-
detect_wrong_difference out = { thrust::raw_pointer_cast(all_differences_correct) };
196-
197-
thrust::adjacent_difference(thrust::device, begin, end, out);
198-
199-
bool all_differences_correct_h = *all_differences_correct;
200-
thrust::device_free(all_differences_correct);
201-
202-
ASSERT_EQUAL(all_differences_correct_h, true);
203-
}
204-
205-
void TestAdjacentDifferenceWithBigIndexes()
206-
{
207-
TestAdjacentDifferenceWithBigIndexesHelper(30);
208-
TestAdjacentDifferenceWithBigIndexesHelper(31);
209-
TestAdjacentDifferenceWithBigIndexesHelper(32);
210-
TestAdjacentDifferenceWithBigIndexesHelper(33);
211-
}
212-
DECLARE_UNITTEST(TestAdjacentDifferenceWithBigIndexes);

‎testing/cpp/adjacent_difference.cu‎

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
#include <unittest/unittest.h>
2+
#include <thrust/adjacent_difference.h>
3+
#include <thrust/execution_policy.h>
4+
#include <thrust/device_malloc.h>
5+
#include <thrust/device_free.h>
6+
7+
struct detect_wrong_difference
8+
{
9+
bool * flag;
10+
11+
__host__ __device__ detect_wrong_difference operator++() const { return *this; }
12+
__host__ __device__ detect_wrong_difference operator*() const { return *this; }
13+
template<typename Difference>
14+
__host__ __device__ detect_wrong_difference operator+(Difference) const { return *this; }
15+
template<typename Index>
16+
__host__ __device__ detect_wrong_difference operator[](Index) const { return *this; }
17+
18+
__device__
19+
void operator=(long long difference) const
20+
{
21+
if (difference != 1)
22+
{
23+
*flag = false;
24+
}
25+
}
26+
};
27+
28+
void TestAdjacentDifferenceWithBigIndexesHelper(int magnitude)
29+
{
30+
thrust::counting_iterator<long long> begin(1);
31+
thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
32+
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
33+
34+
thrust::device_ptr<bool> all_differences_correct = thrust::device_malloc<bool>(1);
35+
*all_differences_correct = true;
36+
37+
detect_wrong_difference out = { thrust::raw_pointer_cast(all_differences_correct) };
38+
39+
thrust::adjacent_difference(thrust::device, begin, end, out);
40+
41+
bool all_differences_correct_h = *all_differences_correct;
42+
thrust::device_free(all_differences_correct);
43+
44+
ASSERT_EQUAL(all_differences_correct_h, true);
45+
}
46+
47+
void TestAdjacentDifferenceWithBigIndexes()
48+
{
49+
TestAdjacentDifferenceWithBigIndexesHelper(30);
50+
TestAdjacentDifferenceWithBigIndexesHelper(31);
51+
TestAdjacentDifferenceWithBigIndexesHelper(32);
52+
TestAdjacentDifferenceWithBigIndexesHelper(33);
53+
}
54+
DECLARE_UNITTEST(TestAdjacentDifferenceWithBigIndexes);

‎testing/cuda/adjacent_difference.cu‎

Lines changed: 58 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -22,36 +22,36 @@ void TestAdjacentDifferenceDevice(ExecutionPolicy exec, const size_t n)
2222
{
2323
thrust::host_vector<T> h_input = unittest::random_samples<T>(n);
2424
thrust::device_vector<T> d_input = h_input;
25-
25+
2626
thrust::host_vector<T> h_output(n);
2727
thrust::device_vector<T> d_output(n);
28-
28+
2929
thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin());
3030
adjacent_difference_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_output.begin());
3131
{
3232
cudaError_t const err = cudaDeviceSynchronize();
3333
ASSERT_EQUAL(cudaSuccess, err);
3434
}
35-
35+
3636
ASSERT_EQUAL(h_output, d_output);
37-
37+
3838
thrust::adjacent_difference(h_input.begin(), h_input.end(), h_output.begin(), thrust::plus<T>());
3939
adjacent_difference_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_output.begin(), thrust::plus<T>());
4040
{
4141
cudaError_t const err = cudaDeviceSynchronize();
4242
ASSERT_EQUAL(cudaSuccess, err);
4343
}
44-
44+
4545
ASSERT_EQUAL(h_output, d_output);
46-
46+
4747
// in-place operation
4848
thrust::adjacent_difference(h_input.begin(), h_input.end(), h_input.begin(), thrust::plus<T>());
4949
adjacent_difference_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_input.begin(), thrust::plus<T>());
5050
{
5151
cudaError_t const err = cudaDeviceSynchronize();
5252
ASSERT_EQUAL(cudaSuccess, err);
5353
}
54-
54+
5555
ASSERT_EQUAL(h_input, h_output); //computed previously
5656
ASSERT_EQUAL(d_input, d_output); //computed previously
5757
}
@@ -77,15 +77,15 @@ void TestAdjacentDifferenceCudaStreams()
7777
{
7878
cudaStream_t s;
7979
cudaStreamCreate(&s);
80-
80+
8181
thrust::device_vector<int> input(3);
8282
thrust::device_vector<int> output(3);
8383
input[0] = 1; input[1] = 4; input[2] = 6;
84-
84+
8585
thrust::adjacent_difference(thrust::cuda::par.on(s), input.begin(), input.end(), output.begin());
8686

8787
cudaStreamSynchronize(s);
88-
88+
8989
ASSERT_EQUAL(output[0], 1);
9090
ASSERT_EQUAL(output[1], 3);
9191
ASSERT_EQUAL(output[2], 2);
@@ -94,3 +94,51 @@ void TestAdjacentDifferenceCudaStreams()
9494
}
9595
DECLARE_UNITTEST(TestAdjacentDifferenceCudaStreams);
9696

97+
struct detect_wrong_difference
98+
{
99+
bool * flag;
100+
101+
__host__ __device__ detect_wrong_difference operator++() const { return *this; }
102+
__host__ __device__ detect_wrong_difference operator*() const { return *this; }
103+
template<typename Difference>
104+
__host__ __device__ detect_wrong_difference operator+(Difference) const { return *this; }
105+
template<typename Index>
106+
__host__ __device__ detect_wrong_difference operator[](Index) const { return *this; }
107+
108+
__device__
109+
void operator=(long long difference) const
110+
{
111+
if (difference != 1)
112+
{
113+
*flag = false;
114+
}
115+
}
116+
};
117+
118+
void TestAdjacentDifferenceWithBigIndexesHelper(int magnitude)
119+
{
120+
thrust::counting_iterator<long long> begin(1);
121+
thrust::counting_iterator<long long> end = begin + (1ll << magnitude);
122+
ASSERT_EQUAL(thrust::distance(begin, end), 1ll << magnitude);
123+
124+
thrust::device_ptr<bool> all_differences_correct = thrust::device_malloc<bool>(1);
125+
*all_differences_correct = true;
126+
127+
detect_wrong_difference out = { thrust::raw_pointer_cast(all_differences_correct) };
128+
129+
thrust::adjacent_difference(thrust::device, begin, end, out);
130+
131+
bool all_differences_correct_h = *all_differences_correct;
132+
thrust::device_free(all_differences_correct);
133+
134+
ASSERT_EQUAL(all_differences_correct_h, true);
135+
}
136+
137+
void TestAdjacentDifferenceWithBigIndexes()
138+
{
139+
TestAdjacentDifferenceWithBigIndexesHelper(30);
140+
TestAdjacentDifferenceWithBigIndexesHelper(31);
141+
TestAdjacentDifferenceWithBigIndexesHelper(32);
142+
TestAdjacentDifferenceWithBigIndexesHelper(33);
143+
}
144+
DECLARE_UNITTEST(TestAdjacentDifferenceWithBigIndexes);

0 commit comments

Comments
 (0)