Use external hip and hipcub headers when HOOMD_GPU_PLATFORM=CUDA

.bumpversion.toml

@@ -6,6 +6,11 @@ filename = "INSTALLING.rst"
 search = "hoomd={current_version}"
 replace = "hoomd={new_version}"
 
+[[tool.bumpversion.files]]
+filename = "BUILDING.rst"
+search = "tree/v{current_version}"
+replace = "tree/v{new_version}"
+
 [[tool.bumpversion.files]]
 filename = "CMakeLists.txt"
 search = 'HOOMD_VERSION_RAW "{current_version}"'

.github/workflows/build_and_test.yaml

@@ -87,12 +87,18 @@ jobs:
           echo 'test_docker_options=--gpus=all' >> "$GITHUB_OUTPUT"
 
           case "${{ inputs.compiler_version }}" in
+            130)
+              echo "image=nvidia/cuda:13.0.2-devel-ubuntu24.04" >> "$GITHUB_OUTPUT";;
+            129)
+              echo "image=nvidia/cuda:12.9.1-devel-ubuntu24.04" >> "$GITHUB_OUTPUT";;
+            128)
+              echo "image=nvidia/cuda:12.8.1-devel-ubuntu24.04" >> "$GITHUB_OUTPUT";;
+            126)
+              echo "image=nvidia/cuda:12.6.3-devel-ubuntu24.04" >> "$GITHUB_OUTPUT";;
             125)
-              echo "image=nvidia/cuda:12.5.0-devel-ubuntu22.04" >> "$GITHUB_OUTPUT";;
+              echo "image=nvidia/cuda:12.5.1-devel-ubuntu24.04" >> "$GITHUB_OUTPUT";;
             124)
               echo "image=nvidia/cuda:12.4.1-devel-ubuntu22.04" >> "$GITHUB_OUTPUT";;
-            122)
-              echo "image=nvidia/cuda:12.2.2-devel-ubuntu22.04" >> "$GITHUB_OUTPUT";;
             *)
               echo "Unknown compiler" && exit 1;;
             esac
@@ -134,6 +140,56 @@ jobs:
         activate-environment: true
         manifest-path: code/pixi.toml
 
+    - name: Clone hip
+      if: ${{ inputs.compiler_family == 'cuda' }}
+      uses: actions/checkout@93cb6efe18208431cddfb8368fd83d5badbf9bfd # v5.0.1
+      with:
+        repository: ROCm/rocm-systems
+        ref: 75ad45d5f131738a4d888499026edfeaa3a11fe3 # hip-version_7.2.53220
+        path: rocm-systems
+
+    - name: Clone hipCUB
+      if: ${{ inputs.compiler_family == 'cuda' }}
+      uses: actions/checkout@93cb6efe18208431cddfb8368fd83d5badbf9bfd # v5.0.1
+      with:
+        repository: ROCm/rocm-libraries
+        ref: 1a01b92fd0971c98258c0ea7811aea2e55ec8698 # rocm-7.1.0
+        sparse-checkout: projects/hipcub
+        path: rocm-libraries
+
+    - name: Patch hip
+      if: ${{ inputs.compiler_family == 'cuda' }}
+      run: |
+        git apply ${GITHUB_WORKSPACE}/code/.github/workflows/fix-cuda-12-rocm-systems.patch
+      working-directory: rocm-systems
+
+    - name: Patch hipcub
+      if: ${{ inputs.compiler_family == 'cuda' }}
+      run: |
+        git apply ${GITHUB_WORKSPACE}/code/.github/workflows/fix-cuda-12-rocm-libraries.patch
+      working-directory: rocm-libraries
+
+    - name: Install hip
+      if: ${{ inputs.compiler_family == 'cuda' }}
+      run: |
+        export CLR_DIR="$(readlink -f rocm-systems/projects/clr)"
+        export HIP_DIR="$(readlink -f rocm-systems/projects/hip)"
+        export HIP_OTHER="$(readlink -f rocm-systems/projects/hipother)"
+        cd "$CLR_DIR"
+        mkdir build
+        cd build
+        cmake -DHIP_COMMON_DIR=$HIP_DIR -DHIP_PLATFORM=nvidia -DCMAKE_INSTALL_PREFIX=${GITHUB_WORKSPACE}/local -DCLR_BUILD_HIP=ON -DCLR_BUILD_OCL=OFF -DHIPNV_DIR=$HIP_OTHER/hipnv ..
+        make install
+
+    - name: Install hipcub
+      if: ${{ inputs.compiler_family == 'cuda' }}
+      run: |
+        mkdir build
+        cd build
+        cmake ../ -Dhip_ROOT=${GITHUB_WORKSPACE}/local -DHIP_PLATFORM=nvidia -DCMAKE_INSTALL_PREFIX=${GITHUB_WORKSPACE}/local
+        make install
+      working-directory: rocm-libraries/projects/hipcub
+
     - name: Set compiler
       if: ${{ inputs.compiler_family != 'cuda' }}
       run: |
@@ -157,8 +213,9 @@ jobs:
                                -DBUILD_MPCD=${BUILD_MD:-"ON"} \
                                -DBUILD_METAL=${BUILD_MD:-"ON"} \
                                -DBUILD_HPMC=${BUILD_HPMC:-"ON"} \
-                               -DCUDA_ARCH_LIST="60;70" \
+                               -DCUDA_ARCH_LIST="80" \
                                -DCMAKE_INSTALL_PREFIX=${GITHUB_WORKSPACE}/install \
+                               -Dhip_ROOT=${GITHUB_WORKSPACE}/local \
                                -DPLUGINS=""
       env:
         ENABLE_GPU: ${{ contains(inputs.config, 'cuda') }}

.github/workflows/fix-cuda-12-rocm-libraries.patch

@@ -0,0 +1,129 @@
+diff --git a/projects/hipcub/CMakeLists.txt b/projects/hipcub/CMakeLists.txt
+index fbcecf0fc4..0fbae70749 100644
+--- a/projects/hipcub/CMakeLists.txt
++++ b/projects/hipcub/CMakeLists.txt
+@@ -174,7 +174,7 @@ math(EXPR hipcub_VERSION_NUMBER "${hipcub_VERSION_MAJOR} * 100000 + ${hipcub_VER
+ include(VerifyCompiler)
+
+ # Get dependencies (except rocm-cmake, included earlier)
+-include(Dependencies)
++# include(Dependencies)
+
+ if(BUILD_ADDRESS_SANITIZER)
+   add_compile_options(-fsanitize=address -shared-libasan)
+diff --git a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_for.hpp b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_for.hpp
+index 0f22c40517..31fab88dd6 100644
+--- a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_for.hpp
++++ b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_for.hpp
+@@ -30,6 +30,7 @@
+ #define HIPCUB_CUB_DEVICE_DEVICE_FOR_HPP_
+
+ #include "../../../config.hpp"
++#if CUDA_VERSION >= 12060
+
+ #include <cub/device/device_for.cuh> // IWYU pragma: export
+
+@@ -194,4 +195,5 @@ HIPCUB_RUNTIME_FUNCTION
+
+ END_HIPCUB_NAMESPACE
+
++#endif
+ #endif // HIPCUB_CUB_DEVICE_DEVICE_FOR_HPP_
+diff --git a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_merge.hpp b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_merge.hpp
+index f314f5a128..e22f08fba9 100644
+--- a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_merge.hpp
++++ b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_merge.hpp
+@@ -30,6 +30,7 @@
+ #define HIPCUB_CUB_DEVICE_DEVICE_MERGE_HPP_
+
+ #include "../../../config.hpp"
++#if CUDA_VERSION >= 12080
+
+ #include <cub/device/device_merge.cuh> // IWYU pragma: export
+
+@@ -104,4 +105,5 @@ struct DeviceMerge
+
+ END_HIPCUB_NAMESPACE
+
++#endif
+ #endif // HIPCUB_CUB_DEVICE_DEVICE_MERGE_HPP_
+diff --git a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_scan.hpp b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_scan.hpp
+index 72ad11f7bc..038fddccdd 100644
+--- a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_scan.hpp
++++ b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_scan.hpp
+@@ -182,6 +182,7 @@ public:
+                                                  stream);
+     }
+
++#if CUDA_VERSION >= 12080
+     template<typename InputIteratorT,
+              typename OutputIteratorT,
+              typename ScanOpT,
+@@ -206,6 +207,7 @@ public:
+                                                                            num_items,
+                                                                            stream));
+     }
++#endif
+
+     template<typename InputIteratorT, typename OutputIteratorT, typename NumItemsT>
+     HIPCUB_RUNTIME_FUNCTION
+diff --git a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_select.hpp b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_select.hpp
+index 6812c5cfeb..33bce4aa23 100644
+--- a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_select.hpp
++++ b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_select.hpp
+@@ -223,6 +223,7 @@ public:
+                   stream);
+     }
+
++#if CUDA_VERSION >= 12060
+     template<typename InputIteratorT,
+              typename FlagIterator,
+              typename OutputIteratorT,
+@@ -329,6 +330,7 @@ public:
+                          select_op,
+                          stream);
+     }
++#endif
+
+     template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT>
+     HIPCUB_RUNTIME_FUNCTION
+diff --git a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_transform.hpp b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_transform.hpp
+index 681a0bbf98..d071c7d859 100644
+--- a/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_transform.hpp
++++ b/projects/hipcub/hipcub/include/hipcub/backend/cub/device/device_transform.hpp
+@@ -30,6 +30,7 @@
+ #define HIBCUB_ROCPRIM_DEVICE_DEVICE_TRANSFORM_HPP_
+
+ #include "../../../config.hpp"
++#if CUDA_VERSION >= 12090
+
+ #include <hip/hip_runtime.h>
+
+@@ -203,4 +204,5 @@ struct DeviceTransform
+
+ END_HIPCUB_NAMESPACE
+
++#endif
+ #endif // HIBCUB_ROCPRIM_DEVICE_DEVICE_TRANSFORM_HPP_
+diff --git a/projects/hipcub/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp b/projects/hipcub/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp
+index fc67d645b1..63540b717f 100644
+--- a/projects/hipcub/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp
++++ b/projects/hipcub/hipcub/include/hipcub/backend/cub/util_temporary_storage.hpp
+@@ -50,10 +50,17 @@ HIPCUB_HOST_DEVICE HIPCUB_FORCEINLINE hipError_t
+                      void* (&allocations)[ALLOCATIONS],
+                      const size_t (&allocation_sizes)[ALLOCATIONS])
+ {
++    #if CUDA_VERSION >= 12090
+     cudaError_t error = ::cub::detail::AliasTemporaries(d_temp_storage,
+                                                         temp_storage_bytes,
+                                                         allocations,
+                                                         allocation_sizes);
++    #else
++    cudaError_t error = ::cub::AliasTemporaries(d_temp_storage,
++                                                        temp_storage_bytes,
++                                                        allocations,
++                                                        allocation_sizes);
++    #endif
+
+     if(cudaSuccess == error)
+     {

.github/workflows/fix-cuda-12-rocm-systems.patch

@@ -0,0 +1,41 @@
+diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+index 8f6c295aab..05036a57fe 100644
+--- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
++++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+@@ -869,7 +869,7 @@ static inline void hipMemcpy2DTocudaMemcpy2D(CUDA_MEMCPY2D* a, const hip_Memcpy2
+   a->Height = (size_t)p->Height;
+ }
+
+-#if CUDA_VERSION >= CUDA_12020
++#if CUDA_VERSION >= 12080
+ typedef enum cudaMemcpyFlags hipMemcpyFlags;
+ #define hipMemcpyFlagDefault cudaMemcpyFlagDefault
+ #define hipMemcpyFlagPreferOverlapWithCompute cudaMemcpyFlagPreferOverlapWithCompute
+@@ -2430,6 +2430,7 @@ inline static hipError_t hipMemcpy2DToArrayAsync(hipArray_t dst, size_t wOffset,
+       cudaMemcpy2DToArrayAsync(dst, wOffset, hOffset, src, spitch, width, height, kind, stream));
+ }
+
++#if CUDA_VERSION >= 12080
+ inline static hipError_t hipMemcpyBatchAsync(void** dsts, void** srcs, size_t* sizes, size_t count,
+                                              hipMemcpyAttributes* attrs, size_t* attrsIdxs,
+                                              size_t numAttrs, size_t* failIdx, hipStream_t stream) {
+@@ -2467,6 +2468,7 @@ inline static hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms* p) {
+ inline static hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms* p, hipStream_t stream) {
+   return hipCUDAErrorTohipError(cudaMemcpy3DPeerAsync(p, stream));
+ }
++#endif
+
+ __HIP_DEPRECATED inline static hipError_t hipMemcpyToArray(hipArray_t dst, size_t wOffset,
+                                                            size_t hOffset, const void* src,
+@@ -3756,9 +3758,11 @@ inline static hipError_t hipLibraryEnumerateKernels(hipKernel_t* kernels, unsign
+   return hipCUResultTohipError(cuLibraryEnumerateKernels(kernels, numKernels, library));
+ }
+
++#if CUDA_VERSION >= 12060
+ inline static hipError_t hipKernelGetLibrary(hipLibrary_t* library, hipKernel_t kernel) {
+   return hipCUResultTohipError(cuKernelGetLibrary(library, kernel));
+ }
++#endif
+
+ inline static hipError_t hipKernelGetName(const char** name, hipKernel_t kernel) {
+   return hipCUResultTohipError(cuKernelGetName(name, kernel));

.github/workflows/test.yaml

@@ -53,13 +53,13 @@ jobs:
         - config: [gcc, 14, -py, 313, -nomd, -nohpmc]
         - config: [gcc, 10, -py, 310, -mpi]
 
-        - config: [cuda, 124, -py, 313, -mpi]
+        - config: [cuda, 125, -py, 313, -mpi]
           validate: true
 
-        - config: [cuda, 124, -py, 313]
+        - config: [cuda, 125, -py, 313]
           validate: true
 
-        - config: [cuda, 124, -py, 313, -mpi, -debug]
+        - config: [cuda, 125, -py, 313, -mpi, -debug]
 
 
   release:
@@ -95,6 +95,9 @@ jobs:
         - config: [gcc, 13, -py, 312, -mpi]
         - config: [gcc, 12, -py, 311, -mpi]
         - config: [gcc, 11, -py, 310, -mpi]
+        - config: [cuda, 126, -py, 313, -mpi]
+        - config: [cuda, 128, -py, 313, -mpi]
+        - config: [cuda, 129, -py, 313, -mpi]
 
   tests_complete:
     name: Unit test

.gitmodules

@@ -7,13 +7,6 @@
 [submodule "hoomd/extern/quickhull"]
 	path = hoomd/extern/quickhull
 	url = https://github.com/glotzerlab/quickhull
-[submodule "hoomd/extern/HIP"]
-	path = hoomd/extern/HIP
-	url = https://github.com/glotzerlab/HIP.git
-[submodule "hoomd/extern/hipCUB"]
-	path = hoomd/extern/hipCUB
-	url = https://github.com/glotzerlab/hipCUB
-	branch = header_only
 [submodule "hoomd/extern/neighbor"]
 	path = hoomd/extern/neighbor
 	url = https://github.com/mphowardlab/neighbor.git

.pre-commit-config.yaml

@@ -36,7 +36,7 @@ repos:
   - id: end-of-file-fixer
     exclude_types: [svg]
   - id: trailing-whitespace
-    exclude_types: [svg]
+    exclude_types: [svg, diff]
   - id: check-json
   - id: check-yaml
     exclude: "\\.clang-format"

BUILDING.rst

@@ -124,29 +124,38 @@ Install additional packages needed to build the documentation:
 
 **For GPU execution** (required when ``ENABLE_GPU=on``):
 
-- **NVIDIA CUDA Toolkit**
+.. tab:: NVIDIA (CUDA)
 
-  *OR*
+    - NVIDIA CUDA Toolkit
+    - hip (`hip installation instructions`_)
+    - hipcub (`hipcub installation instructions`_)
 
-- AMD ROCm
-- HIP [with ``hipcc`` and ``hcc`` as backend]
-- rocFFT
-- rocPRIM
-- rocThrust
-- hipCUB
-- roctracer-dev
+    .. note::
 
-.. note::
+        hip ``hip-version_7.2.53220`` and hipcub ``rocm-7.1.0`` work with CUDA 12.9.
+        Apply the patches in `.github/workflows`_ to add support for CUDA 12.5–12.8.
+
+.. tab:: AMD (HIP)
+
+    - AMD ROCm
+    - HIP [with ``hipcc`` and ``hcc`` as backend]
+    - rocFFT
+    - rocPRIM
+    - rocThrust
+    - hipCUB
+    - roctracer-dev
 
-    When ``ENABLE_GPU=on``, HOOMD-blue will default to CUDA. Set ``HOOMD_GPU_PLATFORM=HIP`` to
-    choose HIP.
+.. _`hip installation instructions`: https://rocmdocs.amd.com/projects/HIP/en/latest/install/build.html
+.. _`hipcub installation instructions`: https://rocm.docs.amd.com/projects/hipCUB/en/latest/install/hipCUB-install-overview.html
+.. _`.github/workflows`: https://github.com/glotzerlab/hoomd-blue/tree/v5.4.0/.github/workflows
 
 **To build the documentation:**
 
 - **sphinx**
 - **furo**
 - **nbsphinx**
 - **ipython**
+- **sphinx-inline-tabs**
 
 .. _Obtain the source: