MVP for disabling nvtx ranges for thrust::seq

thrust/thrust/detail/copy.inl

@@ -26,6 +26,7 @@
 #  pragma system_header
 #endif // no system header
 #include <thrust/detail/copy.h>
+#include <thrust/detail/nvtx_policy.h>
 #include <thrust/system/detail/generic/select_system.h>
 
 // Include all active backend system implementations (generic, sequential, host and device)
@@ -52,7 +53,7 @@ copy(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
      InputIterator last,
      OutputIterator result)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::copy");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::copy");
   using thrust::system::detail::generic::copy;
   return copy(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, result);
 } // end copy()
@@ -62,7 +63,7 @@ template <typename DerivedPolicy, typename InputIterator, typename Size, typenam
 _CCCL_HOST_DEVICE OutputIterator copy_n(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec, InputIterator first, Size n, OutputIterator result)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::copy_n");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::copy_n");
   using thrust::system::detail::generic::copy_n;
   return copy_n(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, n, result);
 } // end copy_n()
@@ -78,7 +79,8 @@ _CCCL_HOST_DEVICE OutputIterator two_system_copy(
   InputIterator last,
   OutputIterator result)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::two_system_copy");
+  _CCCL_NVTX_RANGE_SCOPE_IF(should_enable_nvtx_for_policy<System1>() || should_enable_nvtx_for_policy<System2>(),
+                            "thrust::two_system_copy");
   using thrust::system::detail::generic::select_system;
 
   return thrust::copy(
@@ -98,7 +100,8 @@ _CCCL_HOST_DEVICE OutputIterator two_system_copy_n(
   Size n,
   OutputIterator result)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::two_system_copy_n");
+  _CCCL_NVTX_RANGE_SCOPE_IF(should_enable_nvtx_for_policy<System1>() || should_enable_nvtx_for_policy<System2>(),
+                            "thrust::two_system_copy_n");
   using thrust::system::detail::generic::select_system;
 
   return thrust::copy_n(
@@ -113,9 +116,11 @@ _CCCL_HOST_DEVICE OutputIterator two_system_copy_n(
 template <typename InputIterator, typename OutputIterator>
 OutputIterator copy(InputIterator first, InputIterator last, OutputIterator result)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::copy");
   using System1 = typename thrust::iterator_system<InputIterator>::type;
   using System2 = typename thrust::iterator_system<OutputIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(
+    detail::should_enable_nvtx_for_policy<System1>() || detail::should_enable_nvtx_for_policy<System2>(),
+    "thrust::copy");
 
   System1 system1;
   System2 system2;
@@ -126,9 +131,11 @@ OutputIterator copy(InputIterator first, InputIterator last, OutputIterator resu
 template <typename InputIterator, typename Size, typename OutputIterator>
 OutputIterator copy_n(InputIterator first, Size n, OutputIterator result)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::copy_n");
   using System1 = typename thrust::iterator_system<InputIterator>::type;
   using System2 = typename thrust::iterator_system<OutputIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(
+    detail::should_enable_nvtx_for_policy<System1>() || detail::should_enable_nvtx_for_policy<System2>(),
+    "thrust::copy_n");
 
   System1 system1;
   System2 system2;

thrust/thrust/detail/fill.inl

@@ -26,6 +26,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <thrust/detail/nvtx_policy.h>
 #include <thrust/fill.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
@@ -54,7 +55,7 @@ fill(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
      ForwardIterator last,
      const T& value)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::fill");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::fill");
   using thrust::system::detail::generic::fill;
   return fill(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, value);
 } // end fill()
@@ -64,18 +65,17 @@ template <typename DerivedPolicy, typename OutputIterator, typename Size, typena
 _CCCL_HOST_DEVICE OutputIterator
 fill_n(const thrust::detail::execution_policy_base<DerivedPolicy>& exec, OutputIterator first, Size n, const T& value)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::fill_n");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::fill_n");
   using thrust::system::detail::generic::fill_n;
   return fill_n(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, n, value);
 } // end fill_n()
 
 template <typename ForwardIterator, typename T>
 _CCCL_HOST_DEVICE void fill(ForwardIterator first, ForwardIterator last, const T& value)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::fill");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<ForwardIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "thrust::fill");
+  using thrust::system::detail::generic::select_system;
 
   System system;
 
@@ -85,10 +85,9 @@ _CCCL_HOST_DEVICE void fill(ForwardIterator first, ForwardIterator last, const T
 template <typename OutputIterator, typename Size, typename T>
 _CCCL_HOST_DEVICE OutputIterator fill_n(OutputIterator first, Size n, const T& value)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::fill_n");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<OutputIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "thrust::fill_n");
+  using thrust::system::detail::generic::select_system;
 
   System system;
 

thrust/thrust/detail/for_each.inl

@@ -25,6 +25,7 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
+#include <thrust/detail/nvtx_policy.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
@@ -53,7 +54,7 @@ _CCCL_HOST_DEVICE InputIterator for_each(
   InputIterator last,
   UnaryFunction f)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::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);
@@ -62,9 +63,9 @@ _CCCL_HOST_DEVICE InputIterator for_each(
 template <typename InputIterator, typename UnaryFunction>
 InputIterator for_each(InputIterator first, InputIterator last, UnaryFunction f)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::for_each");
-  using thrust::system::detail::generic::select_system;
   using System = typename thrust::iterator_system<InputIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "thrust::for_each");
+  using thrust::system::detail::generic::select_system;
 
   System system;
   return thrust::for_each(select_system(system), first, last, f);
@@ -75,7 +76,7 @@ template <typename DerivedPolicy, typename InputIterator, typename Size, typenam
 _CCCL_HOST_DEVICE InputIterator for_each_n(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec, InputIterator first, Size n, UnaryFunction f)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::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);
@@ -84,10 +85,9 @@ _CCCL_HOST_DEVICE InputIterator for_each_n(
 template <typename InputIterator, typename Size, typename UnaryFunction>
 InputIterator for_each_n(InputIterator first, Size n, UnaryFunction f)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::for_each_n");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<InputIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "thrust::for_each_n");
+  using thrust::system::detail::generic::select_system;
 
   System system;
   return thrust::for_each_n(select_system(system), first, n, f);

thrust/thrust/detail/generate.inl

@@ -26,6 +26,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <thrust/detail/nvtx_policy.h>
 #include <thrust/generate.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
@@ -54,7 +55,7 @@ generate(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
          ForwardIterator last,
          Generator gen)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::generate");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::generate");
   using thrust::system::detail::generic::generate;
   return generate(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, gen);
 } // end generate()
@@ -64,18 +65,17 @@ template <typename DerivedPolicy, typename OutputIterator, typename Size, typena
 _CCCL_HOST_DEVICE OutputIterator generate_n(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec, OutputIterator first, Size n, Generator gen)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::generate_n");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::generate_n");
   using thrust::system::detail::generic::generate_n;
   return generate_n(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, n, gen);
 } // end generate_n()
 
 template <typename ForwardIterator, typename Generator>
 void generate(ForwardIterator first, ForwardIterator last, Generator gen)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::generate");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<ForwardIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "thrust::generate");
+  using thrust::system::detail::generic::select_system;
 
   System system;
 
@@ -85,10 +85,9 @@ void generate(ForwardIterator first, ForwardIterator last, Generator gen)
 template <typename OutputIterator, typename Size, typename Generator>
 OutputIterator generate_n(OutputIterator first, Size n, Generator gen)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::generate_n");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<OutputIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "thrust::generate_n");
+  using thrust::system::detail::generic::select_system;
 
   System system;
 

thrust/thrust/detail/nvtx_policy.h

@@ -0,0 +1,44 @@
+// SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#pragma once
+
+#include <thrust/detail/config.h>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__type_traits/decay.h>
+#include <cuda/std/__type_traits/is_base_of.h>
+
+THRUST_NAMESPACE_BEGIN
+
+// Forward declarations
+namespace system::detail::sequential
+{
+struct tag;
+template <class>
+struct execution_policy;
+} // namespace system::detail::sequential
+
+namespace detail
+{
+
+// Helper to determine if NVTX should be enabled for a given policy
+// NVTX is DISABLED only for thrust::seq and any policy derived from sequential::execution_policy
+// ENABLED for all other policies (CUDA, OMP, TBB, etc.)
+template <typename DerivedPolicy>
+inline constexpr bool should_enable_nvtx_for_policy()
+{
+  using Policy = ::cuda::std::decay_t<DerivedPolicy>;
+  // This catches thrust::seq, cpp::tag, and any other sequential-based policy
+  return !::cuda::std::is_base_of_v<thrust::system::detail::sequential::execution_policy<Policy>, Policy>;
+}
+
+} // namespace detail
+
+THRUST_NAMESPACE_END

thrust/thrust/detail/reduce.inl

@@ -26,6 +26,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <thrust/detail/nvtx_policy.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/reduce.h>
 #include <thrust/system/detail/generic/select_system.h>
@@ -59,7 +60,7 @@ template <typename DerivedPolicy, typename InputIterator>
 _CCCL_HOST_DEVICE detail::it_value_t<InputIterator>
 reduce(const thrust::detail::execution_policy_base<DerivedPolicy>& exec, InputIterator first, InputIterator last)
 {
-  _CCCL_NVTX_RANGE_SCOPE("thrust::reduce");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "thrust::reduce");
   using thrust::system::detail::generic::reduce;
   return reduce(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last);
 } // end reduce()

thrust/thrust/detail/uninitialized_fill.inl

@@ -26,6 +26,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <thrust/detail/nvtx_policy.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/generic/select_system.h>
 #include <thrust/uninitialized_fill.h>
@@ -54,7 +55,7 @@ _CCCL_HOST_DEVICE void uninitialized_fill(
   ForwardIterator last,
   const T& x)
 {
-  _CCCL_NVTX_RANGE_SCOPE("uninitialized_fill");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "uninitialized_fill");
   using thrust::system::detail::generic::uninitialized_fill;
   return uninitialized_fill(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, x);
 } // end uninitialized_fill()
@@ -64,18 +65,17 @@ template <typename DerivedPolicy, typename ForwardIterator, typename Size, typen
 _CCCL_HOST_DEVICE ForwardIterator uninitialized_fill_n(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec, ForwardIterator first, Size n, const T& x)
 {
-  _CCCL_NVTX_RANGE_SCOPE("uninitialized_fill_n");
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<DerivedPolicy>(), "uninitialized_fill_n");
   using thrust::system::detail::generic::uninitialized_fill_n;
   return uninitialized_fill_n(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, n, x);
 } // end uninitialized_fill_n()
 
 template <typename ForwardIterator, typename T>
 void uninitialized_fill(ForwardIterator first, ForwardIterator last, const T& x)
 {
-  _CCCL_NVTX_RANGE_SCOPE("uninitialized_fill");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<ForwardIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "uninitialized_fill");
+  using thrust::system::detail::generic::select_system;
 
   System system;
 
@@ -85,10 +85,9 @@ void uninitialized_fill(ForwardIterator first, ForwardIterator last, const T& x)
 template <typename ForwardIterator, typename Size, typename T>
 ForwardIterator uninitialized_fill_n(ForwardIterator first, Size n, const T& x)
 {
-  _CCCL_NVTX_RANGE_SCOPE("uninitialized_fill_n");
-  using thrust::system::detail::generic::select_system;
-
   using System = typename thrust::iterator_system<ForwardIterator>::type;
+  _CCCL_NVTX_RANGE_SCOPE_IF(detail::should_enable_nvtx_for_policy<System>(), "uninitialized_fill_n");
+  using thrust::system::detail::generic::select_system;
 
   System system;