diff --git a/CMakeLists.txt b/CMakeLists.txt index 7dd8752..40f2d1d 100644 --- a/CMakeLists.txt +++ b/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} diff --git a/README.md b/README.md index 686445f..2597dbc 100644 --- a/README.md +++ b/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 diff --git a/src/cudaMat4.hpp b/src/cudaMat4.hpp index bc16dda..f8e2495 100644 --- a/src/cudaMat4.hpp +++ b/src/cudaMat4.hpp @@ -1,7 +1,7 @@ #pragma once #include -#include +#include "cuda_to_hip.h" struct cudaMat3 { glm::vec3 x; diff --git a/src/cuda_to_hip.h b/src/cuda_to_hip.h new file mode 100644 index 0000000..119eaad --- /dev/null +++ b/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 + * + * 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 +#include + +// 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 +#include + +#endif diff --git a/src/glm_device.h b/src/glm_device.h new file mode 100644 index 0000000..9c56e35 --- /dev/null +++ b/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 + * + * Include this header instead of 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 +#include +#include diff --git a/src/main.cpp b/src/main.cpp index 7ee7ba2..b5b9aae 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -4,12 +4,12 @@ // Configuration // ================ -#define VISUALIZE 1 +#define VISUALIZE 0 #define STEP true #define CPU false -#define GPU_NAIVE false -#define GPU_OCTREE true -#define MODEL false; +#define GPU_NAIVE true +#define GPU_OCTREE false +#define MODEL 0 #define UNIFORM_GRID 0 #define COHERENT_GRID 0 @@ -75,6 +75,10 @@ bool init(int argc, char **argv) { int gpuDevice = 0; int device_count = 0; cudaGetDeviceCount(&device_count); + + // Set the GPU device first + cudaSetDevice(gpuDevice); + if (gpuDevice > device_count) { std::cout << "Error: GPU device number is greater than the number of devices!" @@ -90,6 +94,7 @@ bool init(int argc, char **argv) { ss << projectName << " [SM " << major << "." << minor << " " << deviceProp.name << "]"; deviceName = ss.str(); +#if VISUALIZE // Window setup stuff glfwSetErrorCallback(errorCallback); @@ -124,12 +129,20 @@ bool init(int argc, char **argv) { // Initialize drawing state initVAO(); - // Default to device ID 0. If you have more than one GPU and want to test a non-default one, - // change the device ID. - cudaGLSetGLDevice(0); + // Register GL buffers with CUDA/HIP using the graphics resource API + cudaError_t err1 = cudaGraphicsGLRegisterBuffer(&positionsResource, boidVBO_positions, cudaGraphicsRegisterFlagsNone); + if (err1 != cudaSuccess) { + std::cerr << "Warning: cudaGraphicsGLRegisterBuffer positions failed: " << cudaGetErrorString(err1) << std::endl; + } + cudaError_t err2 = cudaGraphicsGLRegisterBuffer(&velocitiesResource, boidVBO_velocities, cudaGraphicsRegisterFlagsNone); + if (err2 != cudaSuccess) { + std::cerr << "Warning: cudaGraphicsGLRegisterBuffer velocities failed: " << cudaGetErrorString(err2) << std::endl; + } - cudaGLRegisterBufferObject(boidVBO_positions); - cudaGLRegisterBufferObject(boidVBO_velocities); + updateCamera(); + initShaders(program); + glEnable(GL_DEPTH_TEST); +#endif // Initialize N-body simulation #if CPU @@ -139,15 +152,11 @@ bool init(int argc, char **argv) { #elif GPU_OCTREE ScanMatch::initSimulationGPUOCTREE(TRUE_N, coords); #endif - updateCamera(); - - initShaders(program); - - glEnable(GL_DEPTH_TEST); return true; } +#if VISUALIZE void initVAO() { std::unique_ptr bodies{ new GLfloat[4 * (N_FOR_VIS)] }; @@ -208,20 +217,20 @@ void initShaders(GLuint * program) { } } - //==================================== - // Main loop - //==================================== void runCUDA() { // Map OpenGL buffer object for writing from CUDA on a single GPU // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not // use this buffer - float4 *dptr = NULL; float *dptrVertPositions = NULL; float *dptrVertVelocities = NULL; + size_t size; - cudaGLMapBufferObject((void**)&dptrVertPositions, boidVBO_positions); - cudaGLMapBufferObject((void**)&dptrVertVelocities, boidVBO_velocities); + // Map the graphics resources + cudaGraphicsResource_t resources[2] = {positionsResource, velocitiesResource}; + cudaGraphicsMapResources(2, resources, 0); + cudaGraphicsResourceGetMappedPointer((void**)&dptrVertPositions, &size, positionsResource); + cudaGraphicsResourceGetMappedPointer((void**)&dptrVertVelocities, &size, velocitiesResource); // execute the kernel call Step Here if (STEP) { @@ -233,21 +242,18 @@ void initShaders(GLuint * program) { ScanMatch::stepICPGPU_OCTREE(); #endif } - #if VISUALIZE ScanMatch::copyPointCloudToVBO(dptrVertPositions, dptrVertVelocities, CPU); - #endif // unmap buffer object - - cudaGLUnmapBufferObject(boidVBO_positions); - cudaGLUnmapBufferObject(boidVBO_velocities); + cudaGraphicsUnmapResources(2, resources, 0); } +#endif void mainLoop() { +#if VISUALIZE double fps = 0; double timebase = 0; int frame = 0; int counter = 0; - //ScanMatch::unitTest(); while (!glfwWindowShouldClose(window)) { glfwPollEvents(); @@ -272,7 +278,6 @@ void initShaders(GLuint * program) { glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); - #if VISUALIZE glUseProgram(program[PROG_BOID]); glBindVertexArray(boidVAO); glPointSize((GLfloat)pointSize); @@ -283,11 +288,27 @@ void initShaders(GLuint * program) { glBindVertexArray(0); glfwSwapBuffers(window); - #endif counter++; } glfwDestroyWindow(window); glfwTerminate(); +#else + // Non-visual mode: run a fixed number of ICP iterations + std::cout << "Running ICP in non-visual mode..." << std::endl; + const int NUM_ITERATIONS = 10; + for (int i = 0; i < NUM_ITERATIONS; i++) { + if (STEP) { +#if CPU + ScanMatch::stepICPCPU(); +#elif GPU_NAIVE + ScanMatch::stepICPGPU_NAIVE(); +#elif GPU_OCTREE + ScanMatch::stepICPGPU_OCTREE(); +#endif + } + } + std::cout << "Completed " << NUM_ITERATIONS << " ICP iterations." << std::endl; +#endif } @@ -312,18 +333,23 @@ void initShaders(GLuint * program) { phi += (xpos - lastX) / width; theta -= (ypos - lastY) / height; theta = std::fmax(0.01f, std::fmin(theta, 3.14f)); +#if VISUALIZE updateCamera(); +#endif } else if (rightMousePressed) { zoom += (ypos - lastY) / height; zoom = std::fmax(0.1f, std::fmin(zoom, 5.0f)); +#if VISUALIZE updateCamera(); +#endif } lastX = xpos; lastY = ypos; } +#if VISUALIZE void updateCamera() { cameraPosition.x = zoom * sin(phi) * sin(theta); cameraPosition.z = zoom * cos(theta); @@ -341,3 +367,4 @@ void initShaders(GLuint * program) { glUniformMatrix4fv(location, 1, GL_FALSE, &projection[0][0]); } } +#endif diff --git a/src/main.hpp b/src/main.hpp index 6702736..3257d79 100644 --- a/src/main.hpp +++ b/src/main.hpp @@ -7,8 +7,7 @@ #include #include #include -#include -#include +#include "cuda_to_hip.h" #include #include #include "utilityCore.hpp" @@ -30,6 +29,10 @@ GLuint boidIBO = 0; GLuint displayImage; GLuint program[2]; +// Graphics resources for GPU-GL interop +cudaGraphicsResource_t positionsResource = nullptr; +cudaGraphicsResource_t velocitiesResource = nullptr; + const unsigned int PROG_BOID = 0; const float fovy = (float) (PI / 4); diff --git a/src/octree.h b/src/octree.h index e53de5a..34d7970 100644 --- a/src/octree.h +++ b/src/octree.h @@ -6,15 +6,17 @@ */ #pragma once -#include #include #include #include -#include +#include "cuda_to_hip.h" +// Include Thrust BEFORE glm_device.h (rocThrust backend selection) #include #include #include #include +// Include GLM with device support AFTER Thrust +#include "glm_device.h" #include "utilityCore.hpp" typedef unsigned long long octKey; diff --git a/src/pointcloud.cu b/src/pointcloud.cu index 6553b3d..238b48f 100644 --- a/src/pointcloud.cu +++ b/src/pointcloud.cu @@ -1,6 +1,4 @@ -#pragma once #include "pointcloud.h" -#include #define blockSize 128 #define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__) @@ -206,8 +204,8 @@ void pointcloud::pointCloudToVBOCPU(float *vbodptr_positions, float *vbodptr_rgb //Launching Kernels dim3 fullBlocksPerGrid((N + blockSize - 1) / blockSize); - kernCopyPositionsToVBO << > >(N, tempPos, vbodptr_positions, s_scale, vbo_offset); - kernCopyRGBToVBO << > >(N, tempRGB, vbodptr_rgb, s_scale, vbo_offset); + kernCopyPositionsToVBO <<>>(N, tempPos, vbodptr_positions, s_scale, vbo_offset); + kernCopyRGBToVBO <<>>(N, tempRGB, vbodptr_rgb, s_scale, vbo_offset); utilityCore::checkCUDAErrorWithLine("copyPointCloudToVBO failed!"); cudaDeviceSynchronize(); @@ -274,7 +272,7 @@ void pointcloud::buildCoordsGPU(std::vector coords) { glm::vec3 t(9.0f, 0.f, 0.f); //glm::vec3 t(0.8f, 0.f, 0.f); glm::mat4 rotationMatrix = glm::rotate(angle, axis); - kernRotTrans << > > (dev_pos, rotationMatrix, t, N); + kernRotTrans <<>> (dev_pos, rotationMatrix, t, N); } } diff --git a/src/pointcloud.h b/src/pointcloud.h index e87e8d0..ebaad88 100644 --- a/src/pointcloud.h +++ b/src/pointcloud.h @@ -5,15 +5,17 @@ * @date 2019 */ #pragma once -#include #include #include #include -#include +#include "cuda_to_hip.h" +// Include Thrust BEFORE glm_device.h (rocThrust backend selection) #include #include #include #include +// Include GLM with device support AFTER Thrust +#include "glm_device.h" #include "utilityCore.hpp" class pointcloud { diff --git a/src/scanmatch.cu b/src/scanmatch.cu index 9f34f80..5f63866 100644 --- a/src/scanmatch.cu +++ b/src/scanmatch.cu @@ -1,9 +1,6 @@ -#define GLM_FORCE_CUDA #include -#include #include -#include -#include "utilityCore.hpp" +#include "cuda_to_hip.h" #include "scanmatch.h" #include "svd3.h" diff --git a/src/scanmatch.h b/src/scanmatch.h index 8c5f06d..83c0083 100644 --- a/src/scanmatch.h +++ b/src/scanmatch.h @@ -2,13 +2,14 @@ #include #include +#include +#include +#include "cuda_to_hip.h" +// Include Thrust BEFORE glm_device.h (rocThrust backend selection) #include #include #include #include -#include -#include -#include #include "pointcloud.h" #include "utilityCore.hpp" #include "octree.h" diff --git a/src/svd3.h b/src/svd3.h index a7018a6..faea60f 100644 --- a/src/svd3.h +++ b/src/svd3.h @@ -155,7 +155,7 @@ inline void approximateGivensQuaternion(float a11, float a12, float a22, float & // fast rsqrt function suffices // rsqrt2 (https://code.google.com/p/lppython/source/browse/algorithm/HDcode/newCode/rsqrt.c?r=26) // is even faster but results in too much error - float w = rsqrt(ch*ch + sh*sh); + float w = 1.0f / sqrtf(ch*ch + sh*sh); ch = b ? w*ch : (float)_cstar; sh = b ? w*sh : (float)_sstar; } @@ -281,7 +281,7 @@ void QRGivensQuaternion(float a1, float a2, float &ch, float &sh) ch = fabsf(a1) + fmaxf(rho, epsilon); bool b = a1 < 0; condSwap(b, sh, ch); - float w = rsqrt(ch*ch + sh*sh); + float w = 1.0f / sqrtf(ch*ch + sh*sh); ch *= w; sh *= w; } diff --git a/src/utilityCore.hpp b/src/utilityCore.hpp index 6a675f9..c2b0041 100644 --- a/src/utilityCore.hpp +++ b/src/utilityCore.hpp @@ -4,7 +4,7 @@ #include #include #include -#include +#include "cuda_to_hip.h" #include #include #include