[ROCm] Add HIP/ROCm support for AMD GPUs

CMakeLists.txt

@@ -1,4 +1,4 @@
-cmake_minimum_required(VERSION 3.1)
+cmake_minimum_required(VERSION 3.18)
 project(cis565_ScanMatching)
 
 set_property(GLOBAL PROPERTY USE_FOLDERS ON)
@@ -9,8 +9,8 @@ set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
 set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin)
 
-# Enable C++11 for host code
-set(CMAKE_CXX_STANDARD 11)
+# Enable C++17 for host code (required for rocPRIM/rocThrust)
+set(CMAKE_CXX_STANDARD 17)
 
 # Set a default build type if none was specified
 if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
@@ -20,19 +20,30 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
 endif()
 
 ########################################
-# CUDA Setup
+# HIP/CUDA Setup
 ########################################
-find_package(CUDA 10 REQUIRED)
-include(${CMAKE_MODULE_PATH}/CUDAComputesList.cmake)
-
-list(APPEND CUDA_NVCC_FLAGS ${CUDA_GENERATE_CODE})
-list(APPEND CUDA_NVCC_FLAGS_DEBUG "-g -G")
-set(CUDA_VERBOSE_BUILD ON)
-
-if(WIN32)
-    # Set up include and lib paths
-    set(CUDA_HOST_COMPILER ${CMAKE_CXX_COMPILER} CACHE FILEPATH "Host side compiler used by NVCC" FORCE)
-endif(WIN32)
+option(USE_HIP "Build with HIP for AMD GPUs" OFF)
+
+if(USE_HIP)
+    # enable_language(HIP) honors -DCMAKE_HIP_ARCHITECTURES, else auto-detects the
+    # host GPU, else errors on a no-GPU host.
+    enable_language(HIP)
+
+    # rocThrust is a drop-in for CUDA Thrust
+    find_package(rocthrust REQUIRED)
+else()
+    find_package(CUDA 10 REQUIRED)
+    include(${CMAKE_MODULE_PATH}/CUDAComputesList.cmake)
+
+    list(APPEND CUDA_NVCC_FLAGS ${CUDA_GENERATE_CODE})
+    list(APPEND CUDA_NVCC_FLAGS_DEBUG "-g -G")
+    set(CUDA_VERBOSE_BUILD ON)
+
+    if(WIN32)
+        # Set up include and lib paths
+        set(CUDA_HOST_COMPILER ${CMAKE_CXX_COMPILER} CACHE FILEPATH "Host side compiler used by NVCC" FORCE)
+    endif(WIN32)
+endif()
 ########################################
 
 find_package(OpenGL REQUIRED)
@@ -62,7 +73,9 @@ find_package(GLM REQUIRED)
 include_directories(${GLM_INCLUDE_DIRS})
 
 set(headers
+    src/cuda_to_hip.h
     src/cudaMat4.hpp
+    src/glm_device.h
     src/glslUtility.hpp
     src/scanmatch.h
     src/octree.h
@@ -71,23 +84,37 @@ set(headers
     src/utilityCore.hpp
     )
 
-set(sources
-    src/glslUtility.cpp
+set(cuda_sources
     src/scanmatch.cu
     src/octree.cu
-    src/main.cpp
     src/pointcloud.cu
+    )
+
+set(sources
+    src/glslUtility.cpp
+    src/main.cpp
     src/utilityCore.cpp
     )
 
 list(SORT headers)
+list(SORT cuda_sources)
 list(SORT sources)
 
 source_group(Headers FILES ${headers})
-source_group(Sources FILES ${sources})
-
-cuda_add_executable(${CMAKE_PROJECT_NAME} ${sources} ${headers})
-target_link_libraries(${CMAKE_PROJECT_NAME} ${LIBRARIES})
+source_group(Sources FILES ${sources} ${cuda_sources})
+
+if(USE_HIP)
+    # Mark all sources as HIP language so they compile with hipcc
+    # This ensures HIP headers work correctly in all sources
+    set_source_files_properties(${cuda_sources} ${sources} PROPERTIES LANGUAGE HIP)
+
+    add_executable(${CMAKE_PROJECT_NAME} ${sources} ${cuda_sources} ${headers})
+    target_compile_definitions(${CMAKE_PROJECT_NAME} PRIVATE USE_HIP GLM_FORCE_CUDA)
+    target_link_libraries(${CMAKE_PROJECT_NAME} ${LIBRARIES} roc::rocthrust)
+else()
+    cuda_add_executable(${CMAKE_PROJECT_NAME} ${sources} ${cuda_sources} ${headers})
+    target_link_libraries(${CMAKE_PROJECT_NAME} ${LIBRARIES})
+endif()
 
 add_custom_command(
     TARGET ${CMAKE_PROJECT_NAME}

README.md

@@ -121,7 +121,16 @@ It is recommended that you use Nsight. Nsight is shipped with CUDA. If you set u
 5. Right click and *Refresh* the project.
 6. From the *Run* menu, *Run*. Select "Local C/C++ Application" and the
    `cis565_` binary.
-
+
+### AMD GPUs (ROCm/HIP)
+The project also builds for AMD GPUs through HIP. Configure with `USE_HIP=ON` and select the target architecture with `CMAKE_HIP_ARCHITECTURES` (for example `gfx90a` for CDNA2, `gfx1100` for RDNA3, `gfx1201` for RDNA4). A ROCm install (the compiler and the roc/hip libraries, including rocThrust) is required.
+```bash
+mkdir build && cd build
+cmake .. -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_BUILD_TYPE=Release
+cmake --build . --parallel
+```
+The CUDA sources are translated to HIP through the `src/cuda_to_hip.h` compatibility header, so the same code compiles for NVIDIA (default) and AMD (`USE_HIP=ON`) GPUs.
+
   ## Running the Code
   The file `main.cpp` should have the following at the top:
   ```C

src/cudaMat4.hpp

@@ -1,7 +1,7 @@
 #pragma once
 
 #include <glm/glm.hpp>
-#include <cuda_runtime.h>
+#include "cuda_to_hip.h"
 
 struct cudaMat3 {
     glm::vec3 x;

src/cuda_to_hip.h

@@ -0,0 +1,66 @@
+/**
+ * @file cuda_to_hip.h
+ * @brief CUDA-to-HIP compatibility header for ROCm builds
+ *
+ * Copyright (c) 2026 Advanced Micro Devices, Inc.
+ * @author Jeff Daily <jeff.daily@amd.com>
+ *
+ * On ROCm this header maps CUDA runtime and GL-interop symbols to their HIP
+ * equivalents. On CUDA it is a pass-through to the CUDA headers.
+ *
+ * GLM device function support: GLM 0.9.9 uses __CUDACC__ + CUDA_VERSION>=7000
+ * to apply __host__ __device__ qualifiers. We define these BEFORE including
+ * GLM. rocThrust is designed to work with __HIP__ (defined by hipcc) and will
+ * select the HIP backend correctly as long as it sees __HIP__ before we define
+ * __CUDACC__. Since Thrust's backend selection happens at include time, we
+ * define the GLM macros here (after hip_runtime.h) and rely on include order
+ * in source files: include Thrust headers BEFORE glm headers.
+ *
+ * This port was authored with assistance from Claude, an AI assistant by
+ * Anthropic.
+ */
+#pragma once
+
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_gl_interop.h>
+
+// Runtime API
+#define cudaMalloc                  hipMalloc
+#define cudaFree                    hipFree
+#define cudaMemcpy                  hipMemcpy
+#define cudaMemset                  hipMemset
+#define cudaMemcpyHostToDevice      hipMemcpyHostToDevice
+#define cudaMemcpyDeviceToHost      hipMemcpyDeviceToHost
+#define cudaDeviceSynchronize       hipDeviceSynchronize
+#define cudaThreadSynchronize       hipDeviceSynchronize  // deprecated alias
+
+// Error handling
+#define cudaError_t                 hipError_t
+#define cudaSuccess                 hipSuccess
+#define cudaGetLastError            hipGetLastError
+#define cudaGetErrorString          hipGetErrorString
+
+// Device properties
+#define cudaDeviceProp              hipDeviceProp_t
+#define cudaGetDeviceCount          hipGetDeviceCount
+#define cudaGetDeviceProperties     hipGetDeviceProperties
+#define cudaSetDevice               hipSetDevice
+
+// OpenGL interop (using the modern graphics resource API)
+// The old cudaGL* functions are deprecated; HIP uses hipGraphics* API
+#define cudaGraphicsResource_t       hipGraphicsResource_t
+#define cudaGraphicsGLRegisterBuffer hipGraphicsGLRegisterBuffer
+#define cudaGraphicsRegisterFlagsNone hipGraphicsRegisterFlagsNone
+#define cudaGraphicsMapResources     hipGraphicsMapResources
+#define cudaGraphicsResourceGetMappedPointer hipGraphicsResourceGetMappedPointer
+#define cudaGraphicsUnmapResources   hipGraphicsUnmapResources
+#define cudaGraphicsUnregisterResource hipGraphicsUnregisterResource
+
+#else  // CUDA path
+
+#include <cuda_runtime.h>
+#include <cuda_gl_interop.h>
+
+#endif

src/glm_device.h

@@ -0,0 +1,34 @@
+/**
+ * @file glm_device.h
+ * @brief GLM with device function support for HIP/CUDA
+ *
+ * Copyright (c) 2026 Advanced Micro Devices, Inc.
+ * @author Jeff Daily <jeff.daily@amd.com>
+ *
+ * Include this header instead of <glm/glm.hpp> when GLM is used in device code.
+ * On HIP, this enables __host__ __device__ qualifiers on GLM math functions by
+ * defining __CUDACC__ and CUDA_VERSION before including GLM.
+ *
+ * IMPORTANT: Include Thrust headers BEFORE this header. rocThrust's backend
+ * selection checks __CUDACC__ before __HIP__, so defining __CUDACC__ before
+ * including Thrust would select the CUDA backend incorrectly.
+ */
+#pragma once
+
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+// GLM 0.9.9 uses __CUDACC__ + CUDA_VERSION>=7000 to apply __host__ __device__
+// qualifiers. Define these so GLM treats hipcc like nvcc.
+#ifndef __CUDACC__
+#define __CUDACC__ 1
+#endif
+#ifndef CUDA_VERSION
+#define CUDA_VERSION 8000
+#endif
+#ifndef GLM_FORCE_CUDA
+#define GLM_FORCE_CUDA
+#endif
+#endif
+
+#include <glm/glm.hpp>
+#include <glm/gtc/matrix_transform.hpp>
+#include <glm/gtx/transform.hpp>