From 0bb528866ae2263a7f37aee23b63f275d4bcaf40 Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Mon, 11 May 2026 13:46:44 -0600 Subject: [PATCH] Remove warnings for thrust::less The use of `thrust::less` is deprecated. This is replaced with the more general `cuda::std::less`. The Viskores code is updated to point to the latter class for newer versions of Cuda. This commit was coauthored by Codex AI. --- docs/changelog/thrust-less-warnings.md | 7 + .../internal/DeviceAdapterAlgorithmCuda.h | 48 +++--- .../internal/DeviceAdapterAlgorithmKokkos.h | 36 ++++- viskores/exec/cuda/internal/ExecutionPolicy.h | 150 +++++++++++++----- 4 files changed, 182 insertions(+), 59 deletions(-) create mode 100644 docs/changelog/thrust-less-warnings.md diff --git a/docs/changelog/thrust-less-warnings.md b/docs/changelog/thrust-less-warnings.md new file mode 100644 index 0000000000..ff5f4d7840 --- /dev/null +++ b/docs/changelog/thrust-less-warnings.md @@ -0,0 +1,7 @@ +## Remove warnings for deprecated Thrust utilities + +The use of Thrust function objects such as `thrust::less`, `thrust::equal_to`, +and `thrust::plus`, as well as utilities such as `thrust::distance`, is +deprecated. These are replaced with the more general `cuda::std` equivalents. +The Viskores code is updated to point to the latter classes and functions for +newer versions of Cuda. diff --git a/viskores/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/viskores/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h index aa9f672e0a..9e06f8117d 100644 --- a/viskores/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h +++ b/viskores/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h @@ -457,7 +457,8 @@ struct DeviceAdapterAlgorithm cuda::internal::IteratorBegin(stencil), outputBegin, up); - return static_cast(::thrust::distance(outputBegin, newLast)); + return static_cast( + viskores::exec::cuda::internal::ThrustDistance(outputBegin, newLast)); } catch (...) { @@ -598,7 +599,8 @@ struct DeviceAdapterAlgorithm const OutputPortal& output) { using ValueType = typename ValuesPortal::ValueType; - LowerBoundsPortal(input, values, output, ::thrust::less()); + LowerBoundsPortal( + input, values, output, viskores::exec::cuda::internal::ThrustLess()); } template @@ -606,7 +608,8 @@ struct DeviceAdapterAlgorithm const OutputPortal& values_output) { using ValueType = typename InputPortal::ValueType; - LowerBoundsPortal(input, values_output, values_output, ::thrust::less()); + LowerBoundsPortal( + input, values_output, values_output, viskores::exec::cuda::internal::ThrustLess()); } template @@ -638,7 +641,7 @@ struct DeviceAdapterAlgorithm template VISKORES_CONT static T ReducePortal(const InputPortal& input, T initialValue) { - return ReducePortal(input, initialValue, ::thrust::plus()); + return ReducePortal(input, initialValue, viskores::exec::cuda::internal::ThrustPlus()); } template @@ -722,7 +725,7 @@ struct DeviceAdapterAlgorithm ::thrust::pair result_iterators; - ::thrust::equal_to binaryPredicate; + viskores::exec::cuda::internal::ThrustEqualTo binaryPredicate; using ValueType = typename ValuesPortal::ValueType; viskores::exec::cuda::internal::WrappedBinaryOperator bop( @@ -744,7 +747,8 @@ struct DeviceAdapterAlgorithm cuda::internal::throwAsViskoresException(); } - return static_cast(::thrust::distance(keys_out_begin, result_iterators.first)); + return static_cast( + viskores::exec::cuda::internal::ThrustDistance(keys_out_begin, result_iterators.first)); } template @@ -756,7 +760,7 @@ struct DeviceAdapterAlgorithm return ScanExclusivePortal(input, output, - (::thrust::plus()), + (viskores::exec::cuda::internal::ThrustPlus()), viskores::TypeTraits::ZeroInitialization()); } @@ -814,7 +818,8 @@ struct DeviceAdapterAlgorithm const OutputPortal& output) { using ValueType = typename OutputPortal::ValueType; - return ScanInclusivePortal(input, output, ::thrust::plus()); + return ScanInclusivePortal( + input, output, viskores::exec::cuda::internal::ThrustPlus()); } template @@ -855,8 +860,11 @@ struct DeviceAdapterAlgorithm { using KeyType = typename KeysPortal::ValueType; using ValueType = typename OutputPortal::ValueType; - ScanInclusiveByKeyPortal( - keys, values, output, ::thrust::equal_to(), ::thrust::plus()); + ScanInclusiveByKeyPortal(keys, + values, + output, + viskores::exec::cuda::internal::ThrustEqualTo(), + viskores::exec::cuda::internal::ThrustPlus()); } template values, output, viskores::TypeTraits::ZeroInitialization(), - ::thrust::equal_to(), - ::thrust::plus()); + viskores::exec::cuda::internal::ThrustEqualTo(), + viskores::exec::cuda::internal::ThrustPlus()); } template VISKORES_CONT static void SortPortal(const ValuesPortal& values) { using ValueType = typename ValuesPortal::ValueType; - SortPortal(values, ::thrust::less()); + SortPortal(values, viskores::exec::cuda::internal::ThrustLess()); } template @@ -974,7 +982,7 @@ struct DeviceAdapterAlgorithm VISKORES_CONT static void SortByKeyPortal(const KeysPortal& keys, const ValuesPortal& values) { using ValueType = typename KeysPortal::ValueType; - SortByKeyPortal(keys, values, ::thrust::less()); + SortByKeyPortal(keys, values, viskores::exec::cuda::internal::ThrustLess()); } template @@ -1007,7 +1015,8 @@ struct DeviceAdapterAlgorithm auto begin = cuda::internal::IteratorBegin(values); auto newLast = ::thrust::unique(ThrustCudaPolicyPerThread, begin, cuda::internal::IteratorEnd(values)); - return static_cast(::thrust::distance(begin, newLast)); + return static_cast( + viskores::exec::cuda::internal::ThrustDistance(begin, newLast)); } catch (...) { @@ -1028,7 +1037,8 @@ struct DeviceAdapterAlgorithm auto begin = cuda::internal::IteratorBegin(values); auto newLast = ::thrust::unique( ThrustCudaPolicyPerThread, begin, cuda::internal::IteratorEnd(values), bop); - return static_cast(::thrust::distance(begin, newLast)); + return static_cast( + viskores::exec::cuda::internal::ThrustDistance(begin, newLast)); } catch (...) { @@ -1586,7 +1596,7 @@ struct DeviceAdapterAlgorithm ScanInclusiveByKeyPortal(keysPortal, valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), - ::thrust::equal_to(), + viskores::exec::cuda::internal::ThrustEqualTo(), binary_functor); } @@ -1615,7 +1625,7 @@ struct DeviceAdapterAlgorithm valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), viskores::TypeTraits::ZeroInitialization(), - ::thrust::equal_to(), + viskores::exec::cuda::internal::ThrustEqualTo(), viskores::Add()); } @@ -1651,7 +1661,7 @@ struct DeviceAdapterAlgorithm valuesPortal, output.PrepareForOutput(numberOfValues, DeviceAdapterTagCuda(), token), initialValue, - ::thrust::equal_to(), + viskores::exec::cuda::internal::ThrustEqualTo(), binary_functor); } diff --git a/viskores/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h b/viskores/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h index 47d742a58e..6994e4b18d 100644 --- a/viskores/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h +++ b/viskores/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h @@ -53,8 +53,13 @@ VISKORES_THIRDPARTY_POST_INCLUDE #if defined(VISKORES_USE_KOKKOS_THRUST) #include +#include #include #include +#include +#if defined(__CUDA__) && THRUST_VERSION >= 300100 +#include +#endif #endif namespace viskores @@ -81,6 +86,27 @@ namespace kokkos namespace internal { +// In Thrust 3.1, several std-like utility functions and objects in the +// `thrust` namespace were deprecated in favor of equivalent objects +// provided by Cuda in the `cuda:std` namespace. +#if defined(VISKORES_USE_KOKKOS_THRUST) +#if defined(__CUDA__) && THRUST_VERSION >= 300100 +template +using ThrustSortLess = ::cuda::std::less; +template +using ThrustSortGreater = ::cuda::std::greater; +template +using ThrustEqualTo = ::cuda::std::equal_to; +#else +template +using ThrustSortLess = ::thrust::less; +template +using ThrustSortGreater = ::thrust::greater; +template +using ThrustEqualTo = ::thrust::equal_to; +#endif +#endif + //---------------------------------------------------------------------------- template struct BitFieldToBoolField : public viskores::exec::FunctorBase @@ -866,11 +892,13 @@ struct DeviceAdapterAlgorithm if (std::is_same::value) { - thrust::sort_by_key(keys_begin, keys_end, values_begin, thrust::less()); + thrust::sort_by_key( + keys_begin, keys_end, values_begin, kokkos::internal::ThrustSortLess()); } else { - thrust::sort_by_key(keys_begin, keys_end, values_begin, thrust::greater()); + thrust::sort_by_key( + keys_begin, keys_end, values_begin, kokkos::internal::ThrustSortGreater()); } } @@ -957,7 +985,7 @@ struct DeviceAdapterAlgorithm values_begin, keys_output_begin, values_output_begin, - thrust::equal_to(), + kokkos::internal::ThrustEqualTo(), binary_functor); num_unique_keys = ends.first - keys_output_begin; @@ -1004,7 +1032,7 @@ struct DeviceAdapterAlgorithm values_begin, keys_output_begin, values_output_begin, - thrust::equal_to(), + kokkos::internal::ThrustEqualTo(), binary_functor); num_unique_keys = ends.first - keys_output_begin; diff --git a/viskores/exec/cuda/internal/ExecutionPolicy.h b/viskores/exec/cuda/internal/ExecutionPolicy.h index 183ed0a4c0..9eafcea14d 100644 --- a/viskores/exec/cuda/internal/ExecutionPolicy.h +++ b/viskores/exec/cuda/internal/ExecutionPolicy.h @@ -24,23 +24,81 @@ #include VISKORES_THIRDPARTY_PRE_INCLUDE +#include #include +#include #include #include #include #include +#include +#if THRUST_VERSION >= 300100 +#include +#include +#endif VISKORES_THIRDPARTY_POST_INCLUDE #define ThrustCudaPolicyPerThread ::thrust::cuda::par.on(cudaStreamPerThread) +namespace viskores +{ +namespace exec +{ +namespace cuda +{ +namespace internal +{ + +// In Thrust 3.1, several std-like utility functions and objects in the +// `thrust` namespace were deprecated in favor of equivalent objects +// provided by Cuda in the `cuda:std` namespace. +#if THRUST_VERSION >= 300100 +template +using ThrustLess = ::cuda::std::less; +template +using ThrustGreater = ::cuda::std::greater; +template +using ThrustEqualTo = ::cuda::std::equal_to; +template +using ThrustPlus = ::cuda::std::plus; + +template +VISKORES_EXEC_CONT auto ThrustDistance(Iterator first, Iterator last) + -> decltype(::cuda::std::distance(first, last)) +{ + return ::cuda::std::distance(first, last); +} +#else +template +using ThrustLess = ::thrust::less; +template +using ThrustGreater = ::thrust::greater; +template +using ThrustEqualTo = ::thrust::equal_to; +template +using ThrustPlus = ::thrust::plus; + +template +VISKORES_EXEC_CONT auto ThrustDistance(Iterator first, Iterator last) + -> decltype(::thrust::distance(first, last)) +{ + return ::thrust::distance(first, last); +} +#endif + +} +} +} +} // namespace viskores::exec::cuda::internal + struct viskores_cuda_policy : thrust::device_execution_policy { }; //Specialize the sort call for cuda pointers using less/greater operators. //The purpose of this is that for 32bit types (UInt32,Int32,Float32) thrust -//will call a super fast radix sort only if the operator is thrust::less -//or thrust::greater. +//will call a super fast radix sort only if the operator is the default less +//or greater comparator. VISKORES_SUPPRESS_EXEC_WARNINGS template __host__ __device__ void sort( @@ -50,7 +108,8 @@ __host__ __device__ void sort( viskores::exec::cuda::internal::WrappedBinaryPredicate comp) { //sort for concrete pointers and less than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort(ThrustCudaPolicyPerThread, first, last, thrust::less()); + return thrust::sort( + ThrustCudaPolicyPerThread, first, last, viskores::exec::cuda::internal::ThrustLess()); } VISKORES_SUPPRESS_EXEC_WARNINGS @@ -63,34 +122,43 @@ __host__ __device__ void sort_by_key( viskores::exec::cuda::internal::WrappedBinaryPredicate comp) { //sort for concrete pointers and less than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort_by_key( - ThrustCudaPolicyPerThread, first, last, values_first, thrust::less()); + return thrust::sort_by_key(ThrustCudaPolicyPerThread, + first, + last, + values_first, + viskores::exec::cuda::internal::ThrustLess()); } VISKORES_SUPPRESS_EXEC_WARNINGS template -__host__ __device__ void sort( - const viskores_cuda_policy& exec, - T* first, - T* last, - viskores::exec::cuda::internal::WrappedBinaryPredicate> comp) +__host__ __device__ void sort(const viskores_cuda_policy& exec, + T* first, + T* last, + viskores::exec::cuda::internal::WrappedBinaryPredicate< + T, + viskores::exec::cuda::internal::ThrustLess> comp) { //sort for concrete pointers and less than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort(ThrustCudaPolicyPerThread, first, last, thrust::less()); + return thrust::sort( + ThrustCudaPolicyPerThread, first, last, viskores::exec::cuda::internal::ThrustLess()); } VISKORES_SUPPRESS_EXEC_WARNINGS template -__host__ __device__ void sort_by_key( - const viskores_cuda_policy& exec, - T* first, - T* last, - RandomAccessIterator values_first, - viskores::exec::cuda::internal::WrappedBinaryPredicate> comp) +__host__ __device__ void sort_by_key(const viskores_cuda_policy& exec, + T* first, + T* last, + RandomAccessIterator values_first, + viskores::exec::cuda::internal::WrappedBinaryPredicate< + T, + viskores::exec::cuda::internal::ThrustLess> comp) { //sort for concrete pointers and less than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort_by_key( - ThrustCudaPolicyPerThread, first, last, values_first, thrust::less()); + return thrust::sort_by_key(ThrustCudaPolicyPerThread, + first, + last, + values_first, + viskores::exec::cuda::internal::ThrustLess()); } VISKORES_SUPPRESS_EXEC_WARNINGS @@ -102,7 +170,8 @@ __host__ __device__ void sort( viskores::exec::cuda::internal::WrappedBinaryPredicate comp) { //sort for concrete pointers and greater than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort(ThrustCudaPolicyPerThread, first, last, thrust::greater()); + return thrust::sort( + ThrustCudaPolicyPerThread, first, last, viskores::exec::cuda::internal::ThrustGreater()); } VISKORES_SUPPRESS_EXEC_WARNINGS @@ -115,34 +184,43 @@ __host__ __device__ void sort_by_key( viskores::exec::cuda::internal::WrappedBinaryPredicate comp) { //sort for concrete pointers and greater than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort_by_key( - ThrustCudaPolicyPerThread, first, last, values_first, thrust::greater()); + return thrust::sort_by_key(ThrustCudaPolicyPerThread, + first, + last, + values_first, + viskores::exec::cuda::internal::ThrustGreater()); } VISKORES_SUPPRESS_EXEC_WARNINGS template -__host__ __device__ void sort( - const viskores_cuda_policy& exec, - T* first, - T* last, - viskores::exec::cuda::internal::WrappedBinaryPredicate> comp) +__host__ __device__ void sort(const viskores_cuda_policy& exec, + T* first, + T* last, + viskores::exec::cuda::internal::WrappedBinaryPredicate< + T, + viskores::exec::cuda::internal::ThrustGreater> comp) { //sort for concrete pointers and greater than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort(ThrustCudaPolicyPerThread, first, last, thrust::greater()); + return thrust::sort( + ThrustCudaPolicyPerThread, first, last, viskores::exec::cuda::internal::ThrustGreater()); } VISKORES_SUPPRESS_EXEC_WARNINGS template -__host__ __device__ void sort_by_key( - const viskores_cuda_policy& exec, - T* first, - T* last, - RandomAccessIterator values_first, - viskores::exec::cuda::internal::WrappedBinaryPredicate> comp) +__host__ __device__ void sort_by_key(const viskores_cuda_policy& exec, + T* first, + T* last, + RandomAccessIterator values_first, + viskores::exec::cuda::internal::WrappedBinaryPredicate< + T, + viskores::exec::cuda::internal::ThrustGreater> comp) { //sort for concrete pointers and greater than op //this makes sure that we invoke the thrust radix sort and not merge sort - return thrust::sort_by_key( - ThrustCudaPolicyPerThread, first, last, values_first, thrust::greater()); + return thrust::sort_by_key(ThrustCudaPolicyPerThread, + first, + last, + values_first, + viskores::exec::cuda::internal::ThrustGreater()); } VISKORES_SUPPRESS_EXEC_WARNINGS