From 85ed4f1ac950ca3715473e34cb4aaecb0241c51f Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Thu, 14 Oct 2021 10:52:15 -0700 Subject: [PATCH 01/18] commit for vis --- thrust/detail/internal_functional.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h index 74ff23741..c4876d4f4 100644 --- a/thrust/detail/internal_functional.h +++ b/thrust/detail/internal_functional.h @@ -321,7 +321,10 @@ template >::type operator()(Tuple t) { - thrust::get<1>(t) = f(thrust::get<0>(t)); + typedef typename thrust::tuple_element<1,decltype(t)>::type this_type; + // mstack adding static_cast + thrust::get<1>(t) = static_cast(f(thrust::get<0>(t))); + //thrust::get<1>(t) = f(thrust::get<0>(t)); } }; From e4c52255daaec1a1fb6ad8142b18071a1cd11cae Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:19:52 -0700 Subject: [PATCH 02/18] conversions work --- examples/bucket_sort2d.cu | 4 ++-- examples/cuda/range_view.cu | 2 +- examples/discrete_voronoi.cu | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/examples/bucket_sort2d.cu b/examples/bucket_sort2d.cu index 9e3bb2720..aaa193c5a 100644 --- a/examples/bucket_sort2d.cu +++ b/examples/bucket_sort2d.cu @@ -37,8 +37,8 @@ struct point_to_bucket_index : public thrust::unary_function unsigned int operator()(const vec2& v) const { // find the raster indices of p's bucket - unsigned int x = static_cast(thrust::get<0>(v) * width); - unsigned int y = static_cast(thrust::get<1>(v) * height); + unsigned int x = static_cast(thrust::get<0>(v)) * width; + unsigned int y = static_cast(thrust::get<1>(v)) * height; // return the bucket's linear index return y * width + x; diff --git a/examples/cuda/range_view.cu b/examples/cuda/range_view.cu index 2ede62047..ce14c7165 100644 --- a/examples/cuda/range_view.cu +++ b/examples/cuda/range_view.cu @@ -181,7 +181,7 @@ __host__ __device__ void saxpy(float A, View1 X, View2 Y, View3 Z) { // Z = A * X + Y - const int size = X.size(); + const int size = static_cast(X.size()); thrust::for_each(thrust::device, thrust::make_counting_iterator(0), thrust::make_counting_iterator(size), diff --git a/examples/discrete_voronoi.cu b/examples/discrete_voronoi.cu index bfbf2242d..2d68e1e12 100644 --- a/examples/discrete_voronoi.cu +++ b/examples/discrete_voronoi.cu @@ -229,7 +229,7 @@ int main(void) } display_time(t); - std::cout <<" ( " << seeds.size() / (1e6 * t.elapsed()) << " MPixel/s ) " << std::endl; + std::cout <<" ( " << static_cast(seeds.size()) / (1e6 * t.elapsed()) << " MPixel/s ) " << std::endl; std::cout << "[Device to Host Copy]" << std::endl; t.restart(); From 56924f70ea86413900280d4b7d7c38de2d7c4642 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:20:16 -0700 Subject: [PATCH 03/18] conversions work --- examples/monte_carlo.cu | 4 ++-- examples/monte_carlo_disjoint_sequences.cu | 4 ++-- examples/sort.cu | 2 +- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/examples/monte_carlo.cu b/examples/monte_carlo.cu index 4a11c4de8..221689f91 100644 --- a/examples/monte_carlo.cu +++ b/examples/monte_carlo.cu @@ -56,7 +56,7 @@ struct estimate_pi : public thrust::unary_function sum *= 4.0f; // divide by N - return sum / N; + return sum / static_cast(N); } }; @@ -70,7 +70,7 @@ int main(void) estimate_pi(), 0.0f, thrust::plus()); - estimate /= M; + estimate /= static_cast(M); std::cout << std::setprecision(3); std::cout << "pi is approximately " << estimate << std::endl; diff --git a/examples/monte_carlo_disjoint_sequences.cu b/examples/monte_carlo_disjoint_sequences.cu index 77b0d0086..eab74887e 100644 --- a/examples/monte_carlo_disjoint_sequences.cu +++ b/examples/monte_carlo_disjoint_sequences.cu @@ -62,7 +62,7 @@ struct estimate_pi : public thrust::unary_function sum *= 4.0f; // divide by N - return sum / N; + return sum / static_cast(N); } }; @@ -76,7 +76,7 @@ int main(void) estimate_pi(), 0.0f, thrust::plus()); - estimate /= M; + estimate /= static_cast(M); std::cout << "pi is around " << estimate << std::endl; diff --git a/examples/sort.cu b/examples/sort.cu index 1bbb5d897..97deb9423 100644 --- a/examples/sort.cu +++ b/examples/sort.cu @@ -19,7 +19,7 @@ void initialize(thrust::device_vector& v) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(2, 19); for(size_t i = 0; i < v.size(); i++) - v[i] = dist(rng) / 2.0f; + v[i] = static_cast(dist(rng)) / 2.0f; } void initialize(thrust::device_vector< thrust::pair >& v) From f2e967d55dacb1d512a24340967c89c4b91a1e69 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:20:38 -0700 Subject: [PATCH 04/18] conversions work --- testing/sort_variable_bits.cu | 3 ++- testing/transform_output_iterator.cu | 3 ++- testing/transform_reduce.cu | 4 +++- 3 files changed, 7 insertions(+), 3 deletions(-) diff --git a/testing/sort_variable_bits.cu b/testing/sort_variable_bits.cu index 4192e3da6..953e53659 100644 --- a/testing/sort_variable_bits.cu +++ b/testing/sort_variable_bits.cu @@ -25,7 +25,8 @@ struct TestSortVariableBits size_t mask = (1 << num_bits) - 1; for(size_t i = 0; i < n; i++) - h_keys[i] &= mask; + // mstack adding static_cast + h_keys[i] &= static_cast(mask); thrust::host_vector reference = h_keys; thrust::device_vector d_keys = h_keys; diff --git a/testing/transform_output_iterator.cu b/testing/transform_output_iterator.cu index 403862256..a4756e071 100644 --- a/testing/transform_output_iterator.cu +++ b/testing/transform_output_iterator.cu @@ -48,7 +48,8 @@ void TestMakeTransformOutputIterator(void) Vector output(4); // initialize input - thrust::sequence(input.begin(), input.end(), 1); + // mstack adding static_cast + thrust::sequence(input.begin(), input.end(), static_cast(1)); thrust::copy(input.begin(), input.end(), thrust::make_transform_output_iterator(output.begin(), UnaryFunction())); diff --git a/testing/transform_reduce.cu b/testing/transform_reduce.cu index 3ff3159d6..9e30bf0f0 100644 --- a/testing/transform_reduce.cu +++ b/testing/transform_reduce.cu @@ -120,7 +120,9 @@ void TestTransformReduceCountingIterator(void) thrust::counting_iterator first(1); - T result = thrust::transform_reduce(first, first + 3, thrust::negate(), 0, thrust::plus()); + // mstack adding static_cast + //T result = thrust::transform_reduce(first, first + 3, thrust::negate(), 0, thrust::plus()); + T result = static_cast(thrust::transform_reduce(first, first + 3, thrust::negate(), 0, thrust::plus())); ASSERT_EQUAL(result, -6); } From 9aacdaef1b07ac31fb37916bc3af479af7fb884b Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:21:08 -0700 Subject: [PATCH 05/18] conversions work --- testing/unittest/assertions.h | 6 ++++++ thrust/detail/allocator/allocator_traits.inl | 4 ++-- thrust/detail/function.h | 11 +++++++++++ 3 files changed, 19 insertions(+), 2 deletions(-) diff --git a/testing/unittest/assertions.h b/testing/unittest/assertions.h index 855d705a4..3511f90dc 100644 --- a/testing/unittest/assertions.h +++ b/testing/unittest/assertions.h @@ -118,6 +118,8 @@ template void assert_equal(T1 a, T2 b, const std::string& filename = "unknown", int lineno = -1) { +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion" if(!(a == b)){ unittest::UnitTestFailure f; f << "[" << filename << ":" << lineno << "] "; @@ -125,6 +127,7 @@ void assert_equal(T1 a, T2 b, f << " [type='" << type_name() << "']"; throw f; } +#pragma GCC diagnostic pop } void assert_equal(char a, char b, @@ -144,6 +147,8 @@ template void assert_equal_quiet(const T1& a, const T2& b, const std::string& filename = "unknown", int lineno = -1) { +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion" if(!(a == b)){ unittest::UnitTestFailure f; f << "[" << filename << ":" << lineno << "] "; @@ -151,6 +156,7 @@ void assert_equal_quiet(const T1& a, const T2& b, f << " [type='" << type_name() << "']"; throw f; } +#pragma GCC diagnostic pop } //// diff --git a/thrust/detail/allocator/allocator_traits.inl b/thrust/detail/allocator/allocator_traits.inl index 1d8d92a9c..6f3fc75c8 100644 --- a/thrust/detail/allocator/allocator_traits.inl +++ b/thrust/detail/allocator/allocator_traits.inl @@ -223,7 +223,7 @@ template >::type construct(Alloc &, T *p, const Arg1 &arg1) { - ::new(static_cast(p)) T(arg1); + ::new(static_cast(p)) T(static_cast(arg1)); } #if THRUST_CPP_DIALECT >= 2011 @@ -254,7 +254,7 @@ template >::type construct(Alloc &, T* p, Args&&... args) { - ::new(static_cast(p)) T(THRUST_FWD(args)...); + ::new(static_cast(p)) T(THRUST_FWD(static_cast(args))...); } #endif diff --git a/thrust/detail/function.h b/thrust/detail/function.h index 66e6d4e4e..29eb7344e 100644 --- a/thrust/detail/function.h +++ b/thrust/detail/function.h @@ -19,6 +19,9 @@ #include #include +// mstack adding type_traits +#include + THRUST_NAMESPACE_BEGIN namespace detail @@ -88,8 +91,16 @@ struct wrapped_function inline __host__ __device__ Result operator()(Argument1& x, const Argument2& y) const { + // mstack adding change +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion" return static_cast(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y))); +#pragma GCC diagnostic pop + //using f_type = std::invoke_result_t; + //using f_type = std::result_of::type; + //return static_cast(m_f(static_cast(thrust::raw_reference_cast(x)), + // static_cast(thrust::raw_reference_cast(y)))); } }; // end wrapped_function From 52ce81534f8b3f73c65d6a1f6a50ea19df42f6f7 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:21:32 -0700 Subject: [PATCH 06/18] conversions work --- thrust/detail/internal_functional.h | 5 +++-- thrust/detail/tuple.inl | 8 ++++++++ thrust/detail/vector_base.inl | 2 +- 3 files changed, 12 insertions(+), 3 deletions(-) diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h index c4876d4f4..aae373ded 100644 --- a/thrust/detail/internal_functional.h +++ b/thrust/detail/internal_functional.h @@ -323,8 +323,9 @@ template { typedef typename thrust::tuple_element<1,decltype(t)>::type this_type; // mstack adding static_cast - thrust::get<1>(t) = static_cast(f(thrust::get<0>(t))); - //thrust::get<1>(t) = f(thrust::get<0>(t)); + //thrust::get<1>(t) = static_cast(f(thrust::get<0>(t))); + //thrust::get<1>(t) = f(static_cast(thrust::get<0>(t))); + thrust::get<1>(t) = f(thrust::get<0>(t)); } }; diff --git a/thrust/detail/tuple.inl b/thrust/detail/tuple.inl index 73367ed44..48135d540 100644 --- a/thrust/detail/tuple.inl +++ b/thrust/detail/tuple.inl @@ -359,7 +359,10 @@ template template inline __host__ __device__ cons& operator=( const cons& u ) { +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion" head=u.head; tail=u.tail; return *this; +#pragma GCC diagnostic pop } // must define assignment operator explicitly, implicit version is @@ -467,7 +470,12 @@ template inline __host__ __device__ cons& operator=(const cons& u ) { +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion" head = u.head; +#pragma GCC diagnostic pop + //head = static_cast(u.head); + //head = static_cast(u.head); return *this; } diff --git a/thrust/detail/vector_base.inl b/thrust/detail/vector_base.inl index 915f37699..05c910d5b 100644 --- a/thrust/detail/vector_base.inl +++ b/thrust/detail/vector_base.inl @@ -205,7 +205,7 @@ template IteratorOrIntegralType value, true_type) { - fill_init(n,value); + fill_init(n,static_cast(value)); } // end vector_base::init_dispatch() template From 01b2eec4e0f85a0b2f292de310f9c02fab23640e Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:21:58 -0700 Subject: [PATCH 07/18] conversions work --- thrust/random/detail/subtract_with_carry_engine.inl | 2 +- thrust/random/detail/xor_combine_engine_max.h | 4 ++-- thrust/random/linear_feedback_shift_engine.h | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/thrust/random/detail/subtract_with_carry_engine.inl b/thrust/random/detail/subtract_with_carry_engine.inl index 0cd60960f..62bf66b41 100644 --- a/thrust/random/detail/subtract_with_carry_engine.inl +++ b/thrust/random/detail/subtract_with_carry_engine.inl @@ -63,7 +63,7 @@ template { // XXX we probably need to cache these m_x[m_k] in a register // maybe we need to cache the use of all member variables - int short_index = m_k - short_lag; + int short_index = m_k - static_cast(short_lag); if(short_index < 0) short_index += long_lag; result_type xi; diff --git a/thrust/random/detail/xor_combine_engine_max.h b/thrust/random/detail/xor_combine_engine_max.h index 0756ff9e0..b16158b70 100644 --- a/thrust/random/detail/xor_combine_engine_max.h +++ b/thrust/random/detail/xor_combine_engine_max.h @@ -287,14 +287,14 @@ template::value - 1 + two_to_the_power(w-s1)>::value - 1 >::value; static const result_type m2 = math::min< result_type, result_type(Engine2::max - Engine2::min), - two_to_the_power::value - 1 + two_to_the_power(w-s2)>::value - 1 >::value; static const result_type s = s1 - s2; diff --git a/thrust/random/linear_feedback_shift_engine.h b/thrust/random/linear_feedback_shift_engine.h index a46c6d8ab..9019f591f 100644 --- a/thrust/random/linear_feedback_shift_engine.h +++ b/thrust/random/linear_feedback_shift_engine.h @@ -93,7 +93,7 @@ template static const result_type wordmask = detail::linear_feedback_shift_engine_wordmask< result_type, - w + static_cast(w) >::value; /*! \endcond */ From a0487acf30f7474068077c7496927502672d1771 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 20 Oct 2021 12:22:16 -0700 Subject: [PATCH 08/18] conversions work --- thrust/system/cuda/detail/async/reduce.h | 6 ++++-- thrust/system/detail/sequential/stable_radix_sort.inl | 4 ++-- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/thrust/system/cuda/detail/async/reduce.h b/thrust/system/cuda/detail/async/reduce.h index 03e3dfd1a..88ba34f6c 100644 --- a/thrust/system/cuda/detail/async/reduce.h +++ b/thrust/system/cuda/detail/async/reduce.h @@ -196,9 +196,10 @@ auto async_reduce( , T init , BinaryOp op ) +// mstack adding static_cast THRUST_RETURNS( thrust::system::cuda::detail::async_reduce_n( - policy, first, distance(first, last), init, op + policy, first, static_cast(distance(first, last)), init, op ) ) @@ -333,9 +334,10 @@ auto async_reduce_into( , T init , BinaryOp op ) +// mstack adding static_cast THRUST_RETURNS( thrust::system::cuda::detail::async_reduce_into_n( - policy, first, distance(first, last), output, init, op + policy, first, static_cast(distance(first, last)), output, init, op ) ) diff --git a/thrust/system/detail/sequential/stable_radix_sort.inl b/thrust/system/detail/sequential/stable_radix_sort.inl index 04bf6cdfe..00eb9269e 100644 --- a/thrust/system/detail/sequential/stable_radix_sort.inl +++ b/thrust/system/detail/sequential/stable_radix_sort.inl @@ -161,7 +161,7 @@ template inline __host__ __device__ size_t operator()(KeyType key) { - const EncodedType x = encode(key); + const EncodedType x = static_cast(encode(key)); // note that we mutate the histogram here return histogram[(x >> bit_shift) & BitMask]++; @@ -259,7 +259,7 @@ void radix_sort(sequential::execution_policy &exec, // compute histograms for(size_t i = 0; i < N; i++) { - const EncodedType x = encode(keys1[i]); + const EncodedType x = static_cast(encode(keys1[i])); for(unsigned int j = 0; j < NumHistograms; j++) { From 71ea6e08b720c8376e83268280c9a475c1933341 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Mon, 31 Jan 2022 08:27:28 -0800 Subject: [PATCH 09/18] batch of changes 1 --- cmake/ThrustBuildCompilerTargets.cmake | 3 ++- testing/for_each.cu | 6 ++++-- testing/pair_transform.cu | 3 ++- testing/scan_by_key.cu | 2 +- 4 files changed, 9 insertions(+), 5 deletions(-) diff --git a/cmake/ThrustBuildCompilerTargets.cmake b/cmake/ThrustBuildCompilerTargets.cmake index bf0b31ed4..282d70a31 100644 --- a/cmake/ThrustBuildCompilerTargets.cmake +++ b/cmake/ThrustBuildCompilerTargets.cmake @@ -77,7 +77,7 @@ function(thrust_build_compiler_targets) # "Oh right, this is Visual Studio." list(APPEND cxx_compile_definitions "NOMINMAX") else() - append_option_if_available("-Werror" cxx_compile_options) + #append_option_if_available("-Werror" cxx_compile_options) append_option_if_available("-Wall" cxx_compile_options) append_option_if_available("-Wextra" cxx_compile_options) append_option_if_available("-Winit-self" cxx_compile_options) @@ -86,6 +86,7 @@ function(thrust_build_compiler_targets) append_option_if_available("-Wpointer-arith" cxx_compile_options) append_option_if_available("-Wunused-local-typedef" cxx_compile_options) append_option_if_available("-Wvla" cxx_compile_options) + append_option_if_available("-Wconversion" cxx_compile_options) # Disable GNU extensions (flag is clang only) append_option_if_available("-Wgnu" cxx_compile_options) diff --git a/testing/for_each.cu b/testing/for_each.cu index 8040e5f78..e9e0eb921 100644 --- a/testing/for_each.cu +++ b/testing/for_each.cu @@ -199,7 +199,7 @@ void TestForEach(const size_t n) thrust::host_vector h_input = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_input[i] = ((size_t) h_input[i]) % output_size; + h_input[i] = static_cast(((size_t) h_input[i]) % output_size); thrust::device_vector d_input = h_input; @@ -232,7 +232,7 @@ void TestForEachN(const size_t n) thrust::host_vector h_input = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_input[i] = ((size_t) h_input[i]) % output_size; + h_input[i] = static_cast(((size_t) h_input[i]) % output_size); thrust::device_vector d_input = h_input; @@ -280,6 +280,7 @@ void _TestForEachWithLargeTypes(void) thrust::host_vector< FixedVector > h_data(n); for(size_t i = 0; i < h_data.size(); i++) + //h_data[i] = static_cast(FixedVector(i)); h_data[i] = FixedVector(i); thrust::device_vector< FixedVector > d_data = h_data; @@ -321,6 +322,7 @@ void _TestForEachNWithLargeTypes(void) thrust::host_vector< FixedVector > h_data(n); for(size_t i = 0; i < h_data.size(); i++) + //h_data[i] = static_cast(FixedVector(i)); h_data[i] = FixedVector(i); thrust::device_vector< FixedVector > d_data = h_data; diff --git a/testing/pair_transform.cu b/testing/pair_transform.cu index 612a77af0..3abb375c3 100644 --- a/testing/pair_transform.cu +++ b/testing/pair_transform.cu @@ -20,7 +20,8 @@ struct add_pairs __host__ __device__ Pair1 operator()(const Pair1 &x, const Pair2 &y) { - return thrust::make_pair(x.first + y.first, x.second + y.second); + return thrust::make_pair(x.first + static_cast(y.first), + x.second + static_cast(y.second)); } // end operator() }; // end add_pairs diff --git a/testing/scan_by_key.cu b/testing/scan_by_key.cu index 8d0cd20b9..842f7be0b 100644 --- a/testing/scan_by_key.cu +++ b/testing/scan_by_key.cu @@ -477,7 +477,7 @@ void TestExclusiveScanByKeyInPlace(const size_t n) thrust::host_vector h_vals = unittest::random_integers(n); for(size_t i = 0; i < n; i++) { - h_vals[i] = static_cast(i % 10); + h_vals[i] = static_cast(i % 10); } thrust::device_vector d_vals = h_vals; From 0d41d3cde455531074fcba3fdde9a75af8167718 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Mon, 31 Jan 2022 08:42:22 -0800 Subject: [PATCH 10/18] batch of changes 2 --- testing/set_intersection_by_key.cu | 2 +- testing/set_union_by_key.cu | 2 +- testing/shuffle.cu | 3 ++- testing/sort_by_key_variable_bits.cu | 2 +- thrust/detail/internal_functional.h | 3 +++ 5 files changed, 8 insertions(+), 4 deletions(-) diff --git a/testing/set_intersection_by_key.cu b/testing/set_intersection_by_key.cu index d82ee04ad..1d8e3bb23 100644 --- a/testing/set_intersection_by_key.cu +++ b/testing/set_intersection_by_key.cu @@ -243,7 +243,7 @@ void TestSetIntersectionByKeyMultiset(const size_t n) { int temp = static_cast(*i); temp %= 13; - *i = temp; + *i = static_cast(temp); } thrust::host_vector h_a_key(vec.begin(), vec.begin() + n); diff --git a/testing/set_union_by_key.cu b/testing/set_union_by_key.cu index 7d58ebf4f..3c458a563 100644 --- a/testing/set_union_by_key.cu +++ b/testing/set_union_by_key.cu @@ -263,7 +263,7 @@ void TestSetUnionByKeyMultiset(const size_t n) { int temp = static_cast(*i); temp %= 13; - *i = temp; + *i = static_cast(temp); } thrust::host_vector h_a_key(vec.begin(), vec.begin() + n); diff --git a/testing/shuffle.cu b/testing/shuffle.cu index a5b1c6f29..0b1c224e3 100644 --- a/testing/shuffle.cu +++ b/testing/shuffle.cu @@ -496,7 +496,8 @@ void TestShuffleUniformPermutation() { ASSERT_EQUAL(permutation_counts.size(), total_permutations); double chi_squared = 0.0; - double expected_count = static_cast(num_samples) / total_permutations; + double expected_count = static_cast(num_samples) / + static_cast(total_permutations); for (auto kv : permutation_counts) { chi_squared += std::pow(expected_count - kv.second, 2) / expected_count; } diff --git a/testing/sort_by_key_variable_bits.cu b/testing/sort_by_key_variable_bits.cu index 379160860..f6f32b50e 100644 --- a/testing/sort_by_key_variable_bits.cu +++ b/testing/sort_by_key_variable_bits.cu @@ -24,7 +24,7 @@ struct TestSortByKeyVariableBits { thrust::host_vector h_keys = unittest::random_integers(n); - const T mask = (1 << num_bits) - 1; + const T mask = static_cast((1 << num_bits) - 1); for(size_t i = 0; i < n; i++) h_keys[i] &= mask; diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h index aae373ded..ae82793ce 100644 --- a/thrust/detail/internal_functional.h +++ b/thrust/detail/internal_functional.h @@ -325,7 +325,10 @@ template // mstack adding static_cast //thrust::get<1>(t) = static_cast(f(thrust::get<0>(t))); //thrust::get<1>(t) = f(static_cast(thrust::get<0>(t))); +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion" thrust::get<1>(t) = f(thrust::get<0>(t)); +#pragma GCC diagnostic pop } }; From e49fd143541d0f5a20b590f7227427c509d78383 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Mon, 31 Jan 2022 08:44:02 -0800 Subject: [PATCH 11/18] batch of changes 3 --- thrust/random/detail/normal_distribution_base.h | 2 +- thrust/system/cuda/detail/async/sort.h | 6 +++--- thrust/system/cuda/detail/reduce.h | 4 ++-- thrust/system/detail/generic/sequence.inl | 2 +- thrust/system/detail/sequential/scan.h | 13 +++++++++++-- 5 files changed, 18 insertions(+), 9 deletions(-) diff --git a/thrust/random/detail/normal_distribution_base.h b/thrust/random/detail/normal_distribution_base.h index f67cb7152..244efc458 100644 --- a/thrust/random/detail/normal_distribution_base.h +++ b/thrust/random/detail/normal_distribution_base.h @@ -65,7 +65,7 @@ template } // Convert to floating point in [0,0.5) - RealType p = u*S1 + S2; + RealType p = static_cast(u*S1 + S2); // Apply inverse error function return mean + stddev * S3 * erfcinv(2 * p); diff --git a/thrust/system/cuda/detail/async/sort.h b/thrust/system/cuda/detail/async/sort.h index e8f92d7f7..7a2b3b8d0 100644 --- a/thrust/system/cuda/detail/async/sort.h +++ b/thrust/system/cuda/detail/async/sort.h @@ -302,9 +302,9 @@ invoke_radix_sort( tmp_ptr , tmp_size , keys - , n + , static_cast(n) , 0 - , sizeof(T) * 8 + , (sizeof(T) * 8 , stream , THRUST_DEBUG_SYNC_FLAG ); @@ -328,7 +328,7 @@ invoke_radix_sort( tmp_ptr , tmp_size , keys - , n + , static_cast(n) , 0 , sizeof(T) * 8 , stream diff --git a/thrust/system/cuda/detail/reduce.h b/thrust/system/cuda/detail/reduce.h index 43c85bd0b..201026368 100644 --- a/thrust/system/cuda/detail/reduce.h +++ b/thrust/system/cuda/detail/reduce.h @@ -950,7 +950,7 @@ T reduce_n_impl(execution_policy& policy, >::Dispatch), num_items, (NULL, tmp_size, first, reinterpret_cast(NULL), - num_items_fixed, binary_op, init, stream, + static_cast(num_items_fixed), binary_op, init, stream, THRUST_DEBUG_SYNC_FLAG)); cuda_cub::throw_on_error(status, "after reduction step 1"); @@ -978,7 +978,7 @@ T reduce_n_impl(execution_policy& policy, >::Dispatch), num_items, (tmp_ptr, tmp_size, first, ret_ptr, - num_items_fixed, binary_op, init, stream, + static_cast(num_items_fixed), binary_op, init, stream, THRUST_DEBUG_SYNC_FLAG)); cuda_cub::throw_on_error(status, "after reduction step 2"); diff --git a/thrust/system/detail/generic/sequence.inl b/thrust/system/detail/generic/sequence.inl index 0fe372931..c995321c9 100644 --- a/thrust/system/detail/generic/sequence.inl +++ b/thrust/system/detail/generic/sequence.inl @@ -75,7 +75,7 @@ struct compute_sequence_value:: __host__ __device__ T operator()(std::size_t i) const { - return init + step * static_cast(i); + return static_cast(init + step * static_cast(i)); } }; } diff --git a/thrust/system/detail/sequential/scan.h b/thrust/system/detail/sequential/scan.h index c5fce2475..ae60ac4ef 100644 --- a/thrust/system/detail/sequential/scan.h +++ b/thrust/system/detail/sequential/scan.h @@ -54,6 +54,8 @@ __host__ __device__ // Use the input iterator's value type per https://wg21.link/P0571 using ValueType = typename thrust::iterator_value::type; + // Use for explicit type conversion + using OutputType = typename thrust::iterator_value::type; // wrap binary_op thrust::detail::wrapped_function< @@ -65,10 +67,13 @@ __host__ __device__ { ValueType sum = *first; + //*result = static_cast(*first); *result = *first; for(++first, ++result; first != last; ++first, ++result) *result = sum = wrapped_binary_op(sum,*first); + //*result = wrapped_binary_op(sum,*first); + //sum = wrapped_binary_op(sum,*first); } return result; @@ -93,10 +98,14 @@ __host__ __device__ // Use the initial value type per https://wg21.link/P0571 using ValueType = InitialValueType; + // Use for explicit type conversion + using InputType = typename thrust::iterator_value::type; + // Use for explicit type conversion + using OutputType = typename thrust::iterator_value::type; if(first != last) { - ValueType tmp = *first; // temporary value allows in-situ scan + ValueType tmp = static_cast(*first); // temporary value allows in-situ scan ValueType sum = init; *result = sum; @@ -104,7 +113,7 @@ __host__ __device__ for(++first, ++result; first != last; ++first, ++result) { - tmp = *first; + tmp = static_cast(*first); *result = sum; sum = binary_op(sum, tmp); } From 7f4663d0a9cfc250263e67bbd5920b26c464b936 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Mon, 31 Jan 2022 08:44:26 -0800 Subject: [PATCH 12/18] batch of changes 4 --- thrust/system/detail/sequential/scan_by_key.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/thrust/system/detail/sequential/scan_by_key.h b/thrust/system/detail/sequential/scan_by_key.h index c428c1050..36e4720cf 100644 --- a/thrust/system/detail/sequential/scan_by_key.h +++ b/thrust/system/detail/sequential/scan_by_key.h @@ -117,7 +117,7 @@ __host__ __device__ // first one is init *result = next; - next = binary_op(next, temp_value); + next = static_cast(binary_op(next, temp_value)); for(++first1, ++first2, ++result; first1 != last1; @@ -132,7 +132,7 @@ __host__ __device__ next = init; // reset sum *result = next; - next = binary_op(next, temp_value); + next = static_cast(binary_op(next, temp_value)); temp_key = key; } From fbb93330aa25989a99774d12dac27613de8a5579 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 9 Feb 2022 11:32:48 -0800 Subject: [PATCH 13/18] adding batch of changes 5 --- testing/cuda/for_each.cu | 4 ++-- testing/cuda/gather.cu | 4 ++-- testing/cuda/scan_by_key.cu | 2 +- testing/cuda/scatter.cu | 4 ++-- testing/for_each.cu | 2 +- 5 files changed, 8 insertions(+), 8 deletions(-) diff --git a/testing/cuda/for_each.cu b/testing/cuda/for_each.cu index be6a7738c..59b0ad828 100644 --- a/testing/cuda/for_each.cu +++ b/testing/cuda/for_each.cu @@ -74,7 +74,7 @@ void TestForEachDeviceSeq(const size_t n) thrust::host_vector h_input = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_input[i] = ((size_t) h_input[i]) % output_size; + h_input[i] = static_cast(((size_t) h_input[i]) % output_size); thrust::device_vector d_input = h_input; @@ -105,7 +105,7 @@ void TestForEachDeviceDevice(const size_t n) thrust::host_vector h_input = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_input[i] = ((size_t) h_input[i]) % output_size; + h_input[i] = static_cast(((size_t) h_input[i]) % output_size); thrust::device_vector d_input = h_input; diff --git a/testing/cuda/gather.cu b/testing/cuda/gather.cu index a9a8c9333..685ca921b 100644 --- a/testing/cuda/gather.cu +++ b/testing/cuda/gather.cu @@ -24,7 +24,7 @@ void TestGatherDevice(ExecutionPolicy exec, const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % source_size; + h_map[i] = h_map[i] % static_cast(source_size); thrust::device_vector d_map = h_map; @@ -117,7 +117,7 @@ void TestGatherIfDevice(ExecutionPolicy exec, const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % source_size; + h_map[i] = h_map[i] % static_cast(source_size); thrust::device_vector d_map = h_map; diff --git a/testing/cuda/scan_by_key.cu b/testing/cuda/scan_by_key.cu index e65560edf..d1de791e4 100644 --- a/testing/cuda/scan_by_key.cu +++ b/testing/cuda/scan_by_key.cu @@ -47,7 +47,7 @@ void TestScanByKeyDevice(ExecutionPolicy exec) thrust::host_vector h_vals = unittest::random_integers(n); for(size_t i = 0; i < n; i++) { - h_vals[i] = i % 10; + h_vals[i] = static_cast(i % 10); } thrust::device_vector d_vals = h_vals; diff --git a/testing/cuda/scatter.cu b/testing/cuda/scatter.cu index 52bd9755f..86d5d8b62 100644 --- a/testing/cuda/scatter.cu +++ b/testing/cuda/scatter.cu @@ -24,7 +24,7 @@ void TestScatterDevice(ExecutionPolicy exec) for(size_t i = 0; i < n; i++) { - h_map[i] = h_map[i] % output_size; + h_map[i] = h_map[i] % static_cast(output_size); } thrust::device_vector d_map = h_map; @@ -82,7 +82,7 @@ void TestScatterIfDevice(ExecutionPolicy exec) for(size_t i = 0; i < n; i++) { - h_map[i] = h_map[i] % output_size; + h_map[i] = h_map[i] % static_cast(output_size); } thrust::device_vector d_map = h_map; diff --git a/testing/for_each.cu b/testing/for_each.cu index e9e0eb921..2c7a5212e 100644 --- a/testing/for_each.cu +++ b/testing/for_each.cu @@ -281,7 +281,7 @@ void _TestForEachWithLargeTypes(void) for(size_t i = 0; i < h_data.size(); i++) //h_data[i] = static_cast(FixedVector(i)); - h_data[i] = FixedVector(i); + h_data[i] = FixedVector(static_cast(i)); thrust::device_vector< FixedVector > d_data = h_data; From eae9b81bd822327909445bd762f997294e468dfa Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 9 Feb 2022 11:33:46 -0800 Subject: [PATCH 14/18] adding batch of changes 6 --- testing/gather.cu | 8 ++++---- testing/reduce.cu | 2 +- testing/scan_by_key.cu | 6 +++--- testing/scatter.cu | 8 ++++---- testing/set_difference_by_key.cu | 2 +- testing/set_symmetric_difference.cu | 2 +- 6 files changed, 14 insertions(+), 14 deletions(-) diff --git a/testing/gather.cu b/testing/gather.cu index c164e44b2..1b2a3b42f 100644 --- a/testing/gather.cu +++ b/testing/gather.cu @@ -89,7 +89,7 @@ void TestGather(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % source_size; + h_map[i] = h_map[i] % static_cast(source_size); thrust::device_vector d_map = h_map; @@ -118,7 +118,7 @@ void TestGatherToDiscardIterator(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % source_size; + h_map[i] = h_map[i] % static_cast(source_size); thrust::device_vector d_map = h_map; @@ -244,7 +244,7 @@ void TestGatherIf(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % source_size; + h_map[i] = h_map[i] % static_cast(source_size); thrust::device_vector d_map = h_map; @@ -282,7 +282,7 @@ void TestGatherIfToDiscardIterator(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % source_size; + h_map[i] = h_map[i] % static_cast(source_size); thrust::device_vector d_map = h_map; diff --git a/testing/reduce.cu b/testing/reduce.cu index cb08bc889..e3c976ea1 100644 --- a/testing/reduce.cu +++ b/testing/reduce.cu @@ -11,7 +11,7 @@ template __host__ __device__ T operator()(T lhs, T rhs) const { - return ((lhs % 10) + (rhs % 10)) % 10; + return static_cast(((lhs % 10) + (rhs % 10)) % 10); } }; diff --git a/testing/scan_by_key.cu b/testing/scan_by_key.cu index 842f7be0b..e9f58a9c5 100644 --- a/testing/scan_by_key.cu +++ b/testing/scan_by_key.cu @@ -376,7 +376,7 @@ void TestInclusiveScanByKey(const size_t n) thrust::host_vector h_vals = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_vals[i] = static_cast(i % 10); + h_vals[i] = static_cast(i % 10); thrust::device_vector d_vals = h_vals; thrust::host_vector h_output(n); @@ -406,7 +406,7 @@ void TestExclusiveScanByKey(const size_t n) thrust::host_vector h_vals = unittest::random_integers(n); for(size_t i = 0; i < n; i++) { - h_vals[i] = static_cast(i % 10); + h_vals[i] = static_cast(i % 10); } thrust::device_vector d_vals = h_vals; @@ -443,7 +443,7 @@ void TestInclusiveScanByKeyInPlace(const size_t n) thrust::host_vector h_vals = unittest::random_integers(n); for(size_t i = 0; i < n; i++) { - h_vals[i] = static_cast(i % 10); + h_vals[i] = static_cast(i % 10); } thrust::device_vector d_vals = h_vals; diff --git a/testing/scatter.cu b/testing/scatter.cu index ffd56f27c..b28557604 100644 --- a/testing/scatter.cu +++ b/testing/scatter.cu @@ -98,7 +98,7 @@ void TestScatter(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % output_size; + h_map[i] = h_map[i] % static_cast(output_size); thrust::device_vector d_map = h_map; @@ -124,7 +124,7 @@ void TestScatterToDiscardIterator(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % output_size; + h_map[i] = h_map[i] % static_cast(output_size); thrust::device_vector d_map = h_map; @@ -241,7 +241,7 @@ void TestScatterIf(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % output_size; + h_map[i] = h_map[i] % static_cast(output_size); thrust::device_vector d_map = h_map; @@ -267,7 +267,7 @@ void TestScatterIfToDiscardIterator(const size_t n) thrust::host_vector h_map = unittest::random_integers(n); for(size_t i = 0; i < n; i++) - h_map[i] = h_map[i] % output_size; + h_map[i] = h_map[i] % static_cast(output_size); thrust::device_vector d_map = h_map; diff --git a/testing/set_difference_by_key.cu b/testing/set_difference_by_key.cu index 29dbb68fc..4de23d12d 100644 --- a/testing/set_difference_by_key.cu +++ b/testing/set_difference_by_key.cu @@ -259,7 +259,7 @@ void TestSetDifferenceByKeyMultiset(const size_t n) { int temp = static_cast(*i); temp %= 13; - *i = temp; + *i = static_cast(temp); } thrust::host_vector h_a_key(vec.begin(), vec.begin() + n); diff --git a/testing/set_symmetric_difference.cu b/testing/set_symmetric_difference.cu index dde145fec..b600271f0 100644 --- a/testing/set_symmetric_difference.cu +++ b/testing/set_symmetric_difference.cu @@ -177,7 +177,7 @@ void TestSetSymmetricDifferenceMultiset(const size_t n) { int temp = static_cast(*i); temp %= 13; - *i = temp; + *i = static_cast(temp); } thrust::host_vector h_a(vec.begin(), vec.begin() + n); From 6e0201f0e35b7fbe9ecc965d8d230bac673a22ea Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 9 Feb 2022 11:34:23 -0800 Subject: [PATCH 15/18] adding batch of changes 7 --- testing/set_symmetric_difference_by_key.cu | 2 +- testing/shuffle.cu | 11 +++++++---- thrust/detail/internal_functional.h | 2 ++ thrust/detail/tuple.inl | 6 ++++-- thrust/iterator/transform_iterator.h | 2 +- 5 files changed, 15 insertions(+), 8 deletions(-) diff --git a/testing/set_symmetric_difference_by_key.cu b/testing/set_symmetric_difference_by_key.cu index 98e416af8..1facd2bfd 100644 --- a/testing/set_symmetric_difference_by_key.cu +++ b/testing/set_symmetric_difference_by_key.cu @@ -263,7 +263,7 @@ void TestSetSymmetricDifferenceByKeyMultiset(const size_t n) { int temp = static_cast(*i); temp %= 13; - *i = temp; + *i = static_cast(temp); } thrust::host_vector h_a_key(vec.begin(), vec.begin() + n); diff --git a/testing/shuffle.cu b/testing/shuffle.cu index 0b1c224e3..5dd72b839 100644 --- a/testing/shuffle.cu +++ b/testing/shuffle.cu @@ -383,7 +383,7 @@ void TestFunctionIsBijection(size_t m) { thrust::system::detail::generic::feistel_bijection host_f(m, host_g); thrust::system::detail::generic::feistel_bijection device_f(m, device_g); - if (host_f.nearest_power_of_two() >= std::numeric_limits::max() || m == 0) { + if (static_cast(host_f.nearest_power_of_two()) >= std::numeric_limits::max() || m == 0) { return; } @@ -449,7 +449,8 @@ void TestShuffleKeyPosition() { double expected_average_position = static_cast(m - 1) / 2; double chi_squared = 0.0; for (auto j = 0ull; j < m; j++) { - double average_position = static_cast(index_sum[j]) / num_samples; + double average_position = static_cast(index_sum[j]) / + static_cast(num_samples); chi_squared += std::pow(expected_average_position - average_position, 2) / expected_average_position; } @@ -499,7 +500,8 @@ void TestShuffleUniformPermutation() { double expected_count = static_cast(num_samples) / static_cast(total_permutations); for (auto kv : permutation_counts) { - chi_squared += std::pow(expected_count - kv.second, 2) / expected_count; + chi_squared += std::pow(expected_count - static_cast(kv.second), 2)/ + expected_count; } double p_score = CephesFunctions::cephes_igamc( (double)(total_permutations - 1) / 2.0, chi_squared / 2.0); @@ -578,7 +580,8 @@ void TestShuffleEvenDistribution() { } } - const double expected_occurances = (double)num_samples / shuffle_size; + const double expected_occurances = static_cast(num_samples) / + static_cast(shuffle_size); for (uint64_t i = 0; i < shuffle_size; i++) { double chi_squared_pos = 0.0; double chi_squared_num = 0.0; diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h index ae82793ce..9fd12bf64 100644 --- a/thrust/detail/internal_functional.h +++ b/thrust/detail/internal_functional.h @@ -201,6 +201,8 @@ template T &lvalue = const_cast(x); // this assigns correctly whether x is a true reference or proxy + // mstack problem area + // lvalue = static_cast(gen()); lvalue = gen(); } diff --git a/thrust/detail/tuple.inl b/thrust/detail/tuple.inl index 48135d540..7519316e7 100644 --- a/thrust/detail/tuple.inl +++ b/thrust/detail/tuple.inl @@ -349,7 +349,8 @@ template template inline __host__ __device__ - cons( const cons& u ) : head(u.head), tail(u.tail) {} + cons( const cons& u ) : head(static_cast(u.head)), tail(u.tail) {} + //cons( const cons& u ) : head(u.head), tail(u.tail) {} #if THRUST_CPP_DIALECT >= 2011 cons(const cons &) = default; @@ -459,7 +460,8 @@ template template inline __host__ __device__ - cons( const cons& u ) : head(u.head) {} + cons( const cons& u ) : head(static_cast(u.head)) {} + //cons( const cons& u ) : head(u.head) {} #if THRUST_CPP_DIALECT >= 2011 cons(const cons &) = default; diff --git a/thrust/iterator/transform_iterator.h b/thrust/iterator/transform_iterator.h index 5afb5f37b..f4102605e 100644 --- a/thrust/iterator/transform_iterator.h +++ b/thrust/iterator/transform_iterator.h @@ -312,7 +312,7 @@ template ::type const& x = *this->base(); - return m_f(x); + return static_cast(m_f(x)); } THRUST_DISABLE_MSVC_WARNING_END(4172) From 631ddceeb22466c6c4519466cd779beb7fef442d Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Wed, 9 Feb 2022 11:35:55 -0800 Subject: [PATCH 16/18] adding batch of changes 8 --- thrust/system/cuda/detail/async/sort.h | 2 +- thrust/system/detail/generic/shuffle.inl | 2 +- .../system/detail/sequential/general_copy.h | 5 +++++ thrust/system/detail/sequential/scan.h | 6 +++--- thrust/system/detail/sequential/scan_by_key.h | 20 +++++++++++++------ 5 files changed, 24 insertions(+), 11 deletions(-) diff --git a/thrust/system/cuda/detail/async/sort.h b/thrust/system/cuda/detail/async/sort.h index 7a2b3b8d0..f6e1bee07 100644 --- a/thrust/system/cuda/detail/async/sort.h +++ b/thrust/system/cuda/detail/async/sort.h @@ -304,7 +304,7 @@ invoke_radix_sort( , keys , static_cast(n) , 0 - , (sizeof(T) * 8 + , sizeof(T) * 8 , stream , THRUST_DEBUG_SYNC_FLAG ); diff --git a/thrust/system/detail/generic/shuffle.inl b/thrust/system/detail/generic/shuffle.inl index 91b77351d..236b225b7 100644 --- a/thrust/system/detail/generic/shuffle.inl +++ b/thrust/system/detail/generic/shuffle.inl @@ -107,7 +107,7 @@ class feistel_bijection { __host__ __device__ round_state do_round(const round_state state, const std::uint64_t round) const { - const std::uint32_t new_left = state.right & left_side_mask; + const std::uint32_t new_left = static_cast(state.right & left_side_mask); const std::uint32_t round_function_res = state.left ^ round_function(state.right, key[round]); if (right_side_bits != left_side_bits) { diff --git a/thrust/system/detail/sequential/general_copy.h b/thrust/system/detail/sequential/general_copy.h index 6ea87bbac..2a5f49bee 100644 --- a/thrust/system/detail/sequential/general_copy.h +++ b/thrust/system/detail/sequential/general_copy.h @@ -72,6 +72,9 @@ typename thrust::detail::enable_if< >::type iter_assign(OutputIterator dst, InputIterator src) { + using OutputType = typename thrust::iterator_traits::value_type; + + //*dst = static_cast(*src); *dst = *src; } @@ -84,9 +87,11 @@ typename thrust::detail::disable_if< >::type iter_assign(OutputIterator dst, InputIterator src) { + //using OutputType = typename thrust::iterator_traits::value_type; typedef typename thrust::iterator_value::type value_type; // insert a temporary and hope for the best + //*dst = static_cast(*src); *dst = static_cast(*src); } diff --git a/thrust/system/detail/sequential/scan.h b/thrust/system/detail/sequential/scan.h index ae60ac4ef..290544044 100644 --- a/thrust/system/detail/sequential/scan.h +++ b/thrust/system/detail/sequential/scan.h @@ -67,7 +67,7 @@ __host__ __device__ { ValueType sum = *first; - //*result = static_cast(*first); + *result = static_cast(*first); *result = *first; for(++first, ++result; first != last; ++first, ++result) @@ -109,13 +109,13 @@ __host__ __device__ ValueType sum = init; *result = sum; - sum = binary_op(sum, tmp); + sum = static_cast(binary_op(sum, tmp)); for(++first, ++result; first != last; ++first, ++result) { tmp = static_cast(*first); *result = sum; - sum = binary_op(sum, tmp); + sum = static_cast(binary_op(sum, tmp)); } } diff --git a/thrust/system/detail/sequential/scan_by_key.h b/thrust/system/detail/sequential/scan_by_key.h index 36e4720cf..ed07c3d36 100644 --- a/thrust/system/detail/sequential/scan_by_key.h +++ b/thrust/system/detail/sequential/scan_by_key.h @@ -53,6 +53,7 @@ __host__ __device__ { using KeyType = typename thrust::iterator_traits::value_type; using ValueType = typename thrust::iterator_traits::value_type; + using OutputType = typename thrust::iterator_traits::value_type; // wrap binary_op thrust::detail::wrapped_function< @@ -65,7 +66,7 @@ __host__ __device__ KeyType prev_key = *first1; ValueType prev_value = *first2; - *result = prev_value; + *result = static_cast(prev_value); for(++first1, ++first2, ++result; first1 != last1; @@ -73,11 +74,18 @@ __host__ __device__ { KeyType key = *first1; - if(binary_pred(prev_key, key)) - *result = prev_value = wrapped_binary_op(prev_value,*first2); - else - *result = prev_value = *first2; - + if(binary_pred(prev_key, key)) { + //*result = static_cast(prev_value) = wrapped_binary_op(prev_value,*first2); + //*result = prev_value = wrapped_binary_op(prev_value,*first2); + *result = static_cast(wrapped_binary_op(prev_value,*first2)); + prev_value = static_cast(wrapped_binary_op(prev_value,*first2)); + } + else { + //*result = static_cast(prev_value) = *first2; + //*result = prev_value = *first2; + *result = static_cast(*first2); + prev_value = static_cast(*first2); + } prev_key = key; } } From 22d16e97be886a2cec861ae51f7cc205f2845302 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Tue, 15 Feb 2022 10:19:37 -0800 Subject: [PATCH 17/18] batch of changes 9 --- testing/async/inclusive_scan/counting_iterator.cu | 2 +- thrust/iterator/iterator_adaptor.h | 3 ++- thrust/iterator/transform_iterator.h | 3 ++- thrust/system/detail/sequential/scan.h | 2 +- thrust/system/detail/sequential/scan_by_key.h | 2 +- 5 files changed, 7 insertions(+), 5 deletions(-) diff --git a/testing/async/inclusive_scan/counting_iterator.cu b/testing/async/inclusive_scan/counting_iterator.cu index fe9fdeb80..f156f3833 100644 --- a/testing/async/inclusive_scan/counting_iterator.cu +++ b/testing/async/inclusive_scan/counting_iterator.cu @@ -33,7 +33,7 @@ struct test_counting_iterator { void operator()(std::size_t num_values) const { - num_values = unittest::truncate_to_max_representable(num_values); + num_values = static_cast(unittest::truncate_to_max_representable(num_values)); testing::async::test_policy_overloads>::run(num_values); } }; diff --git a/thrust/iterator/iterator_adaptor.h b/thrust/iterator/iterator_adaptor.h index 67d4866b9..ce942e926 100644 --- a/thrust/iterator/iterator_adaptor.h +++ b/thrust/iterator/iterator_adaptor.h @@ -203,7 +203,8 @@ template(m_iterator + n); + //m_iterator = static_cast(m_iterator + n); + m_iterator = m_iterator + static_cast(n); } __thrust_exec_check_disable__ diff --git a/thrust/iterator/transform_iterator.h b/thrust/iterator/transform_iterator.h index f4102605e..b44782820 100644 --- a/thrust/iterator/transform_iterator.h +++ b/thrust/iterator/transform_iterator.h @@ -312,7 +312,8 @@ template ::type const& x = *this->base(); - return static_cast(m_f(x)); + // The issue is that x needs to be the type specified in m_f, but that is unknown + return m_f(x); } THRUST_DISABLE_MSVC_WARNING_END(4172) diff --git a/thrust/system/detail/sequential/scan.h b/thrust/system/detail/sequential/scan.h index 290544044..193e59bd7 100644 --- a/thrust/system/detail/sequential/scan.h +++ b/thrust/system/detail/sequential/scan.h @@ -67,7 +67,7 @@ __host__ __device__ { ValueType sum = *first; - *result = static_cast(*first); + *result = static_cast(*first); *result = *first; for(++first, ++result; first != last; ++first, ++result) diff --git a/thrust/system/detail/sequential/scan_by_key.h b/thrust/system/detail/sequential/scan_by_key.h index ed07c3d36..270e01484 100644 --- a/thrust/system/detail/sequential/scan_by_key.h +++ b/thrust/system/detail/sequential/scan_by_key.h @@ -66,7 +66,7 @@ __host__ __device__ KeyType prev_key = *first1; ValueType prev_value = *first2; - *result = static_cast(prev_value); + *result = static_cast(prev_value); for(++first1, ++first2, ++result; first1 != last1; From 78222444d1694f8aad8894ead33639de8dc69eb0 Mon Sep 17 00:00:00 2001 From: Matt Stack Date: Mon, 28 Feb 2022 08:34:51 -0800 Subject: [PATCH 18/18] batch of changes 10 --- thrust/iterator/iterator_adaptor.h | 4 ++-- thrust/system/detail/sequential/scan.h | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/thrust/iterator/iterator_adaptor.h b/thrust/iterator/iterator_adaptor.h index ce942e926..505be86d8 100644 --- a/thrust/iterator/iterator_adaptor.h +++ b/thrust/iterator/iterator_adaptor.h @@ -203,8 +203,8 @@ template(m_iterator + n); - m_iterator = m_iterator + static_cast(n); + m_iterator = static_cast(m_iterator + n); + //m_iterator = m_iterator + static_cast(n); } __thrust_exec_check_disable__ diff --git a/thrust/system/detail/sequential/scan.h b/thrust/system/detail/sequential/scan.h index 193e59bd7..290544044 100644 --- a/thrust/system/detail/sequential/scan.h +++ b/thrust/system/detail/sequential/scan.h @@ -67,7 +67,7 @@ __host__ __device__ { ValueType sum = *first; - *result = static_cast(*first); + *result = static_cast(*first); *result = *first; for(++first, ++result; first != last; ++first, ++result)