From e779d2e67340763201e816e9ef3c745ab8c19fa0 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Mon, 4 May 2020 16:57:06 -0400 Subject: [PATCH] Fixed issues in TBB/OpenMP backends on MSVC. - CMake - Define `NOMINMAX` on msvc because aaaargh. - Define `_CRT_SECURE_NO_WARNINGS` for examples (fopen, etc warnings) - Define `iterator_category` for testing iterator. - Add missing include to shuffle unit test - Specialize wrapped_function for void return types - MSVC is not a fan of the pattern "return static_cast(expr);" - Replace buggy SFINAE check with static_assert - SFINAE expression was evaluating wrong type (should have been Tuple, not the first element type of Tuple) - Rather than fix SFINAE expression, switch to a static_assert for the better diagnostic - Replace deprecated `tbb/tbb_thread.h` with ``. - Fix overcounting of initial value in tbb scans. - Apparently reverse_join may be called before operator() - Fix partial sum value type to support edge case from unit test: ``` # testing/scan.cu:260: // float -> float with plus operator (int accumulator) thrust::inclusive_scan(float_input.begin(), float_input.end(), float_output.begin(), thrust::plus()); ASSERT_EQUAL(float_output[0], 1.5); ASSERT_EQUAL(float_output[1], 3.0); ASSERT_EQUAL(float_output[2], 6.0); ASSERT_EQUAL(float_output[3], 10.0); ``` --- CMakeLists.txt | 10 ++ testing/copy.cu | 3 + testing/shuffle.cu | 1 + thrust/detail/function.h | 129 +++++++++++++++------ thrust/detail/internal_functional.h | 9 +- thrust/system/tbb/detail/reduce_by_key.inl | 5 +- thrust/system/tbb/detail/scan.inl | 77 +++++------- 7 files changed, 139 insertions(+), 95 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 5b48717cd7..a4f1cf0985 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -236,6 +236,9 @@ if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") # object files: append_option_if_available("/bigobj" THRUST_CXX_WARNINGS) + # "Oh right, this is Visual Studio." + add_compile_definitions("NOMINMAX") + set(THRUST_TREAT_FILE_AS_CXX "/TP") else () append_option_if_available("-Werror" THRUST_CXX_WARNINGS) @@ -679,6 +682,13 @@ foreach (THRUST_EXAMPLE_SOURCE IN LISTS THRUST_EXAMPLES) endif () endif () + if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}") + # Some examples use unsafe APIs (e.g. fopen) that MSVC will complain about + # unless this is set: + set_target_properties(${THRUST_EXAMPLE} + PROPERTIES COMPILE_DEFINITIONS "_CRT_SECURE_NO_WARNINGS") + endif() + add_test(NAME ${THRUST_EXAMPLE} COMMAND ${CMAKE_COMMAND} -DTHRUST_EXAMPLE=${THRUST_EXAMPLE} diff --git a/testing/copy.cu b/testing/copy.cu index 17c46292c0..8934b877f2 100644 --- a/testing/copy.cu +++ b/testing/copy.cu @@ -724,6 +724,8 @@ struct only_set_when_expected_it __host__ __device__ only_set_when_expected_it operator*() const { return *this; } template __host__ __device__ only_set_when_expected_it operator+(Difference) const { return *this; } + template + __host__ __device__ only_set_when_expected_it operator+=(Difference) const { return *this; } template __host__ __device__ only_set_when_expected_it operator[](Index) const { return *this; } @@ -744,6 +746,7 @@ struct iterator_traits { typedef long long value_type; typedef only_set_when_expected_it reference; + typedef thrust::random_access_device_iterator_tag iterator_category; }; } diff --git a/testing/shuffle.cu b/testing/shuffle.cu index 8c95720712..2d9094b421 100644 --- a/testing/shuffle.cu +++ b/testing/shuffle.cu @@ -2,6 +2,7 @@ #if THRUST_CPP_DIALECT >= 2011 #include +#include #include #include #include diff --git a/thrust/detail/function.h b/thrust/detail/function.h index f1f9e9c943..ee009e8996 100644 --- a/thrust/detail/function.h +++ b/thrust/detail/function.h @@ -24,80 +24,133 @@ namespace thrust namespace detail { - -template - struct wrapped_function +template +struct wrapped_function { // mutable because Function::operator() might be const mutable Function m_f; inline __host__ __device__ wrapped_function() - : m_f() + : m_f() {} inline __host__ __device__ - wrapped_function(const Function &f) - : m_f(f) + wrapped_function(const Function& f) + : m_f(f) {} __thrust_exec_check_disable__ - template + template inline __host__ __device__ - Result operator()(Argument &x) const + Result operator()(Argument& x) const { - // we static cast to Result to handle void Result without error - // in case Function's result is non-void - return static_cast(m_f(thrust::raw_reference_cast(x))); + return m_f(thrust::raw_reference_cast(x)); } __thrust_exec_check_disable__ - template - inline __host__ __device__ Result operator()(const Argument &x) const + template + inline __host__ __device__ + Result operator()(const Argument& x) const { - // we static cast to Result to handle void Result without error - // in case Function's result is non-void - return static_cast(m_f(thrust::raw_reference_cast(x))); + return m_f(thrust::raw_reference_cast(x)); } __thrust_exec_check_disable__ - template - inline __host__ __device__ Result operator()(Argument1 &x, Argument2 &y) const + template + inline __host__ __device__ + Result operator()(Argument1& x, Argument2& y) const { - // we static cast to Result to handle void Result without error - // in case Function's result is non-void - return static_cast(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y))); + return m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); } __thrust_exec_check_disable__ - template - inline __host__ __device__ Result operator()(const Argument1 &x, Argument2 &y) const + template + inline __host__ __device__ + Result operator()(const Argument1& x, Argument2& y) const { - // we static cast to Result to handle void Result without error - // in case Function's result is non-void - return static_cast(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y))); + return m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); } __thrust_exec_check_disable__ - template - inline __host__ __device__ Result operator()(const Argument1 &x, const Argument2 &y) const + template + inline __host__ __device__ + Result operator()(const Argument1& x, const Argument2& y) const { - // we static cast to Result to handle void Result without error - // in case Function's result is non-void - return static_cast(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y))); + return m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); } __thrust_exec_check_disable__ - template - inline __host__ __device__ Result operator()(Argument1 &x, const Argument2 &y) const + template + inline __host__ __device__ + Result operator()(Argument1& x, const Argument2& y) const { - // we static cast to Result to handle void Result without error - // in case Function's result is non-void - return static_cast(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y))); + return m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); } }; // end wrapped_function +// Specialize for void return types: +template +struct wrapped_function +{ + // mutable because Function::operator() might be const + mutable Function m_f; + inline __host__ __device__ + wrapped_function() + : m_f() + {} + + inline __host__ __device__ + wrapped_function(const Function& f) + : m_f(f) + {} + + __thrust_exec_check_disable__ + template + inline __host__ __device__ + void operator()(Argument& x) const + { + m_f(thrust::raw_reference_cast(x)); + } -} // end detail -} // end thrust + __thrust_exec_check_disable__ + template + inline __host__ __device__ + void operator()(const Argument& x) const + { + m_f(thrust::raw_reference_cast(x)); + } + + __thrust_exec_check_disable__ + template + inline __host__ __device__ + void operator()(Argument1& x, Argument2& y) const + { + m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); + } + + __thrust_exec_check_disable__ + template + inline __host__ __device__ + void operator()(const Argument1& x, Argument2& y) const + { + m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); + } + __thrust_exec_check_disable__ + template + inline __host__ __device__ + void operator()(const Argument1& x, const Argument2& y) const + { + m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); + } + __thrust_exec_check_disable__ + template + inline __host__ __device__ + void operator()(Argument1& x, const Argument2& y) const + { + m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)); + } +}; // end wrapped_function +} // namespace detail +} // namespace thrust diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h index 737e75eb4a..c73fab8644 100644 --- a/thrust/detail/internal_functional.h +++ b/thrust/detail/internal_functional.h @@ -23,6 +23,7 @@ #include #include +#include #include #include #include @@ -317,11 +318,11 @@ template __thrust_exec_check_disable__ template inline __host__ __device__ - typename enable_if_non_const_reference_or_tuple_of_iterator_references< - typename thrust::tuple_element<1,Tuple>::type - >::type - operator()(Tuple t) + void operator()(Tuple t) { + THRUST_STATIC_ASSERT_MSG(is_non_const_reference::value || + is_tuple_of_iterator_references::value, + "Expected a non-const reference or thrust::detail::tuple_of_iterator_references"); thrust::get<1>(t) = f(thrust::get<0>(t)); } }; diff --git a/thrust/system/tbb/detail/reduce_by_key.inl b/thrust/system/tbb/detail/reduce_by_key.inl index a9516e4a1c..70933f3070 100644 --- a/thrust/system/tbb/detail/reduce_by_key.inl +++ b/thrust/system/tbb/detail/reduce_by_key.inl @@ -27,8 +27,9 @@ #include #include #include -#include + #include +#include namespace thrust @@ -281,7 +282,7 @@ template(1u, ::tbb::tbb_thread::hardware_concurrency()); + const unsigned int p = thrust::max(1u, std::thread::hardware_concurrency()); // generate O(P) intervals of sequential work // XXX oversubscribing is a tuning opportunity diff --git a/thrust/system/tbb/detail/scan.inl b/thrust/system/tbb/detail/scan.inl index 477c04ee38..e4f7872197 100644 --- a/thrust/system/tbb/detail/scan.inl +++ b/thrust/system/tbb/detail/scan.inl @@ -104,7 +104,12 @@ struct inclusive_body void reverse_join(inclusive_body& b) { - sum = binary_op(b.sum, sum); + // Only accumulate this functor's partial sum if this functor has been + // called at least once -- otherwise we'll over-count the initial value. + if (!first_call) + { + sum = binary_op(b.sum, sum); + } } void assign(inclusive_body& b) @@ -172,8 +177,13 @@ struct exclusive_body void reverse_join(exclusive_body& b) { - sum = binary_op(b.sum, sum); - } + // Only accumulate this functor's partial sum if this functor has been + // called at least once -- otherwise we'll over-count the initial value. + if (!first_call) + { + sum = binary_op(b.sum, sum); + } + } void assign(exclusive_body& b) { @@ -183,8 +193,6 @@ struct exclusive_body } // end scan_detail - - template @@ -194,32 +202,16 @@ template::type - using namespace thrust::detail; typedef typename eval_if< - has_result_type::value, - result_type, - eval_if< - is_output_iterator::value, - thrust::iterator_value, - thrust::iterator_value - > + is_output_iterator::value, + thrust::iterator_value, + thrust::iterator_value >::type ValueType; - - typedef typename thrust::iterator_difference::type Size; - + + typedef typename thrust::iterator_difference::type Size; + Size n = thrust::distance(first, last); if (n != 0) @@ -228,13 +220,12 @@ template(0,n), scan_body); } - + thrust::advance(result, n); return result; } - template::type - using namespace thrust::detail; typedef typename eval_if< - has_result_type::value, - result_type, - eval_if< - is_output_iterator::value, - thrust::iterator_value, - thrust::iterator_value - > + is_output_iterator::value, + thrust::iterator_value, + thrust::iterator_value >::type ValueType; - typedef typename thrust::iterator_difference::type Size; - + typedef typename thrust::iterator_difference::type Size; + Size n = thrust::distance(first, last); if (n != 0) @@ -280,7 +255,7 @@ template(0,n), scan_body); } - + thrust::advance(result, n); return result;