diff --git a/performance/find.test b/performance/find.test new file mode 100644 index 000000000..16bac8da1 --- /dev/null +++ b/performance/find.test @@ -0,0 +1,62 @@ +PREAMBLE = \ + """ + #include + #include + #include + + template + void find_partial(const Vector& v) + { + thrust::find(v.begin(), v.end(), 1); + } + + template + void find_full(const Vector& v) + { + thrust::max_element(v.begin(), v.end()); + } + + template + void reduce_full(const Vector& v) + { + thrust::max_element(v.begin(), v.end()); + } + """ + +INITIALIZE = \ + """ + thrust::host_vector<$InputType> h_input($InputSize, 0); + thrust::device_vector<$InputType> d_input($InputSize, 0); + + size_t pos = $Fraction * $InputSize; + + if (pos < $InputSize) + { + h_input[pos] = 1; + d_input[pos] = 1; + } + + size_t h_index = thrust::find(h_input.begin(), h_input.end(), 1) - h_input.begin(); + size_t d_index = thrust::find(d_input.begin(), d_input.end(), 1) - d_input.begin(); + + ASSERT_EQUAL(h_index, d_index); + """ + +TIME = \ + """ + $Method(d_input); + """ + +FINALIZE = \ + """ + RECORD_TIME(); + RECORD_BANDWIDTH(sizeof($InputType) * double($InputSize)); + """ + +InputTypes = ['int'] +InputSizes = [2**23] +Fractions = [0.01, 0.99] +Methods = ['find_partial', 'find_full', 'reduce_full'] + +TestVariables = [('InputType', InputTypes), ('InputSize', InputSizes), ('Fraction', Fractions), ('Method', Methods)] + diff --git a/testing/find.cu b/testing/find.cu index e4c5bbec7..ed03c1da4 100644 --- a/testing/find.cu +++ b/testing/find.cu @@ -21,16 +21,15 @@ void TestFindSimple(void) vec[0] = 1; vec[1] = 2; vec[2] = 3; - vec[3] = 4; + vec[3] = 3; vec[4] = 5; ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 0) - vec.begin(), 5); ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 1) - vec.begin(), 0); ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 2) - vec.begin(), 1); ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 3) - vec.begin(), 2); - ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 4) - vec.begin(), 3); + ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 4) - vec.begin(), 5); ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 5) - vec.begin(), 4); - ASSERT_EQUAL(thrust::find(vec.begin(), vec.end(), 6) - vec.begin(), 5); } DECLARE_VECTOR_UNITTEST(TestFindSimple); @@ -43,16 +42,15 @@ void TestFindIfSimple(void) vec[0] = 1; vec[1] = 2; vec[2] = 3; - vec[3] = 4; + vec[3] = 3; vec[4] = 5; ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(0)) - vec.begin(), 5); ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(1)) - vec.begin(), 0); ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(2)) - vec.begin(), 1); ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(3)) - vec.begin(), 2); - ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(4)) - vec.begin(), 3); + ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(4)) - vec.begin(), 5); ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(5)) - vec.begin(), 4); - ASSERT_EQUAL(thrust::find_if(vec.begin(), vec.end(), equal_to_value_pred(6)) - vec.begin(), 5); } DECLARE_VECTOR_UNITTEST(TestFindIfSimple); diff --git a/testing/mismatch.cu b/testing/mismatch.cu new file mode 100644 index 000000000..c6f6e2d37 --- /dev/null +++ b/testing/mismatch.cu @@ -0,0 +1,29 @@ +#include +#include + +template +void TestMismatchSimple(void) +{ + typedef typename Vector::value_type T; + + Vector a(4); Vector b(4); + a[0] = 1; b[0] = 1; + a[1] = 2; b[1] = 2; + a[2] = 3; b[2] = 4; + a[3] = 4; b[3] = 3; + + ASSERT_EQUAL(thrust::mismatch(a.begin(), a.end(), b.begin()).first - a.begin(), 2); + ASSERT_EQUAL(thrust::mismatch(a.begin(), a.end(), b.begin()).second - b.begin(), 2); + + b[2] = 3; + + ASSERT_EQUAL(thrust::mismatch(a.begin(), a.end(), b.begin()).first - a.begin(), 3); + ASSERT_EQUAL(thrust::mismatch(a.begin(), a.end(), b.begin()).second - b.begin(), 3); + + b[3] = 4; + + ASSERT_EQUAL(thrust::mismatch(a.begin(), a.end(), b.begin()).first - a.begin(), 4); + ASSERT_EQUAL(thrust::mismatch(a.begin(), a.end(), b.begin()).second - b.begin(), 4); +} +DECLARE_VECTOR_UNITTEST(TestMismatchSimple); + diff --git a/thrust/detail/device/generic/find.inl b/thrust/detail/device/generic/find.inl index 6ed144865..3a6635699 100644 --- a/thrust/detail/device/generic/find.inl +++ b/thrust/detail/device/generic/find.inl @@ -16,13 +16,11 @@ #include -#include #include #include #include #include -// Implementation of find_if() using short-circuiting // Contributed by Erich Elsen namespace thrust @@ -34,6 +32,23 @@ namespace device namespace generic { +template +struct find_if_functor +{ + __host__ __device__ + TupleType operator()(const TupleType& lhs, const TupleType& rhs) const + { + // select the smallest index among true results + if (thrust::get<0>(lhs) && thrust::get<0>(rhs)) + return TupleType(true, min(thrust::get<1>(lhs), thrust::get<1>(rhs))); + else if (thrust::get<0>(lhs)) + return lhs; + else + return rhs; + } +}; + + template InputIterator find_if(InputIterator first, InputIterator last, @@ -48,8 +63,11 @@ InputIterator find_if(InputIterator first, const difference_type n = thrust::distance(first, last); + // this implementation breaks up the sequence into separate intervals + // in an attempt to early-out as soon as a value is found + // TODO incorporate sizeof(InputType) into interval_threshold and round to multiple of 32 - const difference_type interval_threshold = n; //1 << 20; // XXX disabled until performance is sorted out + const difference_type interval_threshold = 1 << 20; const difference_type interval_size = std::min(interval_threshold, n); for(difference_type begin = 0; begin < n; begin += interval_size) @@ -75,11 +93,11 @@ InputIterator find_if(InputIterator first, ) ) + end, result_type(false, end), - thrust::maximum() + find_if_functor() ); // see if we found something - if (thrust::get<1>(result) != end) + if (thrust::get<0>(result)) { return first + thrust::get<1>(result); } diff --git a/thrust/detail/device/generic/mismatch.h b/thrust/detail/device/generic/mismatch.h new file mode 100644 index 000000000..a3da430f1 --- /dev/null +++ b/thrust/detail/device/generic/mismatch.h @@ -0,0 +1,47 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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. + */ + + +/*! \file find.h + * \brief Search for differences between sequences [generic device]. + */ + +#pragma once + +#include + +namespace thrust +{ +namespace detail +{ +namespace device +{ +namespace generic +{ + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred); + +} // end namespace generic +} // end namespace device +} // end namespace detail +} // end namespace thrust + +#include + diff --git a/thrust/detail/device/generic/mismatch.inl b/thrust/detail/device/generic/mismatch.inl new file mode 100644 index 000000000..6cfdbc27b --- /dev/null +++ b/thrust/detail/device/generic/mismatch.inl @@ -0,0 +1,78 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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 +#include +#include +#include +#include + +#include + +#include + +// Contributed by Erich Elsen + +namespace thrust +{ +namespace detail +{ +namespace device +{ +namespace generic +{ + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred) +{ + typedef typename thrust::iterator_traits::difference_type difference_type; + typedef typename thrust::tuple result_type; + + const difference_type n = thrust::distance(first1, last1); + + difference_type offset = thrust::detail::device::find_if + ( + thrust::make_transform_iterator + ( + thrust::make_zip_iterator(thrust::make_tuple(first1, first2)), + thrust::detail::tuple_equal_to(pred) + ), + thrust::make_transform_iterator + ( + thrust::make_zip_iterator(thrust::make_tuple(first1, first2)), + thrust::detail::tuple_equal_to(pred) + ) + n, + thrust::detail::equal_to_value(false) + ) + - + thrust::make_transform_iterator + ( + thrust::make_zip_iterator(thrust::make_tuple(first1, first2)), + thrust::detail::tuple_equal_to(pred) + ); + + return thrust::make_pair(first1 + offset, first2 + offset); +} + +} // end namespace generic +} // end namespace device +} // end namespace detail +} // end namespace thrust + diff --git a/thrust/detail/device/mismatch.h b/thrust/detail/device/mismatch.h new file mode 100644 index 000000000..f229638ab --- /dev/null +++ b/thrust/detail/device/mismatch.h @@ -0,0 +1,47 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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. + */ + + +/*! \file mismatch.h + * \brief Search for differences between sequences [device]. + */ + +#pragma once + +#include + +#include + +namespace thrust +{ +namespace detail +{ +namespace device +{ + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred) +{ + return thrust::detail::device::generic::mismatch(first1, last1, first2, pred); +} + +} // end namespace device +} // end namespace detail +} // end namespace thrust + diff --git a/thrust/detail/dispatch/find.h b/thrust/detail/dispatch/find.h index 522f6733e..e0cea3db1 100644 --- a/thrust/detail/dispatch/find.h +++ b/thrust/detail/dispatch/find.h @@ -15,7 +15,7 @@ */ -/*! \file gather.h +/*! \file find.h * \brief Dispatch layer of the find functions. */ diff --git a/thrust/detail/dispatch/mismatch.h b/thrust/detail/dispatch/mismatch.h new file mode 100644 index 000000000..ba1e8d812 --- /dev/null +++ b/thrust/detail/dispatch/mismatch.h @@ -0,0 +1,83 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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. + */ + + +/*! \file mismatch.h + * \brief Dispatch layer of the mismatch function. + */ + +#pragma once + +#include + +#include +#include + +namespace thrust +{ +namespace detail +{ +namespace dispatch +{ + +//////////////// +// Host Paths // +//////////////// +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred, + thrust::host_space_tag, + thrust::host_space_tag) +{ + return thrust::detail::host::mismatch(first1, last1, first2, pred); +} + + +////////////////// +// Device Paths // +////////////////// +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred, + thrust::device_space_tag, + thrust::device_space_tag) +{ + return thrust::detail::device::mismatch(first1, last1, first2, pred); +} + +/////////////// +// Any Paths // +/////////////// +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred, + thrust::any_space_tag, + thrust::any_space_tag) +{ + // default to device + return thrust::detail::device::mismatch(first1, last1, first2, pred); +} + +} // end namespace dispatch +} // end namespace detail +} // end namespace thrust + diff --git a/thrust/detail/host/mismatch.h b/thrust/detail/host/mismatch.h new file mode 100644 index 000000000..9a7fbc0f1 --- /dev/null +++ b/thrust/detail/host/mismatch.h @@ -0,0 +1,47 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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. + */ + + +/*! \file mismatch.h + * \brief Search for differences between sequences [host]. + */ + +#pragma once + +#include + +#include + +namespace thrust +{ +namespace detail +{ +namespace host +{ + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred) +{ + return std::mismatch(first1, last1, first2, pred); +} + +} // end namespace host +} // end namespace detail +} // end namespace thrust + diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h index c4b3f539c..53d4ba05d 100644 --- a/thrust/detail/internal_functional.h +++ b/thrust/detail/internal_functional.h @@ -24,6 +24,18 @@ namespace thrust namespace detail { +// note that detail::equal_to does not force conversion from T2 -> T1 as equal_to does +template +struct equal_to +{ + template + __host__ __device__ + bool operator()(const T1& lhs, const T2& rhs) const + { + return lhs == rhs; + } +}; + // note that equal_to_value does not force conversion from T2 -> T1 as equal_to does template struct equal_to_value @@ -40,6 +52,24 @@ struct equal_to_value } }; +template +struct tuple_equal_to +{ + typedef bool result_type; + + __host__ __device__ + tuple_equal_to(const Predicate& p) : pred(p) {} + + template + __host__ __device__ + bool operator()(const Tuple& t) const + { + return pred(thrust::get<0>(t), thrust::get<1>(t)); + } + + Predicate pred; +}; + } // end namespace detail } // end namespace thrust diff --git a/thrust/detail/mismatch.inl b/thrust/detail/mismatch.inl new file mode 100644 index 000000000..ccd75c6b2 --- /dev/null +++ b/thrust/detail/mismatch.inl @@ -0,0 +1,52 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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. + */ + + +/*! \file mismatch.inl + * \brief Inline file for mismatch.h + */ + +#include +#include + +#include + +namespace thrust +{ + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2) +{ + typedef typename thrust::iterator_traits::value_type InputType1; + + return thrust::mismatch(first1, last1, first2, thrust::detail::equal_to()); +} + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred) +{ + return thrust::detail::dispatch::mismatch(first1, last1, first2, pred, + typename thrust::iterator_space::type(), + typename thrust::iterator_space::type()); +} + +} // end namespace thrust + diff --git a/thrust/mismatch.h b/thrust/mismatch.h new file mode 100644 index 000000000..e2e29df47 --- /dev/null +++ b/thrust/mismatch.h @@ -0,0 +1,54 @@ +/* + * Copyright 2008-2010 NVIDIA Corporation + * + * 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. + */ + + +/*! \file mismatch.h + * \brief Search for differences between sequences. + */ + +#pragma once + +#include + +namespace thrust +{ + +/*! \addtogroup algorithms + */ + +/*! \addtogroup searching + * \ingroup algorithms + * \{ + */ + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2); + +template +thrust::pair mismatch(InputIterator1 first1, + InputIterator1 last1, + InputIterator2 first2, + BinaryPredicate pred); + +/*! \} // end searching + */ + +} // end namespace thrust + +#include +