[ROCm] Port the core GPU tensor stack to HIP (USE_HIP)

[jeffdaily]

This adds an AMD GPU (ROCm/HIP) build for Open3D's core GPU "tensor" subsystem (cpp/open3d/core, cpp/open3d/t), so the GPU code paths run on AMD in addition to NVIDIA. It is enabled with -DUSE_HIP=ON and validated on Linux AMD GPUs (gfx90a CDNA2 and gfx1100 RDNA3) and on Windows (gfx1201 RDNA4) under the all-clang ROCm toolchain. Building with ROCm is documented in docs/compilation.rst.

The approach keeps the existing CUDA sources in their CUDA spelling: only the .cu translation units take the HIP toolchain, the host C++ is untouched, and the NVIDIA build is byte-for-byte unchanged. USE_HIP=ON reuses the existing BUILD_CUDA_MODULE switch (which already gates the GPU code throughout the tree) and only swaps the language to HIP, the math libraries to hipBLAS/hipSOLVER/hipSPARSE, and retags the .cu sources LANGUAGE HIP. A single compat header (cpp/open3d/core/hip/CUDAToHIP.h), force-included on the HIP translation units, aliases the CUDA-spelled runtime symbols to their HIP equivalents; forwarding shims under hip_compat/ redirect the <cuda.h>/<cuda_runtime.h>/<cub/cub.cuh> includes so no include line is edited.

To review, start with CMakeLists.txt (the USE_HIP enable) and cmake/Open3DSetGlobalProperties.cmake (the .cu LANGUAGE HIP retag, force-include, and THRUST_DEVICE_SYSTEM=HIP), then the compat header, then find_dependencies.cmake (the HIP math-lib targets and the stdgpu HIP backend), then the device/host-guard source fixes, and finally the two test files (cpp/tests/core/HashMap.cpp, cpp/tests/t/geometry/VoxelBlockGrid.cpp) which gate the non-default Slab hashmap backend out under USE_HIP.

Key fixes for the AMD architecture:

• FAISS warp-select (core/nns/kernel): kept kWarpSize=32 and run each 64-lane wavefront as two independent 32-lane bitonic groups (the smallest instantiated warp queue is 32, so kWarpSize=64 would collapse it); replaced the inline-PTX PtxUtils with clang builtins; widened the _shfl*_sync / __any_sync masks to 64-bit. KNN / FixedRadius / Hybrid GPU results match the CPU reference.
• Extended the CUDACC/CUDA_ARCH device-vs-host guards across the kernel headers to also accept HIPCC/HIP_DEVICE_COMPILE, rather than faking CUDACC globally (which would force rocThrust to its CUDA backend); same for MiniVec's FN_SPECIFIERS and the explicit-instantiation attributes in RegistrationImpl.h. CUDACC is deliberately not defined, and THRUST_DEVICE_SYSTEM is pinned to the HIP backend rather than relying on auto-detect.
• Library swaps in LinalgHeadersCUDA.h (cublas/cusolver -> hipblas/hipsolver typed Dn API), guarding the cuSOLVER-only status enum cases.
• stdgpu's HIP backend (the default hashmap) is built via ExternalProject with hipcc as CXX plus a missing-install-file patch. The StdGPU hashmap is wave64-correct as-is: its concurrent duplicate-key dedup elects exactly one winner per key across all 64 lanes (confirmed by repeating the dedup-heavy HashMap tests), so no stdgpu source change is needed.

The hashmap unit tests loop over both GPU backends (Slab and StdGPU). The Slab backend (warp-cooperative, a tid>>5 / __ffs-ballot 32-lane election) is wave64-incorrect on a 64-lane architecture and is out of scope here, so HashMap.cpp and the VoxelBlockGrid test's EnumerateBackends helper exercise only the StdGPU backend under USE_HIP. With Slab excluded, the full HashMap suite and the VoxelBlockGrid TSDF suite (which uses the StdGPU hashmap internally) pass on GPU.

Out of scope for this change (guarded, not regressed): the NPP-based GPU image filters (no ROCm equivalent; the CPU/IPP path is unaffected and the GPU ops report unsupported), the non-default SlabHash backend's wave64 correctness (made to compile only; its tests are excluded under USE_HIP as above), and the CUTLASS-based ML conv ops (ML modules off).

On Windows, USE_HIP forces an all-clang toolchain (enable_language(HIP) does not allow a clang-HIP plus MSVC-host mix), so CMake's MSVC variable is false. A few places keyed Windows behavior on if(MSVC) or assumed an MSVC environment, which kept the ROCm port from building under clang on Windows; these are now gated on WIN32 or the Clang compiler id, leaving the Linux and NVIDIA builds byte-for-byte unchanged: the bundled zlib/libpng imported library names follow WIN32 (they build as zlibstatic / libpng16_static regardless of the frontend), the ignored-nodiscard hipError_t that clang reports under -Wunused-value is demoted to a warning (gcc's -Wunused-result was already relaxed), two assert-only Windows BOOLs in CPUInfo are marked [[maybe_unused]], and the test tbb.dll copy is guarded on the bundled tbb target so USE_SYSTEM_TBB works. Open3D's bundled host third-party that does not build under clang on Windows (TBB, libjpeg-turbo, curl, openssl, embree, zeromq) is supplied through the existing USE_SYSTEM_* options, and the standard Windows defines (_USE_MATH_DEFINES, WIN32) are passed in the build invocation.

This work was done with the assistance of Claude (an AI assistant by Anthropic), which did the CUDA-to-HIP analysis, the compat layer, the build-system changes, and the validation.

Test Plan

Validated on a Linux AMD GPU (Instinct MI250X, one gfx90a GCD, ROCm 7.2.1; the PermuteDevices CPU parameterization is the per-test reference) and on gfx1100 (RDNA3).

cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_PREFIX_PATH=/opt/rocm \
  -DBUILD_PYTORCH_OPS=OFF -DBUILD_TENSORFLOW_OPS=OFF -DBUILD_GUI=OFF \
  -DBUILD_WEBRTC=OFF -DBUILD_EXAMPLES=OFF -DBUILD_BENCHMARKS=OFF \
  -DBUILD_PYTHON_MODULE=OFF -DBUILD_UNIT_TESTS=ON -DBUILD_ISPC_MODULE=OFF \
  -DWITH_IPP=OFF -DCMAKE_BUILD_TYPE=Release -DDEVELOPER_BUILD=ON
cmake --build build -j16 --target tests
HIP_VISIBLE_DEVICES=3 ./build/bin/tests \
  --gtest_filter='*Tensor*:*Reduction*:*MemoryManager*:*NearestNeighbor*:*Knn*:*FixedRadius*:*Registration*:*Feature*'
HIP_VISIBLE_DEVICES=3 ./build/bin/tests --gtest_filter='*HashMap*:*VoxelBlockGrid*'
HIP_VISIBLE_DEVICES=3 ./build/bin/tests --gtest_repeat=30 --gtest_break_on_failure \
  --gtest_filter='*HashMapPermuteDevices.Reserve*:*HashMapPermuteDevices.InsertComplexKeys*:*HashMapPermuteDevices.MultivalueInsertion*:*HashMapPermuteDevices.HashSet*'

Results: Tensor + Reduction + MemoryManager 421/421; NearestNeighborSearch / Knn / FixedRadius / Hybrid; HashMap 20/20 (the four dedup tests stable over 30 repeats); VoxelBlockGrid 14/14 (TSDF integrate / ray casting / IO, stable over 5 repeats); ICP Registration and FPFH Feature all pass.

Also built and run on Windows (Radeon RX 9070 XT, gfx1201 RDNA4) with the all-clang ROCm toolchain, the bundled host third-party supplied via vcpkg, and a dynamic CRT (-DSTATIC_WINDOWS_RUNTIME=OFF):

cmake -G Ninja -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1201 \
  -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -DCMAKE_HIP_COMPILER=clang++ \
  -DSTATIC_WINDOWS_RUNTIME=OFF \
  -DCMAKE_C_FLAGS="-D_USE_MATH_DEFINES" \
  -DCMAKE_CXX_FLAGS="-DNOMINMAX -DWIN32_LEAN_AND_MEAN -D_USE_MATH_DEFINES -DWIN32" \
  -DCMAKE_HIP_FLAGS="-DNOMINMAX -DWIN32_LEAN_AND_MEAN -D_USE_MATH_DEFINES -DWIN32" \
  -DUSE_SYSTEM_TBB=ON -DUSE_SYSTEM_JPEG=ON -DUSE_SYSTEM_CURL=ON -DUSE_SYSTEM_OPENSSL=ON \
  -DUSE_SYSTEM_EMBREE=ON -DUSE_SYSTEM_ZEROMQ=ON \
  -DBUILD_GUI=OFF -DBUILD_WEBRTC=OFF -DBUILD_EXAMPLES=OFF -DBUILD_BENCHMARKS=OFF \
  -DBUILD_PYTHON_MODULE=OFF -DBUILD_UNIT_TESTS=ON -DBUILD_ISPC_MODULE=OFF \
  -DWITH_IPP=OFF -DCMAKE_BUILD_TYPE=Release -DDEVELOPER_BUILD=ON
cmake --build build --target tests
build\bin\tests.exe --gtest_filter='*Tensor*:*Reduction*:*MemoryManager*:*HashMap*:*VoxelBlockGrid*:*NearestNeighbor*:*Knn*:*FixedRadius*:*Registration*:*Feature*'

The architecture-sensitive GPU paths all pass on RDNA4: the FAISS warp-select (NearestNeighbor / Knn / FixedRadius), the StdGPU hashmap (HashMap 20/20), and the TSDF VoxelBlockGrid (14/14). A small number of linear-algebra GPU tests (which call rocSOLVER getrf) are affected by a known ROCm code-object-loading issue on the multi-architecture toolchain build used here; it is not specific to Open3D and is unrelated to this change.

[update-docs]

Thanks for submitting this pull request! The maintainers of this repository would appreciate if you could update the CHANGELOG.md based on your changes.

[Copilot]

Pull request overview

This PR adds a ROCm/HIP build path for Open3D’s GPU “tensor” stack (open3d.core, open3d.t) by compiling the existing CUDA-spelled .cu sources as HIP under a new -DUSE_HIP=ON option, while keeping the NVIDIA/CUDA build behavior unchanged.

Changes:

• Add USE_HIP build option and CMake wiring to enable HIP, retag .cu sources as LANGUAGE HIP, force-include a CUDA→HIP compat header, and pin Thrust to the HIP backend.
• Introduce a small HIP compatibility layer (include shims + symbol aliases) and extend device/host guards across core/t code to treat __HIPCC__ similarly to __CUDACC__.
• Adjust dependencies, third-party builds (stdgpu HIP backend), and tests/docs for ROCm limitations (e.g., NPP image filters unsupported; SlabHash tests skipped under HIP).

Reviewed changes

Copilot reviewed 44 out of 44 changed files in this pull request and generated 8 comments.

Show a summary per file

File	Description
CMakeLists.txt	Adds USE_HIP option, forces GPU-module switches, and enables the HIP language when requested.
cmake/Open3DSetGlobalProperties.cmake	Adds USE_HIP compile definitions, include shims, Thrust backend pinning, and .cu→HIP retagging.
cmake/Open3DShowAndAbortOnWarning.cmake	Relaxes specific warnings for HIP toolchains to avoid -Werror build breaks.
3rdparty/find_dependencies.cmake	Adds HIP math library discovery/targets and disables NPP dependency under HIP.
3rdparty/stdgpu/stdgpu.cmake	Builds stdgpu using its HIP backend via ExternalProject and hip/clang toolchain.
3rdparty/zlib/zlib.cmake	Switches MSVC-gating to WIN32 for correct library naming under clang-on-Windows.
3rdparty/libpng/libpng.cmake	Switches MSVC-gating to WIN32 for correct library naming under clang-on-Windows.
cpp/open3d/core/hip/CUDAToHIP.h	Adds the CUDA→HIP compatibility shim (runtime symbol aliases, masks, Thrust/CUB routing).
cpp/open3d/core/hip/hip_compat/cuda.h	Redirects <cuda.h> to the HIP compat header in HIP builds.
cpp/open3d/core/hip/hip_compat/cuda_runtime.h	Redirects <cuda_runtime.h> to the HIP compat header in HIP builds.
cpp/open3d/core/hip/hip_compat/cuda_runtime_api.h	Redirects <cuda_runtime_api.h> to the HIP compat header in HIP builds.
cpp/open3d/core/hip/hip_compat/cub/cub.cuh	Redirects <cub/cub.cuh> to hipCUB in HIP builds.
cpp/open3d/core/CUDAUtils.h	Makes nvcc-only extended-lambda assertion a no-op under HIP.
cpp/open3d/core/CUDAUtils.cpp	Avoids nodiscard warning-as-error issues on HIP by explicitly discarding cudaGetLastError().
cpp/open3d/core/ParallelFor.h	Extends CUDA compilation guards to include HIP compilation.
cpp/open3d/core/hashmap/Dispatch.h	Extends CUDA compilation guard to include HIP compilation.
cpp/open3d/core/hashmap/CUDA/SlabMacros.h	Widens sync masks for HIP’s 64-bit mask requirements (while noting SlabHash limitations).
cpp/open3d/core/nns/NeighborSearchCommon.h	Extends host/device macro guards to include HIP compilation.
cpp/open3d/core/nns/kernel/WarpShuffle.cuh	Uses a configurable full-warp mask for CUDA/HIP shuffle intrinsics.
cpp/open3d/core/nns/kernel/Select.cuh	Uses full-warp mask for __any_sync in CUDA/HIP.
cpp/open3d/core/nns/kernel/PtxUtils.cuh	Adds HIP/clang-builtin replacements for inline PTX helpers.
cpp/open3d/core/nns/kernel/DeviceDefs.cuh	Documents/implements wave64 handling for FAISS-derived warp-select and uses __syncthreads() on HIP.
cpp/open3d/core/kernel/ReductionCUDA.cu	Makes shuffle mask type/constant HIP-safe (64-bit mask under HIP).
cpp/open3d/core/linalg/LinalgHeadersCUDA.h	Switches CUDA BLAS/SOLVER headers to hipBLAS/hipSOLVER via aliasing in HIP builds.
cpp/open3d/core/linalg/kernel/SVD3x3.h	Extends CUDA guards to include HIP compilation.
cpp/open3d/t/geometry/kernel/GeometryMacros.h	Extends atomic helpers and device guards to include HIP compilation.
cpp/open3d/t/geometry/kernel/ImageImpl.h	Extends CUDA guards to include HIP compilation for image kernels.
cpp/open3d/t/geometry/kernel/NPPImage.cpp	Adds HIP stubs that report NPP GPU image filtering unsupported at runtime.
cpp/open3d/t/geometry/kernel/PointCloudImpl.h	Extends CUDA guards to include HIP compilation for point cloud kernels.
cpp/open3d/t/geometry/kernel/TransformImpl.h	Extends CUDA guards to include HIP compilation for transform kernels.
cpp/open3d/t/geometry/kernel/TriangleMeshImpl.h	Extends CUDA guards to include HIP compilation for triangle mesh kernels.
cpp/open3d/t/geometry/kernel/VoxelBlockGridImpl.h	Extends CUDA guards to include HIP compilation for TSDF/VBG kernels and synchronizations.
cpp/open3d/t/pipelines/kernel/FeatureImpl.h	Extends CUDA guards to include HIP compilation for feature kernels.
cpp/open3d/t/pipelines/kernel/FillInLinearSystemImpl.h	Extends CUDA guards to include HIP compilation for alignment term kernels.
cpp/open3d/t/pipelines/kernel/RGBDOdometryJacobianImpl.h	Extends CUDA guards to include HIP compilation for odometry kernels.
cpp/open3d/t/pipelines/kernel/RegistrationImpl.h	Extends CUDA guards and template instantiation attributes to include HIP compilation.
cpp/open3d/t/pipelines/registration/RobustKernelImpl.h	Extends CUDA guards to include HIP compilation for robust kernels.
cpp/open3d/utility/MiniVec.h	Treats __HIPCC__ like __CUDACC__ for host/device function specifiers.
cpp/open3d/utility/CPUInfo.cpp	Suppresses clang-on-Windows unused-variable warnings for assert-only BOOL results.
cpp/tests/CMakeLists.txt	Avoids copying bundled tbb.dll when using system TBB; links hip runtime for HIP-enabled tests.
cpp/tests/core/HashMap.cpp	Skips the Slab backend in test enumeration when building with HIP.
cpp/tests/t/geometry/VoxelBlockGrid.cpp	Skips the Slab backend for HIP builds in backend enumeration helper.
docs/compilation.rst	Documents the ROCm/HIP build option and basic configuration.
CHANGELOG.md	Adds a changelog entry for the new ROCm/HIP build path.