From ce20c461b35aaa5781724cc0fd2742f8604c4b89 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 11 Jun 2026 21:39:38 +0000 Subject: [PATCH 1/2] [ROCm] Add HIP/ROCm support for AMD GPUs Add AMD GPU support to the CUDA scan-matcher (ICP point-cloud registration) via HIP. A single src/cuda_to_hip.h compatibility header maps the CUDA runtime and OpenGL-interop symbols used by the project to their HIP equivalents, so the same sources compile for NVIDIA (default) and AMD (-DUSE_HIP=ON) GPUs. The GL interop is updated to the modern graphics-resource API (hipGraphicsGLRegisterBuffer and friends), which HIP provides and which deprecates the old cudaGL* entry points. GLM device-function support is enabled through src/glm_device.h, which defines the qualifier macros GLM expects so its math functions get __host__ __device__ under hipcc. rocThrust is a drop-in for CUDA Thrust; it requires C++17, so the host standard is raised accordingly. Include order matters: Thrust must precede GLM so rocThrust selects the HIP backend. To review, start with src/cuda_to_hip.h and the USE_HIP branch of CMakeLists.txt, then the GL interop and device-math call sites in src/main.cpp and the .cu sources. The README documents the AMD build alongside the existing Windows and Linux CUDA instructions. This work was authored with assistance from Claude, an AI assistant by Anthropic. Test Plan: ``` cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_BUILD_TYPE=Release cmake --build build -j HIP_VISIBLE_DEVICES=0 ./build/bin/cis565_ScanMatching # 10 ICP iterations complete, ~440us NN timing per step on gfx90a (MI250X) ``` Built and run on gfx90a (MI250X, CDNA2), gfx1100 (RDNA3), and gfx1201 (RDNA4); the headless ICP run completes 10 iterations on each. --- CMakeLists.txt | 75 ++++++++++++++++++++++++++++------------ README.md | 11 +++++- src/cudaMat4.hpp | 2 +- src/cuda_to_hip.h | 66 +++++++++++++++++++++++++++++++++++ src/glm_device.h | 34 +++++++++++++++++++ src/main.cpp | 83 ++++++++++++++++++++++++++++++--------------- src/main.hpp | 7 ++-- src/octree.h | 6 ++-- src/pointcloud.cu | 8 ++--- src/pointcloud.h | 6 ++-- src/scanmatch.cu | 5 +-- src/scanmatch.h | 7 ++-- src/svd3.h | 4 +-- src/utilityCore.hpp | 2 +- 14 files changed, 243 insertions(+), 73 deletions(-) create mode 100644 src/cuda_to_hip.h create mode 100644 src/glm_device.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 7dd8752..12b6618 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,34 @@ 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) + + # Default to gfx90a if no architecture was specified; override with + # -DCMAKE_HIP_ARCHITECTURES (e.g. gfx1100 for RDNA3, gfx1201 for RDNA4). + if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "") + set(CMAKE_HIP_ARCHITECTURES "gfx90a") + endif() + + # 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 +77,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 +88,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 From 26073aa3ee515b2dfa50db2706c3ce6ae78e4eee Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 23 Jun 2026 22:18:19 +0000 Subject: [PATCH 2/2] [ROCm] Rely on HIP arch auto-detect instead of pinning gfx90a The default-gfx90a block preempted and duplicated CMake's own host-GPU detection. Removing it lets enable_language(HIP) honor an explicit -DCMAKE_HIP_ARCHITECTURES, auto-detect the host GPU, or error on a no-GPU build host. The block was also dead code: it ran after enable_language(HIP), by which point CMAKE_HIP_ARCHITECTURES was already set, so its guard was always false. This work was authored with the Claude AI assistant. --- CMakeLists.txt | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 12b6618..40f2d1d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -25,14 +25,10 @@ endif() 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) - # Default to gfx90a if no architecture was specified; override with - # -DCMAKE_HIP_ARCHITECTURES (e.g. gfx1100 for RDNA3, gfx1201 for RDNA4). - if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "") - set(CMAKE_HIP_ARCHITECTURES "gfx90a") - endif() - # rocThrust is a drop-in for CUDA Thrust find_package(rocthrust REQUIRED) else()