Merge pull request #2102 from glotzerlab/cuda-12.9

Support CUDA 12.9

Author: joaander

CHANGELOG.rst

@@ -19,6 +19,8 @@ Change Log
   (`#2097 <https://github.com/glotzerlab/hoomd-blue/pull/2097>`__).
 * Equations displayed in ``hoomd.hpmc.compute.SDF`` documentation
   (`#2096 <https://github.com/glotzerlab/hoomd-blue/discussions/2096>`__).
+* Support CUDA 12.9
+  (`#2102 <https://github.com/glotzerlab/hoomd-blue/discussions/2102>`__).
 
 5.3.0 (2025-06-26)
 ^^^^^^^^^^^^^^^^^^

hoomd/WarpTools.cuh

@@ -24,6 +24,7 @@
 #else
 #include <cub/warp/warp_reduce.cuh>
 #include <cub/warp/warp_scan.cuh>
+#include <cuda/std/type_traits>
 #endif
 
 #ifndef __CUDACC_RTC__
@@ -89,7 +90,7 @@ class WarpReduce
 #ifdef __HIP_PLATFORM_HCC__
         return Reduce(input, hipcub::Sum());
 #else
-        return Reduce(input, cub::Sum());
+        return Reduce(input, ::cuda::std::plus<> {});
 #endif
         }
 
@@ -108,7 +109,7 @@ class WarpReduce
 #ifdef __HIP_PLATFORM_HCC__
         return Reduce(input, hipcub::Sum(), valid_items);
 #else
-        return Reduce(input, cub::Sum(), valid_items);
+        return Reduce(input, ::cuda::std::plus<> {}, valid_items);
 #endif
         }
 
@@ -154,8 +155,7 @@ class WarpReduce
     typedef hipcub::WarpReduce<T, LOGICAL_WARP_THREADS, PTX_ARCH>
         MyWarpReduce; //!< CUB shuffle-based reduce
 #else
-    typedef cub::WarpReduce<T, LOGICAL_WARP_THREADS, PTX_ARCH>
-        MyWarpReduce; //!< CUB shuffle-based reduce
+    typedef cub::WarpReduce<T, LOGICAL_WARP_THREADS> MyWarpReduce; //!< CUB shuffle-based reduce
 #endif
     typedef typename MyWarpReduce::TempStorage
         TempStorage; //!< Nominal data type for CUB temporary storage
@@ -221,7 +221,7 @@ class WarpScan
 #ifdef __HIP_PLATFORM_HCC__
         InclusiveScan(input, output, hipcub::Sum());
 #else
-        InclusiveScan(input, output, cub::Sum());
+        InclusiveScan(input, output, ::cuda::std::plus<> {});
 #endif
         }
 
@@ -239,7 +239,7 @@ class WarpScan
 #ifdef __HIP_PLATFORM_HCC__
         InclusiveScan(input, output, hipcub::Sum(), aggregate);
 #else
-        InclusiveScan(input, output, cub::Sum(), aggregate);
+        InclusiveScan(input, output, ::cuda::std::plus<> {}, aggregate);
 #endif
         }
 
@@ -295,7 +295,7 @@ class WarpScan
 #ifdef __HIP_PLATFORM_HCC__
         ExclusiveScan(input, output, initial, hipcub::Sum());
 #else
-        ExclusiveScan(input, output, initial, cub::Sum());
+        ExclusiveScan(input, output, initial, ::cuda::std::plus<> {});
 #endif
         }
 
@@ -314,7 +314,7 @@ class WarpScan
 #ifdef __HIP_PLATFORM_HCC__
         ExclusiveScan(input, output, initial, hipcub::Sum(), aggregate);
 #else
-        ExclusiveScan(input, output, initial, cub::Sum(), aggregate);
+        ExclusiveScan(input, output, initial, ::cuda::std::plus<> {}, aggregate);
 #endif
         }
 
@@ -424,7 +424,7 @@ class WarpScan
     typedef hipcub::WarpScan<T, LOGICAL_WARP_THREADS, PTX_ARCH>
         MyWarpScan; //!< CUB shuffle-based scan
 #else
-    typedef cub::WarpScan<T, LOGICAL_WARP_THREADS, PTX_ARCH> MyWarpScan; //!< CUB shuffle-based scan
+    typedef cub::WarpScan<T, LOGICAL_WARP_THREADS> MyWarpScan; //!< CUB shuffle-based scan
 #endif
     typedef typename MyWarpScan::TempStorage
         TempStorage; //!< Nominal data type for CUB temporary storage

hoomd/extern/HIP

@@ -14,6 +14,7 @@
 #pragma GCC diagnostic ignored "-Wconversion"
 #include <cub/device/device_partition.cuh>
 #include <cub/iterator/counting_input_iterator.cuh>
+#include <thrust/iterator/counting_iterator.h>
 #pragma GCC diagnostic pop
 
 namespace hoomd
@@ -155,7 +156,7 @@ cudaError_t mpcd::gpu::partition_particles(void* d_tmp,
                                            unsigned int* d_num_remove,
                                            const unsigned int N)
     {
-    cub::CountingInputIterator<unsigned int> ids(0);
+    thrust::counting_iterator<unsigned int> ids(0);
     cub::DevicePartition::Flagged(d_tmp,
                                   tmp_bytes,
                                   ids,

hoomd/extern/hipCUB

@@ -7,6 +7,7 @@
  */
 
 #include <hipcub/hipcub.hpp>
+#include <thrust/iterator/counting_iterator.h>
 
 #include "RejectionVirtualParticleFillerGPU.cuh"
 
@@ -58,7 +59,7 @@ compact_virtual_particle_indices(void* d_tmp,
                                  unsigned int* d_keep_indices,
                                  unsigned int* d_num_keep)
     {
-    cub::CountingInputIterator<int> itr(0);
+    thrust::counting_iterator<int> itr(0);
     cub::DeviceSelect::Flagged(d_tmp,
                                tmp_bytes,
                                itr,