Merge branch 'main' into stf_unstable_unique_utility_move

Author: caugonnet

cudax/include/cuda/experimental/__stf/internal/cooperative_group_system.cuh

@@ -20,7 +20,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/experimental/__stf/utility/cuda_attributes.cuh>
 #if _CCCL_CUDA_COMPILATION()
 #  include <cooperative_groups.h>
 #endif // _CCCL_CUDA_COMPILATION()

cudax/include/cuda/experimental/__stf/internal/execution_policy.cuh

@@ -24,7 +24,6 @@
 #endif // no system header
 
 #include <cuda/experimental/__stf/utility/core.cuh>
-#include <cuda/experimental/__stf/utility/cuda_attributes.cuh>
 #include <cuda/experimental/__stf/utility/cuda_safe_call.cuh>
 
 #include <cassert>

cudax/include/cuda/experimental/__stf/internal/hashtable_linearprobing.cuh

@@ -27,7 +27,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/experimental/__stf/utility/cuda_attributes.cuh>
 #include <cuda/experimental/__stf/utility/cuda_safe_call.cuh>
 #include <cuda/experimental/__stf/utility/hash.cuh>
 

cudax/include/cuda/experimental/__stf/internal/parallel_for_scope.cuh

@@ -62,9 +62,9 @@ __global__ void loop(const _CCCL_GRID_CONSTANT size_t n, shape_t shape, F f, tup
 
   // This will explode the targs tuple into a pack of data
   // Help the compiler which may not detect that a device lambda is calling a device lambda
-  CUDASTF_NO_DEVICE_STACK
+  _CCCL_DIAG_SUPPRESS_NVHPC(no_device_stack)
   auto const explode_args = [&](auto&... data) {
-    CUDASTF_NO_DEVICE_STACK
+    _CCCL_DIAG_SUPPRESS_NVHPC(no_device_stack)
     auto const explode_coords = [&](auto&&... coords) {
       // No move/forward for `data` because it's used multiple times.
       f(::std::forward<decltype(coords)>(coords)..., data...);
@@ -304,9 +304,9 @@ __global__ void loop_redux(
   // This is used to build the arguments passed to the user-provided lambda function.
 
   // Help the compiler which may not detect that a device lambda is calling a device lambda
-  CUDASTF_NO_DEVICE_STACK
+  _CCCL_DIAG_SUPPRESS_NVHPC(no_device_stack)
   const auto explode_args = [&](auto&&... data) {
-    CUDASTF_NO_DEVICE_STACK
+    _CCCL_DIAG_SUPPRESS_NVHPC(no_device_stack)
     const auto explode_coords = [&](auto&&... coords) {
       // No move/forward for `data` because it's used multiple times.
       f(::std::forward<decltype(coords)>(coords)..., data...);

cudax/include/cuda/experimental/__stf/utility/cartesian_iterator.cuh

@@ -20,7 +20,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/experimental/__stf/utility/cuda_attributes.cuh>
 #include <cuda/experimental/__stf/utility/unittest.cuh>
 
 namespace cuda::experimental::stf::reserved

cudax/include/cuda/experimental/__stf/utility/core.cuh

@@ -24,8 +24,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/experimental/__stf/utility/cuda_attributes.cuh>
-
 #include <cstddef>
 #include <functional>
 #include <limits>

cudax/include/cuda/experimental/__stf/utility/cuda_attributes.cuh

@@ -25,7 +25,6 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/experimental/__stf/utility/cuda_attributes.cuh>
 #include <cuda/experimental/__stf/utility/hash.cuh>
 #include <cuda/experimental/__stf/utility/unittest.cuh>
 
@@ -491,7 +490,7 @@ public:
   _CCCL_HOST_DEVICE coords_t index_to_coords(size_t index) const
   {
     // Help the compiler which may not detect that a device lambda is calling a device lambda
-    CUDASTF_NO_DEVICE_STACK
+    _CCCL_DIAG_SUPPRESS_NVHPC(no_device_stack)
     return make_tuple_indexwise<dimensions>([&](auto i) {
       // included
       const ::std::ptrdiff_t begin_i  = get_begin(i);
@@ -500,7 +499,7 @@ public:
       index /= extent_i;
       return result;
     });
-    CUDASTF_NO_DEVICE_STACK
+    _CCCL_DIAG_SUPPRESS_NVHPC(no_device_stack)
   }
 
 private: