From 05d1469f8616c40b7110168312ee37e0c3563e25 Mon Sep 17 00:00:00 2001 From: Nathan Bell Date: Sun, 14 Nov 2010 18:00:11 -0500 Subject: [PATCH 1/2] fixed TestMerge* failures that were observed on sm_11 (Quadro NVS 160M) --- thrust/detail/device/cuda/block/copy.h | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/thrust/detail/device/cuda/block/copy.h b/thrust/detail/device/cuda/block/copy.h index 434fdd06d..168f76ab9 100644 --- a/thrust/detail/device/cuda/block/copy.h +++ b/thrust/detail/device/cuda/block/copy.h @@ -44,7 +44,7 @@ namespace dispatch template -__device__ + __forceinline__ __device__ RandomAccessIterator2 copy(RandomAccessIterator1 first, RandomAccessIterator1 last, RandomAccessIterator2 result, @@ -65,7 +65,7 @@ __device__ template -__device__ + __forceinline__ __device__ RandomAccessIterator2 copy(RandomAccessIterator1 first, RandomAccessIterator1 last, RandomAccessIterator2 result, @@ -93,13 +93,19 @@ __device__ template -__device__ + __forceinline__ __device__ RandomAccessIterator2 copy(RandomAccessIterator1 first, RandomAccessIterator1 last, RandomAccessIterator2 result) { return detail::dispatch::copy(first, last, result, - typename thrust::detail::dispatch::is_trivial_copy::type()); +#if __CUDA_ARCH__ < 200 + // does not work reliably on pre-Fermi due to "Warning: ... assuming global memory space" issues + false_type() +#else + typename thrust::detail::dispatch::is_trivial_copy::type() +#endif + ); } // end copy() } // end namespace block From 8c818db113e0ee206d12a644f6d9e5e2e3d8ba9b Mon Sep 17 00:00:00 2001 From: Jared Hoberock Date: Sat, 20 Nov 2010 23:19:39 -0800 Subject: [PATCH 2/2] Add unit tests for scalar and vector binary searches with compare. --- testing/binary_search_descending.cu | 129 +++++++++++ testing/binary_search_vector_descending.cu | 237 +++++++++++++++++++++ 2 files changed, 366 insertions(+) create mode 100644 testing/binary_search_descending.cu create mode 100644 testing/binary_search_vector_descending.cu diff --git a/testing/binary_search_descending.cu b/testing/binary_search_descending.cu new file mode 100644 index 000000000..48e44ecbc --- /dev/null +++ b/testing/binary_search_descending.cu @@ -0,0 +1,129 @@ +#include +#include +#include + +#include +#include + +////////////////////// +// Scalar Functions // +////////////////////// + +template +void TestScalarLowerBoundDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::lower_bound(vec.begin(), vec.end(), 0, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::lower_bound(vec.begin(), vec.end(), 1, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), 2, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), 3, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::lower_bound(vec.begin(), vec.end(), 4, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::lower_bound(vec.begin(), vec.end(), 5, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::lower_bound(vec.begin(), vec.end(), 6, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 1, thrust::lower_bound(vec.begin(), vec.end(), 7, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 0, thrust::lower_bound(vec.begin(), vec.end(), 8, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 0, thrust::lower_bound(vec.begin(), vec.end(), 9, thrust::greater())); +} +DECLARE_VECTOR_UNITTEST(TestScalarLowerBoundDescendingSimple); + + +template +void TestScalarUpperBoundDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + ASSERT_EQUAL_QUIET(vec.begin() + 5, thrust::upper_bound(vec.begin(), vec.end(), 0, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::upper_bound(vec.begin(), vec.end(), 1, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::upper_bound(vec.begin(), vec.end(), 2, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), 3, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), 4, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::upper_bound(vec.begin(), vec.end(), 5, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::upper_bound(vec.begin(), vec.end(), 6, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::upper_bound(vec.begin(), vec.end(), 7, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 1, thrust::upper_bound(vec.begin(), vec.end(), 8, thrust::greater())); + ASSERT_EQUAL_QUIET(vec.begin() + 0, thrust::upper_bound(vec.begin(), vec.end(), 9, thrust::greater())); +} +DECLARE_VECTOR_UNITTEST(TestScalarUpperBoundDescendingSimple); + + +template +void TestScalarBinarySearchDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + ASSERT_EQUAL(true, thrust::binary_search(vec.begin(), vec.end(), 0, thrust::greater())); + ASSERT_EQUAL(false, thrust::binary_search(vec.begin(), vec.end(), 1, thrust::greater())); + ASSERT_EQUAL(true, thrust::binary_search(vec.begin(), vec.end(), 2, thrust::greater())); + ASSERT_EQUAL(false, thrust::binary_search(vec.begin(), vec.end(), 3, thrust::greater())); + ASSERT_EQUAL(false, thrust::binary_search(vec.begin(), vec.end(), 4, thrust::greater())); + ASSERT_EQUAL(true, thrust::binary_search(vec.begin(), vec.end(), 5, thrust::greater())); + ASSERT_EQUAL(false, thrust::binary_search(vec.begin(), vec.end(), 6, thrust::greater())); + ASSERT_EQUAL(true, thrust::binary_search(vec.begin(), vec.end(), 7, thrust::greater())); + ASSERT_EQUAL(true, thrust::binary_search(vec.begin(), vec.end(), 8, thrust::greater())); + ASSERT_EQUAL(false, thrust::binary_search(vec.begin(), vec.end(), 9, thrust::greater())); +} +DECLARE_VECTOR_UNITTEST(TestScalarBinarySearchDescendingSimple); + + +template +void TestScalarEqualRangeDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), 0, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), 1, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), 2, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), 3, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), 4, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), 5, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), 6, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 1, thrust::equal_range(vec.begin(), vec.end(), 7, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), 8, thrust::greater()).first); + ASSERT_EQUAL_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), 9, thrust::greater()).first); + + ASSERT_EQUAL_QUIET(vec.begin() + 5, thrust::equal_range(vec.begin(), vec.end(), 0, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), 1, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 4, thrust::equal_range(vec.begin(), vec.end(), 2, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), 3, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), 4, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 3, thrust::equal_range(vec.begin(), vec.end(), 5, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), 6, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 2, thrust::equal_range(vec.begin(), vec.end(), 7, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 1, thrust::equal_range(vec.begin(), vec.end(), 8, thrust::greater()).second); + ASSERT_EQUAL_QUIET(vec.begin() + 0, thrust::equal_range(vec.begin(), vec.end(), 9, thrust::greater()).second); +} +DECLARE_VECTOR_UNITTEST(TestScalarEqualRangeDescendingSimple); + diff --git a/testing/binary_search_vector_descending.cu b/testing/binary_search_vector_descending.cu new file mode 100644 index 000000000..bbb71b5c6 --- /dev/null +++ b/testing/binary_search_vector_descending.cu @@ -0,0 +1,237 @@ +#include +#include +#include + +#include +#include + +////////////////////// +// Vector Functions // +////////////////////// + +// convert xxx_vector to xxx_vector +template +struct vector_like +{ + typedef typename ExampleVector::allocator_type alloc; + typedef typename alloc::template rebind::other new_alloc; + typedef thrust::detail::vector_base type; +}; + +template +void TestVectorLowerBoundDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + Vector input(10); + thrust::sequence(input.begin(), input.end()); + + typedef typename vector_like::type IntVector; + + // test with integral output type + IntVector integral_output(10); + typename IntVector::iterator output_end = thrust::lower_bound(vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), thrust::greater()); + + ASSERT_EQUAL_QUIET(integral_output.end(), output_end); + + ASSERT_EQUAL(4, integral_output[0]); + ASSERT_EQUAL(4, integral_output[1]); + ASSERT_EQUAL(3, integral_output[2]); + ASSERT_EQUAL(3, integral_output[3]); + ASSERT_EQUAL(3, integral_output[4]); + ASSERT_EQUAL(2, integral_output[5]); + ASSERT_EQUAL(2, integral_output[6]); + ASSERT_EQUAL(1, integral_output[7]); + ASSERT_EQUAL(0, integral_output[8]); + ASSERT_EQUAL(0, integral_output[9]); +} +DECLARE_VECTOR_UNITTEST(TestVectorLowerBoundDescendingSimple); + + +template +void TestVectorUpperBoundDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + Vector input(10); + thrust::sequence(input.begin(), input.end()); + + typedef typename vector_like::type IntVector; + + // test with integral output type + IntVector integral_output(10); + typename IntVector::iterator output_end = thrust::upper_bound(vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), thrust::greater()); + + ASSERT_EQUAL_QUIET(output_end, integral_output.end()); + + ASSERT_EQUAL(5, integral_output[0]); + ASSERT_EQUAL(4, integral_output[1]); + ASSERT_EQUAL(4, integral_output[2]); + ASSERT_EQUAL(3, integral_output[3]); + ASSERT_EQUAL(3, integral_output[4]); + ASSERT_EQUAL(3, integral_output[5]); + ASSERT_EQUAL(2, integral_output[6]); + ASSERT_EQUAL(2, integral_output[7]); + ASSERT_EQUAL(1, integral_output[8]); + ASSERT_EQUAL(0, integral_output[9]); +} +DECLARE_VECTOR_UNITTEST(TestVectorUpperBoundDescendingSimple); + + +template +void TestVectorBinarySearchDescendingSimple(void) +{ + typedef typename Vector::value_type T; + + Vector vec(5); + + vec[0] = 8; + vec[1] = 7; + vec[2] = 5; + vec[3] = 2; + vec[4] = 0; + + Vector input(10); + thrust::sequence(input.begin(), input.end()); + + typedef typename vector_like::type BoolVector; + typedef typename vector_like::type IntVector; + + // test with boolean output type + BoolVector bool_output(10); + typename BoolVector::iterator bool_output_end = thrust::binary_search(vec.begin(), vec.end(), input.begin(), input.end(), bool_output.begin(), thrust::greater()); + + ASSERT_EQUAL_QUIET(bool_output_end, bool_output.end()); + + ASSERT_EQUAL(true, bool_output[0]); + ASSERT_EQUAL(false, bool_output[1]); + ASSERT_EQUAL(true, bool_output[2]); + ASSERT_EQUAL(false, bool_output[3]); + ASSERT_EQUAL(false, bool_output[4]); + ASSERT_EQUAL(true, bool_output[5]); + ASSERT_EQUAL(false, bool_output[6]); + ASSERT_EQUAL(true, bool_output[7]); + ASSERT_EQUAL(true, bool_output[8]); + ASSERT_EQUAL(false, bool_output[9]); + + // test with integral output type + IntVector integral_output(10, 2); + typename IntVector::iterator int_output_end = thrust::binary_search(vec.begin(), vec.end(), input.begin(), input.end(), integral_output.begin(), thrust::greater()); + + ASSERT_EQUAL_QUIET(int_output_end, integral_output.end()); + + ASSERT_EQUAL(1, integral_output[0]); + ASSERT_EQUAL(0, integral_output[1]); + ASSERT_EQUAL(1, integral_output[2]); + ASSERT_EQUAL(0, integral_output[3]); + ASSERT_EQUAL(0, integral_output[4]); + ASSERT_EQUAL(1, integral_output[5]); + ASSERT_EQUAL(0, integral_output[6]); + ASSERT_EQUAL(1, integral_output[7]); + ASSERT_EQUAL(1, integral_output[8]); + ASSERT_EQUAL(0, integral_output[9]); +} +DECLARE_VECTOR_UNITTEST(TestVectorBinarySearchDescendingSimple); + + +template +struct TestVectorLowerBoundDescending +{ + void operator()(const size_t n) + { +// XXX an MSVC bug causes problems inside std::stable_sort's implementation: +// std::lower_bound/upper_bound is confused with thrust::lower_bound/upper_bound +#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (THRUST_DEVICE_BACKEND == THRUST_DEVICE_BACKEND_OMP) + KNOWN_FAILURE; +#else + thrust::host_vector h_vec = unittest::random_integers(n); thrust::sort(h_vec.begin(), h_vec.end(), thrust::greater()); + thrust::device_vector d_vec = h_vec; + + thrust::host_vector h_input = unittest::random_integers(2*n); + thrust::device_vector d_input = h_input; + + thrust::host_vector h_output(2*n); + thrust::device_vector d_output(2*n); + + thrust::lower_bound(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), thrust::greater()); + thrust::lower_bound(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), thrust::greater()); + + ASSERT_EQUAL(h_output, d_output); +#endif + } +}; +VariableUnitTest TestVectorLowerBoundDescendingInstance; + + +template +struct TestVectorUpperBoundDescending +{ + void operator()(const size_t n) + { +// XXX an MSVC bug causes problems inside std::stable_sort's implementation: +// std::lower_bound/upper_bound is confused with thrust::lower_bound/upper_bound +#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (THRUST_DEVICE_BACKEND == THRUST_DEVICE_BACKEND_OMP) + KNOWN_FAILURE; +#else + thrust::host_vector h_vec = unittest::random_integers(n); thrust::sort(h_vec.begin(), h_vec.end(), thrust::greater()); + thrust::device_vector d_vec = h_vec; + + thrust::host_vector h_input = unittest::random_integers(2*n); + thrust::device_vector d_input = h_input; + + thrust::host_vector h_output(2*n); + thrust::device_vector d_output(2*n); + + thrust::upper_bound(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), thrust::greater()); + thrust::upper_bound(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), thrust::greater()); + + ASSERT_EQUAL(h_output, d_output); +#endif + } +}; +VariableUnitTest TestVectorUpperBoundDescendingInstance; + +template +struct TestVectorBinarySearchDescending +{ + void operator()(const size_t n) + { +// XXX an MSVC bug causes problems inside std::stable_sort's implementation: +// std::lower_bound/upper_bound is confused with thrust::lower_bound/upper_bound +#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC) && (THRUST_DEVICE_BACKEND == THRUST_DEVICE_BACKEND_OMP) + KNOWN_FAILURE; +#else + thrust::host_vector h_vec = unittest::random_integers(n); thrust::sort(h_vec.begin(), h_vec.end(), thrust::greater()); + thrust::device_vector d_vec = h_vec; + + thrust::host_vector h_input = unittest::random_integers(2*n); + thrust::device_vector d_input = h_input; + + thrust::host_vector h_output(2*n); + thrust::device_vector d_output(2*n); + + thrust::binary_search(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), h_output.begin(), thrust::greater()); + thrust::binary_search(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), d_output.begin(), thrust::greater()); + + ASSERT_EQUAL(h_output, d_output); +#endif + } +}; +VariableUnitTest TestVectorBinarySearchDescendingInstance; +