performance-unnecessary-value-param

.clang-tidy

@@ -22,7 +22,16 @@ Checks:
   - '-*'
   - 'performance-*'
   - '-performance-noexcept-swap'
+  # TODO(jfaibussowit): re-enable for CCCL 4.0
   - '-performance-enum-size'
+  # The majority of thrust/CUB require value semantics because they may be used on device,
+  # where converting to const-ref is detrimental to performance because of possible
+  # register spilling. Also, for device functions, the overwhelming majority of types are
+  # PODs or trivially copyable, so moving them (another potential way to fix this) gives
+  # no additional perf improvement and only leads to more verbosity (not to mention
+  # compile-time instantiating all those move templates).
+  #
+  # All this is to say it probably is not worth the hassle of enabling this.
   - '-performance-unnecessary-value-param'
   # END REMOVE ME
   # HICPP is 99% aliased to other checks (mostly modernize-* and bugprone-*). We don't

c/parallel/test/test_util.h

@@ -917,7 +917,8 @@ struct name_source_t
 };
 
 template <class ValueT, class StateT>
-iterator_t<ValueT, StateT> make_iterator(name_source_t state, operation_t advance, operation_t dereference)
+iterator_t<ValueT, StateT>
+make_iterator(const name_source_t& state, const operation_t& advance, const operation_t& dereference)
 {
   iterator_t<ValueT, StateT> it;
   it.state_name                = state.name;
@@ -977,7 +978,7 @@ inline std::tuple<std::string, std::string, std::string> make_random_access_iter
 
 template <class ValueT>
 iterator_t<ValueT, random_access_iterator_state_t<ValueT>> make_random_access_iterator(
-  iterator_kind kind, std::string_view value_type, std::string prefix = "", std::string transform = "")
+  iterator_kind kind, std::string_view value_type, std::string_view prefix = "", std::string_view transform = "")
 {
   std::string iterator_state_name = std::format("{0}state_t", prefix);
   std::string advance_fn_name     = std::format("{0}advance", prefix);

cub/.clang-tidy

@@ -1,21 +1,5 @@
 ---
-Checks:
-  # TODO: BEGIN REMOVE ME
-  - '-*'
-  - 'performance-implicit-conversion-in-loop'
-  # END REMOVE ME
-      # 'modernize-*,
-      #  -modernize-use-equals-default,
-      #  -modernize-concat-nested-namespaces,
-      #  -modernize-use-trailing-return-type'
-
-      # -modernize-use-equals-default        # auto-fix is broken (doesn't insert =default correctly)
-      # -modernize-concat-nested-namespaces  # auto-fix is broken (can delete code)
-      # -modernize-use-trailing-return-type  # just a preference
-
-WarningsAsErrors: ''
-HeaderFilterRegex: ''
-FormatStyle:     none
+InheritParentConfig: true
 CheckOptions:
  - key:             modernize-loop-convert.MaxCopySize
    value:           '16'

cudax/test/execution/common/checked_receiver.cuh

@@ -29,7 +29,7 @@ struct checked_value_receiver
   using receiver_concept = cudax_async::receiver_t;
 
   _CCCL_HOST_DEVICE checked_value_receiver(Values... values)
-      : _values{values...}
+      : _values{::cuda::std::move(values)...}
   {}
 
   _CCCL_HOST_DEVICE checked_value_receiver(checked_value_receiver&& other) noexcept
@@ -106,7 +106,7 @@ struct checked_error_receiver
   }
 
   template <class Ty>
-  _CCCL_HOST_DEVICE void set_error(Ty ty) && noexcept
+  _CCCL_HOST_DEVICE void set_error(const Ty& ty) && noexcept
   {
     if constexpr (::cuda::std::is_same_v<Error, Ty>)
     {

cudax/test/execution/test_bulk.cu

@@ -147,7 +147,7 @@ void bulk_keeps_values_type_from_input_sender()
   check_value_types<types<double>>(ex::just(4.2) //
                                    | ex::bulk(ex::par, n, [] _CCCL_HOST_DEVICE(int, double) {}));
   check_value_types<types<double, string>>(ex::just(4.2, string{}) //
-                                           | ex::bulk(ex::par, n, [] _CCCL_HOST_DEVICE(int, double, string) {}));
+                                           | ex::bulk(ex::par, n, [] _CCCL_HOST_DEVICE(int, double, const string&) {}));
 }
 
 void bulk_chunked_keeps_values_type_from_input_sender()
@@ -158,7 +158,7 @@ void bulk_chunked_keeps_values_type_from_input_sender()
   check_value_types<types<double>>(ex::just(4.2) //
                                    | ex::bulk_chunked(ex::par, n, [] _CCCL_HOST_DEVICE(int, int, double) {}));
   check_value_types<types<double, string>>(
-    ex::just(4.2, string{}) | ex::bulk_chunked(ex::par, n, [] _CCCL_HOST_DEVICE(int, int, double, string) {}));
+    ex::just(4.2, string{}) | ex::bulk_chunked(ex::par, n, [] _CCCL_HOST_DEVICE(int, int, double, const string&) {}));
 }
 
 void bulk_unchunked_keeps_values_type_from_input_sender()
@@ -170,7 +170,7 @@ void bulk_unchunked_keeps_values_type_from_input_sender()
                                    | ex::bulk_unchunked(ex::par, n, [] _CCCL_HOST_DEVICE(int, double) {}));
   check_value_types<types<double, string>>(
     ex::just(4.2, string{}) //
-    | ex::bulk_unchunked(ex::par, n, [] _CCCL_HOST_DEVICE(int, double, string) {}));
+    | ex::bulk_unchunked(ex::par, n, [] _CCCL_HOST_DEVICE(int, double, const string&) {}));
 }
 
 void bulk_keeps_error_types_from_input_sender()

cudax/test/execution/test_then.cu

@@ -268,7 +268,7 @@ C2H_TEST("then can be customized early", "[adaptors][then]")
 {
   // The customization will return a different value
   dummy_scheduler<then_test_domain> sched;
-  auto snd = ex::just(string{"hello"}) | ex::continues_on(sched) | ex::then([](string x) {
+  auto snd = ex::just(string{"hello"}) | ex::continues_on(sched) | ex::then([](const string& x) {
                return x + ", world";
              });
   wait_for_value(std::move(snd), string{"ciao"});
@@ -279,7 +279,7 @@ C2H_TEST("then can be customized late", "[adaptors][then]")
   // The customization will return a different value
   dummy_scheduler<then_test_domain> sched;
   auto snd = ex::just(string{"hello"})
-           | ex::on(sched, ex::then([](string x) {
+           | ex::on(sched, ex::then([](const string& x) {
                       return x + ", world";
                     }))
            | ex::write_env(ex::prop{ex::get_scheduler, dummy_scheduler()});

libcudacxx/test/libcudacxx/std/numerics/numeric.ops/transform.reduce/pstl_transform_reduce_binary.cu

@@ -43,7 +43,7 @@
 inline constexpr int size = 1000;
 
 template <class Policy, class Iter>
-void test_transform_reduce(const Policy policy, const thrust::device_vector<int>& Input1, Iter input2)
+void test_transform_reduce(const Policy& policy, const thrust::device_vector<int>& Input1, Iter input2)
 {
   // N * (N + 1) / 2 for the first N integrals
   // 0 for multiplying by 1

libcudacxx/test/libcudacxx/std/numerics/numeric.ops/transform.reduce/pstl_transform_reduce_unary.cu

@@ -45,7 +45,7 @@ struct plus_one
 };
 
 template <class Policy, class Iter>
-void test_transform_reduce(const Policy policy, Iter input1)
+void test_transform_reduce(const Policy& policy, Iter input1)
 {
   // N * (N + 1) / 2 for the first N integrals
   // N for plus_one

nvbench_helper/nvbench_helper/nvbench_helper.cuh

@@ -182,7 +182,7 @@ NVBENCH_DECLARE_TYPE_STRINGS(bit_entropy, "BE", "bit entropy");
   }
 }
 
-[[nodiscard]] inline bit_entropy str_to_entropy(std::string str)
+[[nodiscard]] inline bit_entropy str_to_entropy(const std::string& str)
 {
   if (str == "1.000")
   {

thrust/testing/copy.cu

@@ -213,7 +213,7 @@ DECLARE_VECTOR_UNITTEST(TestCopyListTo);
 template <typename T>
 struct is_even
 {
-  _CCCL_HOST_DEVICE bool operator()(T x)
+  _CCCL_HOST_DEVICE bool operator()(const T& x)
   {
     return (x & 1) == 0;
   }
@@ -222,7 +222,7 @@ struct is_even
 template <typename T>
 struct is_true
 {
-  _CCCL_HOST_DEVICE bool operator()(T x)
+  _CCCL_HOST_DEVICE bool operator()(const T& x)
   {
     return x ? true : false;
   }
@@ -231,7 +231,7 @@ struct is_true
 template <typename T>
 struct mod_3
 {
-  _CCCL_HOST_DEVICE unsigned int operator()(T x)
+  _CCCL_HOST_DEVICE unsigned int operator()(const T& x)
   {
     return x % 3;
   }

thrust/testing/generate.cu

@@ -2,6 +2,8 @@
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 
+#include <cuda/std/utility>
+
 #include <unittest/unittest.h>
 
 _CCCL_DIAG_PUSH
@@ -12,9 +14,9 @@ struct return_value
 {
   T val;
 
-  return_value() {}
-  return_value(T v)
-      : val(v)
+  _CCCL_HOST_DEVICE return_value() {};
+  _CCCL_HOST_DEVICE return_value(T v)
+      : val(::cuda::std::move(v))
   {}
 
   _CCCL_HOST_DEVICE T operator()(void)

thrust/testing/unittest/assertions.h

@@ -117,7 +117,7 @@ struct value_type<THRUST_NS_QUALIFIER::device_reference<T>>
 ////
 // check scalar values
 template <typename T1, typename T2>
-void assert_equal(T1 a, T2 b, const std::string& filename = "unknown", int lineno = -1)
+void assert_equal(const T1& a, const T2& b, const std::string& filename = "unknown", int lineno = -1)
 {
   if (!(a == b))
   {

thrust/testing/unittest/special_types.h

@@ -2,6 +2,8 @@
 
 #include <thrust/execution_policy.h>
 
+#include <cuda/std/utility>
+
 #include <iosfwd>
 
 template <typename T, unsigned int N>
@@ -81,8 +83,8 @@ struct key_value
   {}
 
   _CCCL_HOST_DEVICE key_value(key_type k, value_type v)
-      : key(k)
-      , value(v)
+      : key(::cuda::std::move(k))
+      , value(::cuda::std::move(v))
   {}
 
   _CCCL_HOST_DEVICE bool operator<(const key_value& rhs) const

thrust/thrust/detail/copy_if.inl

@@ -16,6 +16,8 @@
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
 
+#include <cuda/std/__utility/move.h>
+
 // Include all active backend system implementations (generic, sequential, host and device)
 #include <thrust/system/detail/generic/copy_if.h>
 #include <thrust/system/detail/sequential/copy_if.h>
@@ -43,7 +45,8 @@ _CCCL_HOST_DEVICE OutputIterator copy_if(
 {
   _CCCL_NVTX_RANGE_SCOPE("thrust::copy_if");
   using thrust::system::detail::generic::copy_if;
-  return copy_if(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, result, pred);
+  return copy_if(
+    thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, result, ::cuda::std::move(pred));
 } // end copy_if()
 
 _CCCL_EXEC_CHECK_DISABLE
@@ -62,7 +65,13 @@ _CCCL_HOST_DEVICE OutputIterator copy_if(
 {
   _CCCL_NVTX_RANGE_SCOPE("thrust::copy_if");
   using thrust::system::detail::generic::copy_if;
-  return copy_if(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, stencil, result, pred);
+  return copy_if(
+    thrust::detail::derived_cast(thrust::detail::strip_const(exec)),
+    first,
+    last,
+    stencil,
+    result,
+    ::cuda::std::move(pred));
 } // end copy_if()
 
 template <typename InputIterator, typename OutputIterator, typename Predicate>
@@ -77,7 +86,7 @@ OutputIterator copy_if(InputIterator first, InputIterator last, OutputIterator r
   System1 system1;
   System2 system2;
 
-  return thrust::copy_if(select_system(system1, system2), first, last, result, pred);
+  return thrust::copy_if(select_system(system1, system2), first, last, result, ::cuda::std::move(pred));
 } // end copy_if()
 
 template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename Predicate>
@@ -95,7 +104,8 @@ copy_if(InputIterator1 first, InputIterator1 last, InputIterator2 stencil, Outpu
   System2 system2;
   System3 system3;
 
-  return thrust::copy_if(select_system(system1, system2, system3), first, last, stencil, result, pred);
+  return thrust::copy_if(
+    select_system(system1, system2, system3), first, last, stencil, result, ::cuda::std::move(pred));
 } // end copy_if()
 
 THRUST_NAMESPACE_END

thrust/thrust/detail/find.inl

@@ -15,6 +15,8 @@
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
 
+#include <cuda/std/__utility/move.h>
+
 // Include all active backend system implementations (generic, sequential, host and device)
 #include <thrust/system/detail/generic/find.h>
 #include <thrust/system/detail/sequential/find.h>
@@ -54,7 +56,7 @@ _CCCL_HOST_DEVICE InputIterator find_if(
 {
   _CCCL_NVTX_RANGE_SCOPE("thrust::find_if");
   using thrust::system::detail::generic::find_if;
-  return find_if(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, pred);
+  return find_if(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, ::cuda::std::move(pred));
 } // end find_if()
 
 _CCCL_EXEC_CHECK_DISABLE
@@ -67,7 +69,8 @@ _CCCL_HOST_DEVICE InputIterator find_if_not(
 {
   _CCCL_NVTX_RANGE_SCOPE("thrust::find_if_not");
   using thrust::system::detail::generic::find_if_not;
-  return find_if_not(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, pred);
+  return find_if_not(
+    thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, ::cuda::std::move(pred));
 } // end find_if_not()
 
 template <typename InputIterator, typename T>
@@ -93,7 +96,7 @@ InputIterator find_if(InputIterator first, InputIterator last, Predicate pred)
 
   System system;
 
-  return thrust::find_if(select_system(system), first, last, pred);
+  return thrust::find_if(select_system(system), first, last, ::cuda::std::move(pred));
 }
 
 template <typename InputIterator, typename Predicate>
@@ -106,7 +109,7 @@ InputIterator find_if_not(InputIterator first, InputIterator last, Predicate pre
 
   System system;
 
-  return thrust::find_if_not(select_system(system), first, last, pred);
+  return thrust::find_if_not(select_system(system), first, last, ::cuda::std::move(pred));
 }
 
 THRUST_NAMESPACE_END

thrust/thrust/detail/for_each.inl

@@ -17,6 +17,8 @@
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
 
+#include <cuda/std/__utility/move.h>
+
 // Include all active backend system implementations (generic, sequential, host and device)
 #include <thrust/system/detail/generic/for_each.h>
 #include <thrust/system/detail/sequential/for_each.h>
@@ -44,7 +46,7 @@ _CCCL_HOST_DEVICE InputIterator for_each(
   _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::for_each");
   using thrust::system::detail::generic::for_each;
 
-  return for_each(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, f);
+  return for_each(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, ::cuda::std::move(f));
 }
 
 template <typename InputIterator, typename UnaryFunction>
@@ -55,7 +57,7 @@ InputIterator for_each(InputIterator first, InputIterator last, UnaryFunction f)
   using thrust::system::detail::generic::select_system;
 
   System system;
-  return thrust::for_each(select_system(system), first, last, f);
+  return thrust::for_each(select_system(system), first, last, ::cuda::std::move(f));
 } // end for_each()
 
 _CCCL_EXEC_CHECK_DISABLE
@@ -66,7 +68,7 @@ _CCCL_HOST_DEVICE InputIterator for_each_n(
   _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::for_each_n");
   using thrust::system::detail::generic::for_each_n;
 
-  return for_each_n(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, n, f);
+  return for_each_n(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, n, ::cuda::std::move(f));
 } // end for_each_n()
 
 template <typename InputIterator, typename Size, typename UnaryFunction>
@@ -77,7 +79,7 @@ InputIterator for_each_n(InputIterator first, Size n, UnaryFunction f)
   using thrust::system::detail::generic::select_system;
 
   System system;
-  return thrust::for_each_n(select_system(system), first, n, f);
+  return thrust::for_each_n(select_system(system), first, n, ::cuda::std::move(f));
 } // end for_each_n()
 
 THRUST_NAMESPACE_END