From 5abe276676a8d40a55f46df8724967d671e94288 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Fri, 19 Jun 2026 04:09:10 +0000 Subject: [PATCH 1/2] [ROCm] Add AMD GPU support (HIP software + HIPRT backends) Add the ability to build and run barney's GPU rendering on AMD GPUs with HIP for ROCm, alongside the existing CUDA/OptiX path. The support is additive and gated behind a new USE_HIP CMake option (default OFF), so the CUDA/OptiX build is unchanged when it is not enabled. How to review, in order: 1. Build wiring (CMakeLists.txt, rtcore/CMakeLists.txt, barney/CMakeLists.txt): under USE_HIP, enable_language(HIP) and set the HIP architecture from CMAKE_HIP_ARCHITECTURES (default gfx90a when unset), and retag the reused CUDA backend sources LANGUAGE HIP. The new backend slots into the existing rtcore backend selection without disturbing the CUDA/OptiX path. 2. The compat shim (rtcore/cudaCommon/cuda_to_hip.h): force-included only on the HIP build, it maps the small set of CUDA runtime spellings the shared backend uses to their hipXxx equivalents and otherwise falls through to , so the CUDA translation units are unchanged. 3. Software ray-tracing backend: the CUDA device code is reused under HIP, traversing the cuBQL BVH; no vendor RT hardware is required. OptiX is NVIDIA-only and is not used in this configuration. 4. HIPRT hardware-traversal backend (rtcore/hiprt/, -DBARNEY_BACKEND_HIPRT=ON): the AMD analogue of the OptiX backend, derived from the CUDA backend; it builds HIPRT scenes and maps barney's anyHit/closestHit/intersect programs onto HIPRT's filter and custom-geometry function tables. 5. cmake/Findhiprt.cmake locates the HIPRT SDK as a discovered dependency (never vendored), the way barney finds OptiX. Building for AMD GPUs: add -DUSE_HIP=ON and select the architecture with -DCMAKE_HIP_ARCHITECTURES= (defaults to gfx90a); add -DBARNEY_BACKEND_HIPRT=ON -Dhiprt_ROOT= for HIPRT traversal. Validated on Linux (gfx90a CDNA2, gfx1100 RDNA3) and Windows (gfx1201 RDNA4): both backends build, and the bundled validation scenes (triangles, spheres, cylinders, instances, and the opaque/transparent path-traced scenes) render with pixel statistics identical to the CUDA reference; the HIPRT custom-geometry path is pixel-identical to the software backend. The CUDA/OptiX build is unchanged when USE_HIP is off. This work was authored with the assistance of Claude (Anthropic). Test Plan: ``` # software backend (gfx90a) cmake -S . -B build-hip -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_BUILD_TYPE=Release cmake --build build-hip -j # HIPRT hardware backend cmake -S . -B build-hiprt -DUSE_HIP=ON -DBARNEY_BACKEND_HIPRT=ON \ -Dhiprt_ROOT= -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_BUILD_TYPE=Release cmake --build build-hiprt -j # render the validation scenes and compare pixel stats to the CUDA reference ``` Signed-off-by: Jeff Daily --- CMakeLists.txt | 75 +++++- README.md | 17 ++ anari/CMakeLists.txt | 15 +- barney/CMakeLists.txt | 48 +++- barney/LocalContext.cpp | 44 +++- barney/api/barney.cu | 18 +- barney/include/barney.h | 3 + cmake/Findhiprt.cmake | 64 +++++ rtcore/AppInterface.h | 5 + rtcore/CMakeLists.txt | 85 ++++++- rtcore/ComputeInterface.h | 5 + rtcore/TraceInterface.h | 4 + rtcore/common/rtcore-common.h | 2 +- rtcore/cuda/TraceInterface.h | 4 + rtcore/cudaCommon/ComputeInterface.h | 4 +- rtcore/cudaCommon/cuda-common.h | 2 +- rtcore/cudaCommon/cuda-helper.h | 2 +- rtcore/cudaCommon/cuda_to_hip.h | 105 ++++++++ rtcore/hiprt/AppInterface.h | 37 +++ rtcore/hiprt/Buffer.cpp | 48 ++++ rtcore/hiprt/Buffer.h | 31 +++ rtcore/hiprt/ComputeInterface.h | 26 ++ rtcore/hiprt/Device.cpp | 96 +++++++ rtcore/hiprt/Device.h | 97 ++++++++ rtcore/hiprt/Geom.cpp | 36 +++ rtcore/hiprt/Geom.h | 68 +++++ rtcore/hiprt/GeomType.cpp | 38 +++ rtcore/hiprt/GeomType.h | 64 +++++ rtcore/hiprt/Group.cpp | 357 +++++++++++++++++++++++++++ rtcore/hiprt/Group.h | 126 ++++++++++ rtcore/hiprt/ProgramInterface.h | 255 +++++++++++++++++++ rtcore/hiprt/TraceInterface.h | 286 +++++++++++++++++++++ rtcore/hiprt/TraceKernel.cpp | 112 +++++++++ rtcore/hiprt/TraceKernel.h | 43 ++++ 34 files changed, 2203 insertions(+), 19 deletions(-) create mode 100644 cmake/Findhiprt.cmake create mode 100644 rtcore/cudaCommon/cuda_to_hip.h create mode 100644 rtcore/hiprt/AppInterface.h create mode 100644 rtcore/hiprt/Buffer.cpp create mode 100644 rtcore/hiprt/Buffer.h create mode 100644 rtcore/hiprt/ComputeInterface.h create mode 100644 rtcore/hiprt/Device.cpp create mode 100644 rtcore/hiprt/Device.h create mode 100644 rtcore/hiprt/Geom.cpp create mode 100644 rtcore/hiprt/Geom.h create mode 100644 rtcore/hiprt/GeomType.cpp create mode 100644 rtcore/hiprt/GeomType.h create mode 100644 rtcore/hiprt/Group.cpp create mode 100644 rtcore/hiprt/Group.h create mode 100644 rtcore/hiprt/ProgramInterface.h create mode 100644 rtcore/hiprt/TraceInterface.h create mode 100644 rtcore/hiprt/TraceKernel.cpp create mode 100644 rtcore/hiprt/TraceKernel.h diff --git a/CMakeLists.txt b/CMakeLists.txt index ff6d5c8d..47dc973a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -32,18 +32,54 @@ set(BARNEY_VERSION ${BARNEY_VERSION_MAJOR}.${BARNEY_VERSION_MINOR}.${BARNEY_VERS project(Barney VERSION ${BARNEY_VERSION} LANGUAGES C CXX) include(GNUInstallDirs) -set(BARNEY_RTC_EXT +set(BARNEY_RTC_EXT "" CACHE STRING "Experimental/external rtcore config.") mark_as_advanced(USE_EXT) -set(BARNEY_DEVICE_NAME +# barney's own cmake modules (e.g. Findhiprt.cmake for the HIPRT backend) +list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake") + +# AMD GPU support: compile barney's OptiX-free software ray-tracing backend +# (rtcore/cuda + rtcore/cudaCommon) through HIP/ROCm. When USE_HIP is ON the +# .cu sources are compiled with hipcc (LANGUAGE HIP); the CUDA/OptiX build is +# unaffected (this whole path is inert unless USE_HIP is requested). +option(USE_HIP "Build GPU code with HIP for AMD GPUs" OFF) + +set(BARNEY_DEVICE_NAME "barney" CACHE STRING "Name of device to be built.") if ((NOT BARNEY_RTC_EXT STREQUAL "") AND (EXISTS "${BARNEY_RTC_EXT}")) set(BARNEY_HAVE_EXT ON) include(${BARNEY_RTC_EXT}/config.cmake) +elseif (USE_HIP) + enable_language(HIP) + if (NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "") + set(CMAKE_HIP_ARCHITECTURES "gfx90a") + endif() + message("#barney: building GPU code with HIP for CMAKE_HIP_ARCHITECTURES=${CMAKE_HIP_ARCHITECTURES}") + # the cuda backend's GPU code is reused under HIP; OptiX is NVIDIA-only. + set(BARNEY_HAVE_HIP ON) + set(BARNEY_HAVE_CUDA OFF) + set(CMAKE_CUDA_ARCHITECTURES) + if (WIN32) + # On Windows, CMake's Windows-Clang platform module injects -fuse-ld=lld-link + # into HIP link commands, but the AMD clang driver rejects it when doing HIP + # device-link (--hip-link); lld-link is the default host linker already. + set(CMAKE_HIP_USING_LINKER_DEFAULT "") + endif() + + # Optional hardware-RT backend on AMD GPUs via AMD HIPRT. HIPRT supplies the + # BVH build + ray traversal (hardware-accelerated on RDNA2+, software on CDNA + # such as gfx90a); barney keeps its function-pointer shading dispatch. HIPRT + # is a discovered dependency (find via hiprt_ROOT / HIPRT_PATH), never + # vendored -- mirrors how barney finds OptiX/OIDN. + option(BARNEY_BACKEND_HIPRT "Enable HIPRT hardware-RT backend (AMD)?" OFF) + if (BARNEY_BACKEND_HIPRT) + find_package(hiprt REQUIRED) + message("#barney: HIPRT backend enabled (hiprt at ${hiprt_LIBRARY})") + endif() elseif (NOT BARNEY_DISABLE_CUDA) include(CheckLanguage) check_language(CUDA) @@ -122,6 +158,16 @@ target_link_libraries(barney_config INTERFACE owl-config ) +if (BARNEY_HAVE_HIP) + # let the host (.cpp) translation units that include the backend's compat + # header (cuda_to_hip.h) select the HIP path; the .cu/HIP passes already see + # __HIPCC__, but plain-C++ TUs need this define. + target_compile_definitions(barney_config INTERFACE USE_HIP=1 BARNEY_HAVE_HIP=1) + # hip::host carries the ROCm include dirs and __HIP_PLATFORM_AMD__ so the + # host C++ TUs that call the hip runtime (hipMalloc/hipMemcpy/...) build. + find_package(hip REQUIRED) + target_link_libraries(barney_config INTERFACE hip::host) +endif() target_include_directories(barney_config INTERFACE $ $ @@ -142,6 +188,19 @@ endif() if (BARNEY_HAVE_EXT) # leave it to the external plugins to configure themselves - not # currently working +elseif (BARNEY_HAVE_HIP) + # on AMD GPUs the default is the OptiX-free software tracer (the cuda backend); + # OptiX has no HIP equivalent here, embree is the CPU fallback. When the HIPRT + # backend is requested it is the GPU backend instead (one rtc backend per + # build); cuda_common (textures) is still built because the HIPRT backend + # reuses it, exactly as the optix backend does. + if (BARNEY_BACKEND_HIPRT) + set(BARNEY_BACKEND_CUDA OFF) + else() + option(BARNEY_BACKEND_CUDA "Enable software-tracer GPU backend?" ON) + endif() + set(BARNEY_BACKEND_OPTIX OFF) + set(BARNEY_BACKEND_EMBREE OFF) elseif (BARNEY_HAVE_CUDA) option(BARNEY_BACKEND_CUDA "Enable (native-)CUDA Backend?" OFF) option(BARNEY_BACKEND_OPTIX "Enable OptiX Backend?" ON) @@ -157,6 +216,9 @@ option(BARNEY_DISABLE_DENOISING "DISable denoising" OFF) option(BARNEY_HAVE_NANOVDB "Include Support for NanoVDB" ON) option(BARNEY_USE_EXTERNAL_CUBQL "Use External CuBQL dir" OFF) +set(BARNEY_EXTERNAL_CUBQL_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../cuBQL" CACHE PATH + "Path to an external cuBQL source tree (when BARNEY_USE_EXTERNAL_CUBQL is ON)") +set(BARNEY_CUBQL_BVH_WIDTH 4 CACHE INT "CuBQL BVH width") # ================================================================== if (BARNEY_DISABLE_DENOISING) @@ -184,18 +246,21 @@ endif() # ------------------------------------------------------------------ if (NOT (TARGET cuBQL)) - if (BARNEY_HAVE_CUDA) + if (BARNEY_HAVE_HIP) + # build cuBQL's GPU code through HIP for the same arch(s) as barney + set(CUBQL_USE_HIP ON) + elseif (BARNEY_HAVE_CUDA) else() set(CUBQL_DISABLE_CUDA ON) endif() if (BARNEY_USE_EXTERNAL_CUBQL) - add_subdirectory(../cuBQL EXCLUDE_FROM_ALL builddir_cuBQL) + add_subdirectory(${BARNEY_EXTERNAL_CUBQL_DIR} EXCLUDE_FROM_ALL builddir_cuBQL) else() add_subdirectory(submodules/cuBQL EXCLUDE_FROM_ALL) endif() endif() -if (BARNEY_HAVE_CUDA) +if (BARNEY_HAVE_CUDA OR BARNEY_HAVE_HIP) option(BARNEY_CUBQL_HOST "Use CUBQL host builder" OFF) else() set(BARNEY_CUBQL_HOST ON) diff --git a/README.md b/README.md index afd0ec9e..1d8624b9 100644 --- a/README.md +++ b/README.md @@ -113,6 +113,16 @@ For CUDA/OptiX Acceleration, it also requires: - `OptiX`, as part of OWL. See documentation in OWL (https://github.com/NVIDIA/owl) for where to get, and how to best install for OWL to easily find it) + +For AMD GPU (ROCm/HIP) Acceleration, it instead requires: + +- `ROCm` (https://rocm.docs.amd.com/), version 7 and up, which provides `hipcc` + and the HIP runtime. The same GPU code that builds for CUDA is compiled with + HIP; OptiX is NVIDIA-only and is not used in this configuration. + +- optionally, `HIPRT` (https://github.com/GPUOpen-LibrariesAndSDKs/HIPRT) for + hardware-accelerated ray traversal. Without it, barney uses its own + (cuBQL-based) software traversal. For MPI-based data-parallel rendering: @@ -152,6 +162,13 @@ cmake --install .. [ --config Release ] By default Barney builds without MPI support; to enable this add `-DBARNEY_MPI=ON` to the cmake config command. +To build for AMD GPUs with HIP/ROCm instead of CUDA/OptiX, add `-DUSE_HIP=ON` +and select the target architecture with `-DCMAKE_HIP_ARCHITECTURES=` (for +example `gfx90a` for CDNA2, or `gfx1100` for RDNA3); when left unset it defaults +to `gfx90a`. To use HIPRT hardware traversal, also add `-DBARNEY_BACKEND_HIPRT=ON +-Dhiprt_ROOT=`. The CUDA/OptiX build is unchanged when `USE_HIP` +is off. + # Examples of Supported Geometry and Volume Types diff --git a/anari/CMakeLists.txt b/anari/CMakeLists.txt index 6014308f..01588d1c 100644 --- a/anari/CMakeLists.txt +++ b/anari/CMakeLists.txt @@ -97,9 +97,18 @@ if(WIN32) ${BARNEY_DEVICE_NAME}_static anari::helium ) - set_target_properties(anari_library_${BARNEY_DEVICE_NAME} PROPERTIES - LINK_FLAGS "/WHOLEARCHIVE:${BARNEY_DEVICE_NAME}_static" - ) + if (BARNEY_HAVE_HIP) + # When linking through the HIP device-link wrapper (clang++ --hip-link), + # /WHOLEARCHIVE must be passed via -Xlinker so clang forwards it to lld-link + # rather than treating it as an input file. + set_target_properties(anari_library_${BARNEY_DEVICE_NAME} PROPERTIES + LINK_FLAGS "-Xlinker /WHOLEARCHIVE:${BARNEY_DEVICE_NAME}_static.lib" + ) + else() + set_target_properties(anari_library_${BARNEY_DEVICE_NAME} PROPERTIES + LINK_FLAGS "/WHOLEARCHIVE:${BARNEY_DEVICE_NAME}_static.lib" + ) + endif() else() target_link_libraries(anari_library_${BARNEY_DEVICE_NAME} PRIVATE $ diff --git a/barney/CMakeLists.txt b/barney/CMakeLists.txt index 0513f089..d9ad17ef 100644 --- a/barney/CMakeLists.txt +++ b/barney/CMakeLists.txt @@ -33,6 +33,14 @@ macro(set_library_properties tgt) CUDA_RESOLVE_DEVICE_SYMBOLS ON CUDA_VISIBILITY_PRESET hidden ) +if (BARNEY_HAVE_HIP) + # objects were compiled -fgpu-rdc, so the final device-link of this target + # (and anything that links it) must pass -fgpu-rdc too, otherwise hipcc skips + # the device link and the __hip_gpubin_handle symbols stay undefined. + set_target_properties(${tgt} PROPERTIES HIP_SEPARABLE_COMPILATION ON) + target_link_options(${tgt} PUBLIC -fgpu-rdc --hip-link) + target_compile_options(${tgt} PRIVATE $<$:-fgpu-rdc>) +endif() if (APPLE) set_target_properties(${tgt} PROPERTIES INSTALL_RPATH "$loader_path") else() @@ -321,7 +329,35 @@ if (BARNEY_BACKEND_CUDA) set_library_properties(barney_mpi_cuda) endif() endif() - + +# ------------------------------------------------------------------ +# barney_hiprt: the same barney host + device programs compiled against the +# HIPRT (hardware-RT) rtcore backend. Mirrors the barney_cuda instantiation; +# HIPRT replaces the BVH build + traversal, barney keeps its function-pointer +# program dispatch, so the program set device-links exactly as the software +# backend (-fgpu-rdc, set by set_library_properties). +# ------------------------------------------------------------------ +if (BARNEY_BACKEND_HIPRT) + add_library(barney_hiprt_programs STATIC ${DEVICE_PROGRAM_SOURCES}) + target_link_libraries(barney_hiprt_programs PUBLIC barney_config barney_rtc_hiprt) + target_compile_definitions(barney_hiprt_programs PRIVATE -DBARNEY_DEVICE_PROGRAM=1) + set_library_properties(barney_hiprt_programs) + + add_library(barney_hiprt STATIC ${HOST_SOURCES}) + target_link_libraries(barney_hiprt PUBLIC barney_config barney_rtc_hiprt) + target_link_libraries(barney_hiprt PRIVATE + barney_hiprt_programs + barney_config + ) + set_library_properties(barney_hiprt) + + if (BARNEY_MPI) + add_library(barney_mpi_hiprt ${MPI_SOURCES}) + target_link_libraries(barney_mpi_hiprt PUBLIC barney_hiprt MPI::MPI_C) + set_library_properties(barney_mpi_hiprt) + endif() +endif() + # ============================================================================= # create frontends for barney and barney_static @@ -427,6 +463,16 @@ if (BARNEY_BACKEND_CUDA) endif() endif() +# link hiprt backend (if enabled) to all frontends +if (BARNEY_BACKEND_HIPRT) + target_link_libraries(${BARNEY_DEVICE_NAME} PRIVATE $) + target_link_libraries(${BARNEY_DEVICE_NAME}_static PRIVATE $) + if (BARNEY_MPI) + target_link_libraries(${BARNEY_DEVICE_NAME}_mpi PUBLIC $) + target_link_libraries(${BARNEY_DEVICE_NAME}_mpi_static PUBLIC $) + endif() +endif() + # ################################################################## diff --git a/barney/LocalContext.cpp b/barney/LocalContext.cpp index 89bbcfe7..be942ac7 100644 --- a/barney/LocalContext.cpp +++ b/barney/LocalContext.cpp @@ -129,7 +129,49 @@ namespace barney_api { // Context *ctx = new BARNEY_NS::LocalContext(dgIDs,gpuIDs); // return ctx; // } - } + } +#endif +#if BARNEY_RTC_HIPRT + extern "C" { + Context *createContext_hiprt(const std::vector &dgIDs, + int numGPUs, const int *gpuIDs) + { + if (FromEnv::get()->logBackend) + std::cout << "#bn: creating *hiprt (AMD hardware-RT)* context" << std::endl; + int numDGs = dgIDs.size(); + if (numGPUs == -1) { + BARNEY_CUDA_CALL(GetDeviceCount(&numGPUs)); + } + +#if ALLOW_OVERSUBSCRIBE + std::vector fakeIDs; + if (numGPUs < numDGs) { + for (int i=0;i localSlots(dgIDs.size()); + for (int lsIdx=0;lsIdx &dgIDs, int numGPUs, const int *gpuIDs); #endif +#if BARNEY_BACKEND_HIPRT + barney_api::Context * + createContext_hiprt(const std::vector &dgIDs, + int numGPUs, const int *gpuIDs); +#endif #if BARNEY_MPI # if BARNEY_BACKEND_EMBREE barney_api::Context * @@ -868,6 +873,8 @@ namespace barney_api { return (BNContext)createContext_optix(dataGroupIDs,numGPUs,_gpuIDs); #elif BARNEY_BACKEND_CUDA return (BNContext)createContext_cuda(dataGroupIDs,numGPUs,_gpuIDs); +#elif BARNEY_BACKEND_HIPRT + return (BNContext)createContext_hiprt(dataGroupIDs,numGPUs,_gpuIDs); #else throw std::runtime_error ("explicitly asked for GPU backend, " @@ -898,7 +905,16 @@ namespace barney_api { << e.what() << ")" << std::endl; } #endif - + +#if BARNEY_BACKEND_HIPRT + try { + return (BNContext)createContext_hiprt(dataGroupIDs,numGPUs,_gpuIDs); + } catch (std::exception &e) { + std::cerr << "#barney(warn): could not create hiprt backend (reason: " + << e.what() << ")" << std::endl; + } +#endif + # if BARNEY_BACKEND_EMBREE return (BNContext)createContext_embree(dataGroupIDs); #endif diff --git a/barney/include/barney.h b/barney/include/barney.h index 6f9300a2..12d250b6 100644 --- a/barney/include/barney.h +++ b/barney/include/barney.h @@ -17,6 +17,9 @@ /*! whether the embree backend has been included in this build */ #cmakedefine01 BARNEY_BACKEND_EMBREE +/*! whether the HIPRT (AMD hardware-RT) backend has been included in this build */ +#cmakedefine01 BARNEY_BACKEND_HIPRT + /*! whether this build of barney has support for MPI based rendering; i.e., wether barney/barney_mpi.h and (if anari is enabled) anari_library_barney_mpi exist */ diff --git a/cmake/Findhiprt.cmake b/cmake/Findhiprt.cmake new file mode 100644 index 00000000..7c0c1b22 --- /dev/null +++ b/cmake/Findhiprt.cmake @@ -0,0 +1,64 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# \author Jeff Daily + +# Locate the AMD HIPRT SDK (https://github.com/GPUOpen-LibrariesAndSDKs/HIPRT). +# Mirrors how barney finds OptiX/OIDN: a discovered dependency, never vendored, +# pointed at by hiprt_ROOT (or the HIPRT_PATH environment variable). HIPRT +# supplies the ray-tracing BVH build + traversal used by the rtcore/hiprt +# backend; on a GPU without hardware RT units (e.g. gfx90a) it falls back to a +# correct software BVH traversal. +# +# Sets: +# hiprt_FOUND +# hiprt::hiprt imported target (include dir + libhiprt) +# hiprt_INCLUDE_DIR +# hiprt_LIBRARY + +find_path(hiprt_INCLUDE_DIR + NAMES hiprt/hiprt.h + HINTS + ${hiprt_ROOT} + $ENV{hiprt_ROOT} + $ENV{HIPRT_PATH} + ) + +# HIPRT names its shared library libhiprt.so (e.g. libhiprt0300164.so); +# accept any so the consumer need not hardcode the exact build version. +find_library(hiprt_LIBRARY + NAMES hiprt hiprt0300264 hiprt0300164 hiprt0300064 + HINTS + ${hiprt_ROOT} + $ENV{hiprt_ROOT} + $ENV{HIPRT_PATH} + PATH_SUFFIXES + dist/bin/Release + dist/bin/Debug + bin + lib + ) + +if (NOT hiprt_LIBRARY AND hiprt_INCLUDE_DIR) + file(GLOB _hiprt_candidates + "${hiprt_INCLUDE_DIR}/dist/bin/Release/libhiprt*.so" + "${hiprt_INCLUDE_DIR}/dist/bin/Debug/libhiprt*.so" + "${hiprt_INCLUDE_DIR}/lib/libhiprt*.so" + "${hiprt_INCLUDE_DIR}/bin/libhiprt*.so") + if (_hiprt_candidates) + list(GET _hiprt_candidates 0 hiprt_LIBRARY) + endif() +endif() + +include(FindPackageHandleStandardArgs) +find_package_handle_standard_args(hiprt + REQUIRED_VARS hiprt_INCLUDE_DIR hiprt_LIBRARY) + +if (hiprt_FOUND AND NOT TARGET hiprt::hiprt) + add_library(hiprt::hiprt UNKNOWN IMPORTED) + set_target_properties(hiprt::hiprt PROPERTIES + IMPORTED_LOCATION "${hiprt_LIBRARY}" + INTERFACE_INCLUDE_DIRECTORIES "${hiprt_INCLUDE_DIR}") +endif() + +mark_as_advanced(hiprt_INCLUDE_DIR hiprt_LIBRARY) diff --git a/rtcore/AppInterface.h b/rtcore/AppInterface.h index 68e4645f..e787b489 100644 --- a/rtcore/AppInterface.h +++ b/rtcore/AppInterface.h @@ -13,6 +13,11 @@ namespace rtc { using namespace rtc::cuda; } #endif +#if BARNEY_RTC_HIPRT +# include "hiprt/AppInterface.h" +namespace rtc { using namespace rtc::hiprt; } +#endif + #if BARNEY_RTC_OPTIX # include "optix/AppInterface.h" namespace rtc { using namespace rtc::optix; } diff --git a/rtcore/CMakeLists.txt b/rtcore/CMakeLists.txt index 5cd77908..fe081ac5 100644 --- a/rtcore/CMakeLists.txt +++ b/rtcore/CMakeLists.txt @@ -7,7 +7,20 @@ set(CMAKE_CUDA_VISIBILITY_PRESET hidden) set(RTCORE_SOURCES ) - +# On AMD GPUs the cuda backend's .cu sources are compiled as HIP. Mark the +# given sources LANGUAGE HIP so hipcc handles the device code (no effect on the +# CUDA/OptiX build, where this macro is never invoked). +macro(rtc_mark_hip_sources) + foreach(src ${ARGN}) + get_filename_component(ext "${src}" EXT) + if (ext STREQUAL ".cu") + set_source_files_properties(${src} PROPERTIES + LANGUAGE HIP + COMPILE_OPTIONS "-fgpu-rdc") + endif() + endforeach() +endmacro() + macro(rtc_library_properties tgt) set_target_properties(${tgt} PROPERTIES @@ -147,7 +160,13 @@ else() endif() # ================================================================== -if (BARNEY_BACKEND_CUDA OR BARNEY_BACKEND_OPTIX) +if (BARNEY_BACKEND_CUDA OR BARNEY_BACKEND_OPTIX OR BARNEY_BACKEND_HIPRT) + if (BARNEY_HAVE_HIP) + rtc_mark_hip_sources( + cudaCommon/Device.cu + cudaCommon/Texture.cu + cudaCommon/TextureData.cu) + endif() set(BACKEND_CUDA_COMMON_SOURCES cudaCommon/Device.h cudaCommon/Device.cu @@ -216,6 +235,11 @@ endif() # ================================================================== if (BARNEY_BACKEND_CUDA) + if (BARNEY_HAVE_HIP) + rtc_mark_hip_sources( + cuda/cudaDevice.cu + cuda/cudaGroup.cu) + endif() set(BACKEND_CUDA_SOURCES cuda/Device.h cuda/cudaDevice.cu @@ -256,6 +280,47 @@ else() endif() +# ================================================================== +if (BARNEY_BACKEND_HIPRT) + message("Enabling HIPRT hardware-RT backend") + # the hiprt backend's sources are HIP TUs (they launch kernels and pull in + # HIPRT's inline device traversal); compile them with hipcc and -fgpu-rdc so + # the function-pointer SBT device-links exactly as the software backend. + set_source_files_properties( + hiprt/Device.cpp + hiprt/Group.cpp + hiprt/Buffer.cpp + hiprt/Geom.cpp + hiprt/GeomType.cpp + hiprt/TraceKernel.cpp + PROPERTIES LANGUAGE HIP COMPILE_OPTIONS "-fgpu-rdc") + add_library(barney_rtc_hiprt STATIC + hiprt/Device.h + hiprt/Device.cpp + hiprt/Buffer.h + hiprt/Buffer.cpp + hiprt/Geom.h + hiprt/Geom.cpp + hiprt/GeomType.h + hiprt/GeomType.cpp + hiprt/Group.h + hiprt/Group.cpp + hiprt/TraceKernel.h + hiprt/TraceKernel.cpp + ) + target_link_libraries(barney_rtc_hiprt PUBLIC + barney_rtc_config + barney_rtc_cuda_common + hiprt::hiprt + ) + target_compile_definitions(barney_rtc_hiprt + PUBLIC -DBARNEY_RTC_HIPRT=1) + rtc_library_properties(barney_rtc_hiprt) +else() + add_library(barney_rtc_hiprt INTERFACE) +endif() + + # ================================================================== # ================================================================== add_library(barney_rtc INTERFACE) @@ -268,13 +333,27 @@ endif() if (BARNEY_BACKEND_CUDA) target_link_libraries(barney_rtc INTERFACE barney_rtc_cuda) endif() +if (BARNEY_BACKEND_HIPRT) + target_link_libraries(barney_rtc INTERFACE barney_rtc_hiprt) +endif() if (COMMAND configure_cuda_source) # macro already defined else() macro(configure_cuda_source src) - if (NOT BARNEY_HAVE_CUDA) + if (BARNEY_HAVE_HIP) + # on AMD GPUs the same .cu/.dev.cu sources are compiled with hipcc; + # -fgpu-rdc gives relocatable device code so the function-pointer SBT + # dispatch resolves at device-link time (the analogue of nvcc -rdc=true). + get_filename_component(ext "${src}" EXT) + if (("${ext}" STREQUAL ".cu") OR ("${ext}" STREQUAL ".dev.cu")) + set_source_files_properties(${src} PROPERTIES + LANGUAGE HIP + COMPILE_OPTIONS "-fgpu-rdc" + ) + endif() + elseif (NOT BARNEY_HAVE_CUDA) get_filename_component(ext "${src}" EXT) if (("${ext}" STREQUAL ".cu") OR ("${ext}" STREQUAL ".dev.cu")) message("setting language of ${src} to CXX") diff --git a/rtcore/ComputeInterface.h b/rtcore/ComputeInterface.h index 723146fb..ce68275a 100644 --- a/rtcore/ComputeInterface.h +++ b/rtcore/ComputeInterface.h @@ -17,6 +17,11 @@ namespace rtc { } #endif +#if BARNEY_RTC_HIPRT +# include "hiprt/ComputeInterface.h" +namespace rtc { using namespace rtc::hiprt; } +#endif + #if BARNEY_RTC_OPTIX # include "optix/ComputeInterface.h" namespace rtc { using namespace rtc::optix; } diff --git a/rtcore/TraceInterface.h b/rtcore/TraceInterface.h index 15529ddf..2889b64d 100644 --- a/rtcore/TraceInterface.h +++ b/rtcore/TraceInterface.h @@ -21,3 +21,7 @@ # include "rtcore/cuda/TraceInterface.h" #endif +#if BARNEY_RTC_HIPRT +# include "rtcore/hiprt/TraceInterface.h" +#endif + diff --git a/rtcore/common/rtcore-common.h b/rtcore/common/rtcore-common.h index 5fbef011..30ba1308 100644 --- a/rtcore/common/rtcore-common.h +++ b/rtcore/common/rtcore-common.h @@ -14,7 +14,7 @@ #include #include "barney/barney.h" -# ifdef __CUDACC__ +# if defined(__CUDACC__) || defined(__HIPCC__) // # ifdef __CUDA_ARCH__ # define RTC_DEVICE_CODE 1 # endif diff --git a/rtcore/cuda/TraceInterface.h b/rtcore/cuda/TraceInterface.h index d7dbaa75..8483b888 100644 --- a/rtcore/cuda/TraceInterface.h +++ b/rtcore/cuda/TraceInterface.h @@ -252,6 +252,10 @@ namespace rtc { if (acceptedSBT && acceptedSBT->ch) { current = accepted; this->geomData = (acceptedSBT+1); + // Restore currentInstance for closestHit transform calls: + // leaveBlas() zeroed it, but closestHit needs it for + // object-to-world space transforms. + this->currentInstance = model->instanceRecords + accepted.instID; acceptedSBT->ch(*this); } } diff --git a/rtcore/cudaCommon/ComputeInterface.h b/rtcore/cudaCommon/ComputeInterface.h index 282535de..ac36fdc3 100644 --- a/rtcore/cudaCommon/ComputeInterface.h +++ b/rtcore/cudaCommon/ComputeInterface.h @@ -42,7 +42,7 @@ namespace rtc { struct ComputeInterface { // #if RTC_DEVICE_CODE -#ifdef __CUDACC__ +#if defined(__CUDACC__) || defined(__HIPCC__) inline __device__ vec3ui launchIndex() const { return getThreadIdx() + getBlockIdx() * getBlockDim(); @@ -64,7 +64,7 @@ namespace rtc { // INLINE IMPLEMENTATION // ================================================================== // #if RTC_DEVICE_CODE -#ifdef __CUDACC__ +#if defined(__CUDACC__) || defined(__HIPCC__) // ------------------------------------------------------------------ // cuda texturing // ------------------------------------------------------------------ diff --git a/rtcore/cudaCommon/cuda-common.h b/rtcore/cudaCommon/cuda-common.h index b5521bf9..2016bf0a 100644 --- a/rtcore/cudaCommon/cuda-common.h +++ b/rtcore/cudaCommon/cuda-common.h @@ -5,7 +5,7 @@ #pragma once #include "rtcore/common/rtcore-common.h" -#include +#include "rtcore/cudaCommon/cuda_to_hip.h" #ifdef __CUDACC__ # include # include diff --git a/rtcore/cudaCommon/cuda-helper.h b/rtcore/cudaCommon/cuda-helper.h index d916ffaf..acfef52b 100644 --- a/rtcore/cudaCommon/cuda-helper.h +++ b/rtcore/cudaCommon/cuda-helper.h @@ -8,7 +8,7 @@ #ifdef __GNUC__ # include #endif -#include "cuda_runtime.h" +#include "rtcore/cudaCommon/cuda_to_hip.h" #ifdef _WIN32 #include diff --git a/rtcore/cudaCommon/cuda_to_hip.h b/rtcore/cudaCommon/cuda_to_hip.h new file mode 100644 index 00000000..a53ae8be --- /dev/null +++ b/rtcore/cudaCommon/cuda_to_hip.h @@ -0,0 +1,105 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily +// +// Single compatibility shim that lets barney's CUDA software ray-tracing +// backend (rtcore/cuda + rtcore/cudaCommon) compile and run on AMD GPUs +// through HIP/ROCm. This is the only header in the backend that knows about +// HIP: it includes the HIP runtime and aliases the small set of CUDA runtime +// spellings the backend actually uses to their hipXxx equivalents, so the +// CUDA sources stay in CUDA spelling and the CUDA/OptiX build is byte +// identical (this whole header is inert unless compiled by hipcc). + +#pragma once + +#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP) || defined(__HIPCC__) + +// Pull in the host libc memory routines before the HIP runtime so that +// host-side memcpy/memset resolve to the C library, not a device overload +// the HIP runtime may bring into scope. +#include +#include +#include + +// ------------------------------------------------------------------ +// runtime error type + status +// ------------------------------------------------------------------ +using cudaError_t = hipError_t; +#define cudaSuccess hipSuccess +#define cudaErrorPeerAccessAlreadyEnabled hipErrorPeerAccessAlreadyEnabled +#define cudaGetErrorString hipGetErrorString +#define cudaGetLastError hipGetLastError + +// ------------------------------------------------------------------ +// device / stream management (used by BARNEY_CUDA_CALL(cuda##X) sites +// via token paste, plus a handful of bare spellings) +// ------------------------------------------------------------------ +#define cudaSetDevice hipSetDevice +#define cudaGetDevice hipGetDevice +#define cudaGetDeviceCount hipGetDeviceCount +#define cudaGetDeviceProperties hipGetDeviceProperties +#define cudaDeviceProp hipDeviceProp_t +#define cudaDeviceSynchronize hipDeviceSynchronize +#define cudaDeviceCanAccessPeer hipDeviceCanAccessPeer +#define cudaDeviceEnablePeerAccess hipDeviceEnablePeerAccess + +using cudaStream_t = hipStream_t; +#define cudaStreamNonBlocking hipStreamNonBlocking +#define cudaStreamCreate hipStreamCreate +#define cudaStreamCreateWithFlags hipStreamCreateWithFlags +#define cudaStreamSynchronize hipStreamSynchronize +#define cudaStreamDestroy hipStreamDestroy + +// ------------------------------------------------------------------ +// memory +// ------------------------------------------------------------------ +#define cudaMalloc hipMalloc +#define cudaMallocManaged hipMallocManaged +#define cudaMallocHost hipHostMalloc +#define cudaFree hipFree +#define cudaFreeHost hipHostFree +#define cudaMemcpy hipMemcpy +#define cudaMemcpyAsync hipMemcpyAsync +#define cudaMemsetAsync hipMemsetAsync +#define cudaMemcpyDefault hipMemcpyDefault +#define cudaMemcpyHostToDevice hipMemcpyHostToDevice +#define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost + +// ------------------------------------------------------------------ +// textures and arrays (cudaArray-backed, no pitched 2D binds) +// ------------------------------------------------------------------ +using cudaTextureObject_t = hipTextureObject_t; +using cudaArray_t = hipArray_t; +using cudaChannelFormatDesc = hipChannelFormatDesc; +using cudaResourceDesc = hipResourceDesc; +using cudaTextureDesc = hipTextureDesc; +using cudaExtent = hipExtent; +using cudaMemcpy3DParms = hipMemcpy3DParms; +using cudaTextureFilterMode = hipTextureFilterMode; +using cudaTextureAddressMode = hipTextureAddressMode; +using cudaTextureReadMode = hipTextureReadMode; + +#define cudaCreateChannelDesc hipCreateChannelDesc +#define cudaCreateTextureObject hipCreateTextureObject +#define cudaDestroyTextureObject hipDestroyTextureObject +#define cudaMallocArray hipMallocArray +#define cudaMalloc3DArray hipMalloc3DArray +#define cudaFreeArray hipFreeArray +#define cudaMemcpy2DToArray hipMemcpy2DToArray +#define cudaMemcpy3D hipMemcpy3D +#define make_cudaPitchedPtr make_hipPitchedPtr + +#define cudaResourceTypeArray hipResourceTypeArray +#define cudaReadModeElementType hipReadModeElementType +#define cudaReadModeNormalizedFloat hipReadModeNormalizedFloat +#define cudaFilterModePoint hipFilterModePoint +#define cudaFilterModeLinear hipFilterModeLinear +#define cudaAddressModeWrap hipAddressModeWrap +#define cudaAddressModeClamp hipAddressModeClamp +#define cudaAddressModeBorder hipAddressModeBorder +#define cudaAddressModeMirror hipAddressModeMirror + +#else +# include +#endif diff --git a/rtcore/hiprt/AppInterface.h b/rtcore/hiprt/AppInterface.h new file mode 100644 index 00000000..d52884a3 --- /dev/null +++ b/rtcore/hiprt/AppInterface.h @@ -0,0 +1,37 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/hiprt/Device.h" +#include "rtcore/cuda/Buffer.h" +#include "rtcore/cudaCommon/ComputeKernel.h" +#include "rtcore/hiprt/TraceKernel.h" +#include "rtcore/cuda/Geom.h" +#include "rtcore/hiprt/Group.h" +#include "rtcore/cudaCommon/TextureData.h" +#include "rtcore/cudaCommon/Texture.h" + +namespace rtc { + namespace hiprt { + + using rtc::cuda_common::enablePeerAccess; + using rtc::cuda_common::getPhysicalDeviceHash; + + using rtc::cuda_common::ComputeKernel1D; + using rtc::cuda_common::ComputeKernel2D; + using rtc::cuda_common::ComputeKernel3D; + + using rtc::cuda_common::Texture; + using rtc::cuda_common::TextureData; + } +} + +#define RTC_IMPORT_USER_GEOM(moduleName,typeName,DD,has_ah,has_ch) \ + extern ::rtc::hiprt::GeomType *createGeomType_##typeName(::rtc::Device *); + +#define RTC_IMPORT_TRIANGLES_GEOM(moduleName,typeName,DD,has_ah,has_ch) \ + extern rtc::hiprt::GeomType *createGeomType_##typeName(rtc::Device *); diff --git a/rtcore/hiprt/Buffer.cpp b/rtcore/hiprt/Buffer.cpp new file mode 100644 index 00000000..91663859 --- /dev/null +++ b/rtcore/hiprt/Buffer.cpp @@ -0,0 +1,48 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#include "rtcore/hiprt/Buffer.h" + +namespace rtc { + namespace hiprt { + + using cuda_common::SetActiveGPU; + + Buffer::Buffer(cuda_common::Device *device, + size_t numBytes, + const void *initValues) + : device(device) + { + if (numBytes == 0) return; + SetActiveGPU forDuration(device); + BARNEY_CUDA_CALL(Malloc((void**)&d_data,numBytes)); + if (initValues) + BARNEY_CUDA_CALL(Memcpy(d_data,initValues,numBytes,cudaMemcpyDefault)); + } + + Buffer::~Buffer() + { + if (!d_data) return; + BARNEY_CUDA_CALL_NOTHROW(Free(d_data)); + } + + void Buffer::upload(const void *data, size_t numBytes, size_t offset) + { + if (!d_data) return; + SetActiveGPU forDuration(device); + BARNEY_CUDA_CALL(Memcpy(((char *)d_data)+offset,data,numBytes,cudaMemcpyDefault)); + } + + void Buffer::resize(size_t numBytes) + { + SetActiveGPU forDuration(device); + if (d_data) + BARNEY_CUDA_CALL(Free(d_data)); + BARNEY_CUDA_CALL(Malloc((void**)&d_data,numBytes)); + } + + } +} diff --git a/rtcore/hiprt/Buffer.h b/rtcore/hiprt/Buffer.h new file mode 100644 index 00000000..8ca24012 --- /dev/null +++ b/rtcore/hiprt/Buffer.h @@ -0,0 +1,31 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/cudaCommon/Device.h" + +namespace rtc { + namespace hiprt { + + // a plain device buffer; identical to the cuda backend's Buffer but bound to + // cuda_common::Device so the hiprt and cuda backends do not share a Device + // type (each backend owns its rtc::Device). + struct Buffer { + Buffer(cuda_common::Device *device, + size_t numBytes, + const void *initValues); + virtual ~Buffer(); + + void *getDD() const { return d_data; } + void upload(const void *data, size_t numBytes, size_t offset=0); + void resize(size_t newNumBytes); + void *d_data = 0; + cuda_common::Device *const device; + }; + + } +} diff --git a/rtcore/hiprt/ComputeInterface.h b/rtcore/hiprt/ComputeInterface.h new file mode 100644 index 00000000..bc232586 --- /dev/null +++ b/rtcore/hiprt/ComputeInterface.h @@ -0,0 +1,26 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/cudaCommon/ComputeInterface.h" + +namespace rtc { + namespace hiprt { +#if RTC_DEVICE_CODE + using cuda_common::ComputeInterface; + using cuda_common::tex1D; + using cuda_common::tex2D; + using cuda_common::tex3D; + using cuda_common::fatomicMin; + using cuda_common::fatomicMax; +#endif + } +} + +# define __rtc_global __global__ +# define __rtc_launch(myRTC,kernel,nb,bs,...) \ + { rtc::hiprt::SetActiveGPU forDuration(myRTC); if (nb) kernel<<stream>>>(rtc::hiprt::ComputeInterface(), __VA_ARGS__); } diff --git a/rtcore/hiprt/Device.cpp b/rtcore/hiprt/Device.cpp new file mode 100644 index 00000000..96e495f7 --- /dev/null +++ b/rtcore/hiprt/Device.cpp @@ -0,0 +1,96 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#include "rtcore/hiprt/Device.h" +#include "rtcore/hiprt/Group.h" +#include "rtcore/hiprt/TraceKernel.h" +#include "rtcore/hiprt/Geom.h" +#include "rtcore/hiprt/GeomType.h" +#include "rtcore/hiprt/Buffer.h" + +#include +#include + +namespace rtc { + namespace hiprt { + + static void hiprtCheck(hiprtError e, const char *where) + { + if (e != hiprtSuccess) + throw std::runtime_error(std::string("HIPRT error in ")+where + +" code "+std::to_string((int)e)); + } +#define HIPRT_CALL(call,where) rtc::hiprt::hiprtCheck(call,where) + + Device::Device(int physicalGPU) + : cuda_common::Device(physicalGPU) + { + SetActiveGPU forDuration(this); + + // Bind HIPRT to THIS device's HIP context/device so every BVH build and + // the trace kernels see the same device pointers barney allocates (this + // avoids a "fresh context => degenerate BVH, zero hits" failure mode). These + // are HIP driver-API calls used only on the HIP build, so they are spelled + // directly rather than through the cuda_to_hip compat aliases. + hipDevice_t hdev = 0; + (void)hipDeviceGet(&hdev, physicalID); + hipCtx_t hctx = nullptr; + if (hipCtxGetCurrent(&hctx) != hipSuccess || hctx == nullptr) { + (void)hipDevicePrimaryCtxRetain(&hctx, hdev); + (void)hipCtxSetCurrent(hctx); + } + + hiprtContextCreationInput ci{}; + ci.deviceType = hiprtDeviceAMD; + ci.ctxt = hctx; + ci.device = hdev; + HIPRT_CALL(hiprtCreateContext(HIPRT_API_VERSION, ci, hiprtCtx), + "hiprtCreateContext"); + } + + Device::~Device() + { + if (hiprtCtx) + hiprtDestroyContext(hiprtCtx); + } + + rtc::AccelHandle getAccelHandle(Group *ig) + { return ig->getDD(); } + + void Device::freeGroup(Group *g) + { delete g; } + + Denoiser *Device::createDenoiser() + { return nullptr; } + + Buffer *Device::createBuffer(size_t numBytes, const void *initValues) + { return new Buffer(this,numBytes,initValues); } + + void Device::freeBuffer(Buffer *b) + { delete b; } + + void Device::freeGeomType(GeomType *gt) + { delete gt; } + + void Device::freeGeom(Geom *g) + { delete g; } + + Group *Device::createTrianglesGroup(const std::vector &geoms) + { return new TrianglesGeomGroup(this,geoms); } + + Group *Device::createUserGeomsGroup(const std::vector &geoms) + { return new UserGeomGroup(this,geoms); } + + Group *Device::createInstanceGroup(const std::vector &groups, + const std::vector &instIDs, + const std::vector &xfms) + { return new InstanceGroup(this,groups,instIDs,xfms); } + + void Device::buildPipeline() {} + void Device::buildSBT() {} + + } +} diff --git a/rtcore/hiprt/Device.h b/rtcore/hiprt/Device.h new file mode 100644 index 00000000..4b2f5fea --- /dev/null +++ b/rtcore/hiprt/Device.h @@ -0,0 +1,97 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/cudaCommon/Device.h" + +// host-side HIPRT API; the device-side traversal lives in TraceInterface.h +#include + +namespace rtc { + namespace hiprt { + + using cuda_common::SetActiveGPU; + + // forward decls; the geometry/SBT/program data model is backend-generic + // compute (no OptiX, no HIPRT) and lives in this backend's Geom/GeomType/ + // Buffer (cuda_common::Device-based). Only BVH build + ray traversal differ + // (HIPRT here, the cuBQL software walk in the cuda backend), so only Group + // and the device-side TraceInterface are HIPRT-specific. + struct Buffer; + struct Geom; + struct GeomType; + struct TrianglesGeom; + struct UserGeom; + struct UserGeomType; + struct TrianglesGeomType; + struct Group; + struct Denoiser; + struct TraceKernel2D; + + using rtc::cuda_common::Texture; + using rtc::cuda_common::TextureData; + + using cuda_common::float2; + using cuda_common::float3; + using cuda_common::float4; + using cuda_common::int2; + using cuda_common::int3; + using cuda_common::int4; + using cuda_common::load; + using cuda_common::TextureObject; + + struct Device : public cuda_common::Device { + Device(int physicalGPU); + virtual ~Device(); + + std::string toString() const + { return "rtc::hiprt::Device(physical="+std::to_string(physicalID)+")"; } + + /*! the HIPRT context, created over this Device's HIP device + stream; + shared by every Group's BVH build and the trace kernels so they see + the same device pointers (binding to this Device's HIP context avoids + a "fresh context => degenerate BVH" failure mode). */ + hiprtContext hiprtCtx = nullptr; + + void freeGeomType(GeomType *); + void freeGeom(Geom *); + + Group *createTrianglesGroup(const std::vector &geoms); + Group *createUserGeomsGroup(const std::vector &geoms); + Group *createInstanceGroup(const std::vector &groups, + const std::vector &instIDs, + const std::vector &xfms); + void freeGroup(Group *); + void buildPipeline(); + void buildSBT(); + Buffer *createBuffer(size_t numBytes, const void *initValues = 0); + void freeBuffer(Buffer *); + Denoiser *createDenoiser(); + }; + + /*! HIPRT has no denoiser; mirror the cuda backend's no-op passthrough (the + renderer runs un-denoised, exactly as on the software backend). OIDN-GPU + on ROCm is a deferred enhancement. */ + struct Denoiser + { + Denoiser(Device* device) : device(device) {} + ~Denoiser() = default; + void resize(vec2i dims) { outputDims = dims; } + void run(float blendFactor) {} + Device* const device; + + vec4f *in_rgba = 0; + vec4f *out_rgba = 0; + vec3f *in_normal = 0; + + bool upscaleMode = false; + vec2i outputDims = {0,0}; + }; + + rtc::AccelHandle getAccelHandle(Group *ig); + } +} diff --git a/rtcore/hiprt/Geom.cpp b/rtcore/hiprt/Geom.cpp new file mode 100644 index 00000000..29dd8457 --- /dev/null +++ b/rtcore/hiprt/Geom.cpp @@ -0,0 +1,36 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#include "rtcore/hiprt/Geom.h" +#include "rtcore/hiprt/GeomType.h" + +namespace rtc { + namespace hiprt { + + Geom::Geom(GeomType *gt) + : gt(gt), data(gt->sizeOfDD) + {} + + Geom::~Geom() {} + + void Geom::setDD(const void *dd) + { memcpy(data.data(),dd,data.size()); } + + TrianglesGeom::TrianglesGeom(GeomType *gt) : Geom(gt) {} + + void TrianglesGeom::setVertices(Buffer *vertices, int numVertices) + { this->vertices = vertices; this->numVertices = numVertices; } + + void TrianglesGeom::setIndices(Buffer *indices, int numIndices) + { this->indices = indices; this->numIndices = numIndices; } + + UserGeom::UserGeom(GeomType *gt) : Geom(gt) {} + + void UserGeom::setPrimCount(int primCount) + { this->primCount = primCount; } + + } +} diff --git a/rtcore/hiprt/Geom.h b/rtcore/hiprt/Geom.h new file mode 100644 index 00000000..63e6e9c0 --- /dev/null +++ b/rtcore/hiprt/Geom.h @@ -0,0 +1,68 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/hiprt/Buffer.h" +#include "rtcore/hiprt/GeomType.h" + +namespace rtc { + namespace hiprt { + + struct GeomType; + + struct Geom { + Geom(GeomType *gt); + virtual ~Geom(); + void setDD(const void *dd); + + // the per-geom shading record: the same function-pointer SBT the software + // backend uses. HIPRT yields the hit; barney's megakernel calls these. + struct SBTHeader { + AHProg ah; + CHProg ch; + union { + struct { + const vec3f *vertices; + const vec3i *indices; + } triangles; + struct { + IntersectProg intersect; + } user; + }; + }; + + virtual void setPrimCount(int primCount) = 0; + virtual void setVertices(Buffer *vertices, int numVertices) = 0; + virtual void setIndices(Buffer *indices, int numIndices) = 0; + + GeomType *const gt; + std::vector data; + }; + + struct TrianglesGeom : public Geom { + TrianglesGeom(GeomType *gt); + void setPrimCount(int primCount) override { assert(0); } + void setVertices(Buffer *vertices, int numVertices) override; + void setIndices(Buffer *indices, int numIndices) override; + + Buffer *vertices = 0; + int numVertices = 0; + Buffer *indices = 0; + int numIndices = 0; + }; + + struct UserGeom : public Geom { + UserGeom(GeomType *gt); + void setPrimCount(int primCount) override; + void setVertices(Buffer *vertices, int numVertices) override { assert(0); } + void setIndices(Buffer *indices, int numIndices) override { assert(0); } + + int primCount = 0; + }; + + } +} diff --git a/rtcore/hiprt/GeomType.cpp b/rtcore/hiprt/GeomType.cpp new file mode 100644 index 00000000..1a5c7b00 --- /dev/null +++ b/rtcore/hiprt/GeomType.cpp @@ -0,0 +1,38 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#include "rtcore/hiprt/GeomType.h" +#include "rtcore/hiprt/Geom.h" + +namespace rtc { + namespace hiprt { + + GeomType::GeomType(cuda_common::Device *device, size_t sizeOfDD) + : device(device), sizeOfDD(sizeOfDD) + {} + + UserGeomType::UserGeomType(cuda_common::Device *device, + size_t sizeOfDD, + BoundsKernel bounds, + IntersectProg intersect, + AHProg ah, + CHProg ch) + : GeomType(device,sizeOfDD), + bounds(bounds), intersect(intersect), ah(ah), ch(ch) + {} + + TrianglesGeomType::TrianglesGeomType(cuda_common::Device *device, + size_t sizeOfDD, + AHProg ah, + CHProg ch) + : GeomType(device,sizeOfDD), ah(ah), ch(ch) + {} + + Geom *UserGeomType::createGeom() { return new UserGeom(this); } + Geom *TrianglesGeomType::createGeom() { return new TrianglesGeom(this); } + + } +} diff --git a/rtcore/hiprt/GeomType.h b/rtcore/hiprt/GeomType.h new file mode 100644 index 00000000..f16d09f3 --- /dev/null +++ b/rtcore/hiprt/GeomType.h @@ -0,0 +1,64 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/cudaCommon/Device.h" + +namespace rtc { + namespace hiprt { + + struct Geom; + struct Device; + struct TraceInterface; + + typedef void (*BoundsKernel)(Device *device, + const void *dd, + box3f *boundsArray, + int numPrims); + typedef void (*IntersectProg)(TraceInterface &ti); + typedef void (*AHProg)(TraceInterface &ti); + typedef void (*CHProg)(TraceInterface &ti); + + struct GeomType { + GeomType(cuda_common::Device *device, size_t sizeOfDD); + virtual ~GeomType() = default; + virtual Geom *createGeom() = 0; + + cuda_common::Device *const device; + size_t const sizeOfDD; + }; + + struct UserGeomType : public GeomType { + UserGeomType(cuda_common::Device *device, + size_t sizeOfDD, + BoundsKernel bounds, + IntersectProg intersect, + AHProg ah, + CHProg ch); + virtual ~UserGeomType() = default; + Geom *createGeom() override; + + BoundsKernel const bounds; + IntersectProg const intersect; + AHProg const ah; + CHProg const ch; + }; + + struct TrianglesGeomType : public GeomType { + TrianglesGeomType(cuda_common::Device *device, + size_t sizeOfDD, + AHProg ah, + CHProg ch); + virtual ~TrianglesGeomType() = default; + Geom *createGeom() override; + + AHProg const ah; + CHProg const ch; + }; + + } +} diff --git a/rtcore/hiprt/Group.cpp b/rtcore/hiprt/Group.cpp new file mode 100644 index 00000000..726fd6dc --- /dev/null +++ b/rtcore/hiprt/Group.cpp @@ -0,0 +1,357 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily +// +// HIPRT acceleration-structure build for the hardware-RT backend. Each geom +// group becomes one HIPRT geometry (BLAS): a hiprtTriangleMeshPrimitive for +// triangle geoms, or a hiprtAABBListPrimitive (per-prim AABBs + a func-table +// intersect thunk) for user geoms. The instance group becomes a hiprtScene +// (TLAS). We keep barney's per-prim (geomID, primID) remap and per-geom SBT so +// the device-side shading dispatch is identical to the software backend; HIPRT +// supplies only the BVH and the traversal that yields the geometry-local primID. + +#include "rtcore/hiprt/Group.h" +#include "rtcore/hiprt/Geom.h" +#include "rtcore/hiprt/GeomType.h" +#include "rtcore/hiprt/Buffer.h" + +#include +#include +#include + +namespace rtc { + namespace hiprt { + + static void hiprtCheck(hiprtError e, const char *where) + { + if (e != hiprtSuccess) + throw std::runtime_error(std::string("HIPRT error in ")+where + +" code "+std::to_string((int)e)); + } +#define HC(call) rtc::hiprt::hiprtCheck(call,#call) + + // ------------------------------------------------------------------ + Group::Group(Device *device) : device(device) {} + Group::~Group() {} + + // ------------------------------------------------------------------ + GeomGroup::GeomGroup(Device *device, const std::vector &geoms) + : Group(device), geoms(geoms) + {} + + GeomGroup::~GeomGroup() + { + SetActiveGPU forDuration(device); + if (geom) hiprtDestroyGeometry(device->hiprtCtx, geom); + if (d_buildTmp) BARNEY_CUDA_CALL_NOTHROW(Free(d_buildTmp)); + if (sbt) BARNEY_CUDA_CALL_NOTHROW(Free(sbt)); + if (prims) BARNEY_CUDA_CALL_NOTHROW(Free(prims)); + if (d_vertices) BARNEY_CUDA_CALL_NOTHROW(Free(d_vertices)); + if (d_indices) BARNEY_CUDA_CALL_NOTHROW(Free(d_indices)); + if (d_aabbs) BARNEY_CUDA_CALL_NOTHROW(Free(d_aabbs)); + } + + // assemble the per-geom SBT (shared with the software backend's layout) and + // the (geomID, primID) prim remap, uploaded to `sbt`/`prims`. Returns the + // total prim count. + int GeomGroup_buildSBTandPrims(GeomGroup *gg, bool isTriangles) + { + Device *device = gg->device; + gg->sbtEntrySize = 0; + size_t align = sizeof(float4); + for (size_t i=0;igeoms.size();i++) { + gg->sbtEntrySize = std::max(gg->sbtEntrySize, gg->geoms[i]->gt->sizeOfDD); + gg->sbtEntrySize = align*((gg->sbtEntrySize+align-1)/align); + } + gg->sbtEntrySize += sizeof(Geom::SBTHeader); + + std::vector hostSBT(gg->geoms.size()*gg->sbtEntrySize); + uint8_t *p = hostSBT.data(); + for (size_t i=0;igeoms.size();i++) { + Geom *geom = gg->geoms[i]; + Geom::SBTHeader *header = (Geom::SBTHeader *)p; + if (isTriangles) { + TrianglesGeom *tg = (TrianglesGeom *)geom; + header->ah = ((TrianglesGeomType*)tg->gt)->ah; + header->ch = ((TrianglesGeomType*)tg->gt)->ch; + header->triangles.vertices = (const vec3f*)tg->vertices->getDD(); + header->triangles.indices = (const vec3i*)tg->indices->getDD(); + } else { + UserGeom *ug = (UserGeom *)geom; + header->ah = ((UserGeomType*)ug->gt)->ah; + header->ch = ((UserGeomType*)ug->gt)->ch; + header->user.intersect = ((UserGeomType*)ug->gt)->intersect; + } + memcpy(p+sizeof(Geom::SBTHeader),geom->data.data(),geom->data.size()); + p += gg->sbtEntrySize; + } + if (gg->sbt) BARNEY_CUDA_CALL(Free(gg->sbt)); + BARNEY_CUDA_CALL(Malloc((void**)&gg->sbt,hostSBT.size())); + BARNEY_CUDA_CALL(Memcpy(gg->sbt,hostSBT.data(),hostSBT.size(),cudaMemcpyDefault)); + + // count prims and build the (geomID, primID) remap, in geom order (which + // is exactly the order we concatenate triangles / AABBs into the BLAS, so + // HIPRT's geometry-local primID indexes straight into prims[]). + std::vector hostPrims; + for (size_t i=0;igeoms.size();i++) { + int count = isTriangles + ? ((TrianglesGeom*)gg->geoms[i])->numIndices + : ((UserGeom*)gg->geoms[i])->primCount; + for (int j=0;jnumPrims = (int)hostPrims.size(); + if (gg->prims) BARNEY_CUDA_CALL(Free(gg->prims)); + if (gg->numPrims) { + BARNEY_CUDA_CALL(Malloc((void**)&gg->prims, + gg->numPrims*sizeof(GeomGroup::Prim))); + BARNEY_CUDA_CALL(Memcpy(gg->prims,hostPrims.data(), + gg->numPrims*sizeof(GeomGroup::Prim), + cudaMemcpyDefault)); + } + return gg->numPrims; + } + + static void buildHiprtGeometry(GeomGroup *gg, hiprtGeometryBuildInput &bi) + { + Device *device = gg->device; + hiprtBuildOptions bo{}; + bo.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + size_t tempSize = 0; + HC(hiprtGetGeometryBuildTemporaryBufferSize(device->hiprtCtx,bi,bo,tempSize)); + if (tempSize > gg->buildTmpSize) { + if (gg->d_buildTmp) BARNEY_CUDA_CALL(Free(gg->d_buildTmp)); + gg->d_buildTmp = nullptr; + if (tempSize) BARNEY_CUDA_CALL(Malloc(&gg->d_buildTmp,tempSize)); + gg->buildTmpSize = tempSize; + } + if (gg->geom) { hiprtDestroyGeometry(device->hiprtCtx,gg->geom); gg->geom=nullptr; } + HC(hiprtCreateGeometry(device->hiprtCtx,bi,bo,gg->geom)); + HC(hiprtBuildGeometry(device->hiprtCtx,hiprtBuildOperationBuild,bi,bo, + gg->d_buildTmp,device->stream,gg->geom)); + device->sync(); + } + + // ------------------------------------------------------------------ + TrianglesGeomGroup::TrianglesGeomGroup(Device *device, + const std::vector &geoms) + : GeomGroup(device,geoms) + {} + + void TrianglesGeomGroup::buildAccel() + { + SetActiveGPU forDuration(device); + GeomGroup_buildSBTandPrims(this,/*isTriangles*/true); + + // concatenate all geoms' triangles into one vertex+index buffer; HIPRT + // primID then indexes straight into prims[]. + std::vector hostVerts; + std::vector hostIdx; + for (size_t i=0;i v(tg->numVertices); + std::vector idx(tg->numIndices); + BARNEY_CUDA_CALL(Memcpy(v.data(),tg->vertices->getDD(), + tg->numVertices*sizeof(vec3f),cudaMemcpyDefault)); + BARNEY_CUDA_CALL(Memcpy(idx.data(),tg->indices->getDD(), + tg->numIndices*sizeof(vec3i),cudaMemcpyDefault)); + for (auto &vv : v) hostVerts.push_back(vv); + for (auto t : idx) hostIdx.push_back(vec3i{t.x+base,t.y+base,t.z+base}); + } + + if (d_vertices) { BARNEY_CUDA_CALL(Free(d_vertices)); d_vertices=nullptr; } + if (d_indices) { BARNEY_CUDA_CALL(Free(d_indices)); d_indices=nullptr; } + BARNEY_CUDA_CALL(Malloc(&d_vertices,hostVerts.size()*sizeof(vec3f))); + BARNEY_CUDA_CALL(Memcpy(d_vertices,hostVerts.data(), + hostVerts.size()*sizeof(vec3f),cudaMemcpyDefault)); + BARNEY_CUDA_CALL(Malloc(&d_indices,hostIdx.size()*sizeof(vec3i))); + BARNEY_CUDA_CALL(Memcpy(d_indices,hostIdx.data(), + hostIdx.size()*sizeof(vec3i),cudaMemcpyDefault)); + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = d_vertices; + mesh.vertexCount = (uint32_t)hostVerts.size(); + mesh.vertexStride = sizeof(vec3f); + mesh.triangleIndices= d_indices; + mesh.triangleCount = (uint32_t)hostIdx.size(); + mesh.triangleStride = sizeof(vec3i); + + hiprtGeometryBuildInput bi{}; + bi.type = hiprtPrimitiveTypeTriangleMesh; + bi.primitive.triangleMesh = mesh; + bi.geomType = 0; + buildHiprtGeometry(this,bi); + } + + GeomGroup::DeviceRecord TrianglesGeomGroup::getRecord() + { + DeviceRecord r; + r.sbt = sbt; + r.prims = prims; + r.sbtEntrySize = (uint32_t)sbtEntrySize; + r.isTrianglesGroup = 1; + return r; + } + + // ------------------------------------------------------------------ + UserGeomGroup::UserGeomGroup(Device *device, const std::vector &geoms) + : GeomGroup(device,geoms) + {} + + void UserGeomGroup::buildAccel() + { + SetActiveGPU forDuration(device); + GeomGroup_buildSBTandPrims(this,/*isTriangles*/false); + + // per-prim AABBs via each geom type's bounds kernel, concatenated in geom + // order; HIPRT builds an AABB-list BLAS over them and calls the func-table + // intersect thunk (see TraceKernel.cpp) per candidate prim. + box3f *d_bounds = nullptr; + if (numPrims) BARNEY_CUDA_CALL(Malloc((void**)&d_bounds,numPrims*sizeof(box3f))); + size_t ofs = 0; + for (size_t i=0;iprimCount; + if (!count) continue; + void *mySBT = sbt + i*sbtEntrySize; + void *geomData = ((Geom::SBTHeader *)mySBT)+1; + ((UserGeomType*)geom->gt)->bounds(device,geomData, + d_bounds+ofs,count); + ofs += count; + } + device->sync(); + + // HIPRT's AABB list wants packed (lower.xyz, upper.xyz) float pairs; owl's + // box3f is exactly {vec3f lower; vec3f upper;} so the layout already + // matches a hiprtFloat4-pair-free tight 6-float AABB. + if (d_aabbs) { BARNEY_CUDA_CALL(Free(d_aabbs)); d_aabbs=nullptr; } + d_aabbs = d_bounds; // box3f == 6 contiguous floats per prim + + hiprtAABBListPrimitive list{}; + list.aabbCount = (uint32_t)numPrims; + list.aabbStride = sizeof(box3f); + list.aabbs = d_aabbs; + + hiprtGeometryBuildInput bi{}; + bi.type = hiprtPrimitiveTypeAABBList; + bi.primitive.aabbList = list; + bi.geomType = 0; + buildHiprtGeometry(this,bi); + } + + GeomGroup::DeviceRecord UserGeomGroup::getRecord() + { + DeviceRecord r; + r.sbt = sbt; + r.prims = prims; + r.sbtEntrySize = (uint32_t)sbtEntrySize; + r.isTrianglesGroup = 0; + return r; + } + + // ------------------------------------------------------------------ + InstanceGroup::InstanceGroup(Device *device, + const std::vector &groups, + const std::vector &instanceIDs, + const std::vector &xfms) + : Group(device), groups(groups), instanceIDs(instanceIDs), xfms(xfms) + {} + + InstanceGroup::~InstanceGroup() + { + SetActiveGPU forDuration(device); + if (scene) hiprtDestroyScene(device->hiprtCtx,scene); + if (d_sceneTmp) BARNEY_CUDA_CALL_NOTHROW(Free(d_sceneTmp)); + if (d_instances) BARNEY_CUDA_CALL_NOTHROW(Free(d_instances)); + if (d_frames) BARNEY_CUDA_CALL_NOTHROW(Free(d_frames)); + if (d_instanceRecords) BARNEY_CUDA_CALL_NOTHROW(Free(d_instanceRecords)); + if (d_deviceRecord) BARNEY_CUDA_CALL_NOTHROW(Free(d_deviceRecord)); + } + + void InstanceGroup::setTransforms(const std::vector &newXfms) + { xfms = newXfms; } + + void InstanceGroup::buildAccel() + { + SetActiveGPU forDuration(device); + int numInstances = (int)groups.size(); + + // per-instance shading records (transforms + the group's SBT/prim record). + std::vector hostRecs(numInstances); + std::vector hostInst(numInstances); + std::vector hostFrames(numInstances); + for (int i=0;igetRecord(); + hostRecs[i].ID = instanceIDs.empty()? (uint32_t)i : (uint32_t)instanceIDs[i]; + + hostInst[i].type = hiprtInstanceTypeGeometry; + hostInst[i].geometry = gg->getGeometry(); + + // owl affine3f: linear l (vx,vy,vz columns) + translation p. HIPRT's + // hiprtFrameMatrix is a 3x4 row-major object->world transform. + const affine3f &m = xfms[i]; + hostFrames[i].matrix[0][0]=m.l.vx.x; hostFrames[i].matrix[0][1]=m.l.vy.x; + hostFrames[i].matrix[0][2]=m.l.vz.x; hostFrames[i].matrix[0][3]=m.p.x; + hostFrames[i].matrix[1][0]=m.l.vx.y; hostFrames[i].matrix[1][1]=m.l.vy.y; + hostFrames[i].matrix[1][2]=m.l.vz.y; hostFrames[i].matrix[1][3]=m.p.y; + hostFrames[i].matrix[2][0]=m.l.vx.z; hostFrames[i].matrix[2][1]=m.l.vy.z; + hostFrames[i].matrix[2][2]=m.l.vz.z; hostFrames[i].matrix[2][3]=m.p.z; + } + + if (d_instanceRecords) { BARNEY_CUDA_CALL(Free(d_instanceRecords)); d_instanceRecords=nullptr; } + if (numInstances) { + BARNEY_CUDA_CALL(Malloc((void**)&d_instanceRecords,numInstances*sizeof(InstanceRecord))); + BARNEY_CUDA_CALL(Memcpy(d_instanceRecords,hostRecs.data(), + numInstances*sizeof(InstanceRecord),cudaMemcpyDefault)); + } + if (d_instances) { BARNEY_CUDA_CALL(Free(d_instances)); d_instances=nullptr; } + if (d_frames) { BARNEY_CUDA_CALL(Free(d_frames)); d_frames=nullptr; } + BARNEY_CUDA_CALL(Malloc(&d_instances,numInstances*sizeof(hiprtInstance))); + BARNEY_CUDA_CALL(Memcpy(d_instances,hostInst.data(), + numInstances*sizeof(hiprtInstance),cudaMemcpyDefault)); + BARNEY_CUDA_CALL(Malloc(&d_frames,numInstances*sizeof(hiprtFrameMatrix))); + BARNEY_CUDA_CALL(Memcpy(d_frames,hostFrames.data(), + numInstances*sizeof(hiprtFrameMatrix),cudaMemcpyDefault)); + + hiprtSceneBuildInput si{}; + si.instanceCount = (uint32_t)numInstances; + si.instances = d_instances; + si.instanceTransformHeaders = nullptr; // one frame per instance + si.instanceFrames = d_frames; + si.frameType = hiprtFrameTypeMatrix; + + hiprtBuildOptions bo{}; + bo.buildFlags = hiprtBuildFlagBitPreferFastBuild; + size_t tempSize = 0; + HC(hiprtGetSceneBuildTemporaryBufferSize(device->hiprtCtx,si,bo,tempSize)); + if (tempSize > sceneTmpSize) { + if (d_sceneTmp) BARNEY_CUDA_CALL(Free(d_sceneTmp)); + d_sceneTmp = nullptr; + if (tempSize) BARNEY_CUDA_CALL(Malloc(&d_sceneTmp,tempSize)); + sceneTmpSize = tempSize; + } + if (scene) { hiprtDestroyScene(device->hiprtCtx,scene); scene=nullptr; } + HC(hiprtCreateScene(device->hiprtCtx,si,bo,scene)); + HC(hiprtBuildScene(device->hiprtCtx,hiprtBuildOperationBuild,si,bo, + d_sceneTmp,device->stream,scene)); + device->sync(); + + DeviceRecord dd; + dd.scene = scene; + dd.instanceRecords = d_instanceRecords; + if (!d_deviceRecord) + BARNEY_CUDA_CALL(Malloc((void**)&d_deviceRecord,sizeof(DeviceRecord))); + BARNEY_CUDA_CALL(Memcpy(d_deviceRecord,&dd,sizeof(DeviceRecord),cudaMemcpyDefault)); + device->sync(); + d_accel = d_deviceRecord; + } + + } +} diff --git a/rtcore/hiprt/Group.h b/rtcore/hiprt/Group.h new file mode 100644 index 00000000..662766ee --- /dev/null +++ b/rtcore/hiprt/Group.h @@ -0,0 +1,126 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/hiprt/Device.h" +#include "rtcore/hiprt/Geom.h" +#include + +namespace rtc { + namespace hiprt { + + struct Device; + + struct Group { + Group(Device *device); + virtual ~Group(); + virtual void buildAccel() = 0; + virtual void refitAccel() { buildAccel(); } + virtual void setTransforms(const std::vector &) {} + AccelHandle getDD() const { return (AccelHandle)d_accel; } + + Device *const device; + void *d_accel = 0; + }; + + /*! a geometry group (BLAS) -- one HIPRT geometry built over the + concatenated primitives of all its geoms. We keep barney's per-prim + (geomID, primID) remap and per-geom SBT exactly as the cuda backend + does, so the device-side shading dispatch (intersect/anyHit/closestHit + function pointers) is identical; HIPRT only supplies the BVH + traversal + that yields the geometry-local primID. */ + struct GeomGroup : public Group { + struct Prim { int geomID; int primID; }; + + /*! the device-visible per-group record stitched into each instance + record; the trace kernel reads sbt+prims to dispatch shading. Field + order matches rtc::cuda::GeomGroup::DeviceRecord's leading fields. */ + struct DeviceRecord { + uint8_t *sbt; + Prim *prims; + uint32_t sbtEntrySize; + uint32_t isTrianglesGroup; + }; + + GeomGroup(Device *device, const std::vector &geoms); + ~GeomGroup(); + + virtual DeviceRecord getRecord() = 0; + hiprtGeometry getGeometry() const { return geom; } + + const std::vector geoms; + + uint8_t *sbt = 0; + size_t sbtEntrySize = 0; + int numPrims = 0; + Prim *prims = 0; + + // HIPRT BLAS for this group, plus the reused build-temp scratch. + hiprtGeometry geom = nullptr; + void *d_buildTmp = nullptr; + size_t buildTmpSize = 0; + // device-side geometry inputs HIPRT references during traversal + void *d_vertices = nullptr; + void *d_indices = nullptr; + void *d_aabbs = nullptr; + }; + + /*! the top-level instance group (TLAS) -- a HIPRT scene over the geom + groups' BLASes. getDD() yields a DeviceRecord with the hiprtScene handle + and the per-instance records the trace kernel uses to map a hit back to + a barney shading dispatch. */ + struct InstanceGroup : public Group { + /*! per-instance record; layout mirrors the leading fields of + rtc::cuda::InstanceGroup::InstanceRecord so the device shading path is + identical. */ + struct InstanceRecord { + affine3f worldToObjectXfm; + affine3f objectToWorldXfm; + GeomGroup::DeviceRecord group; + uint32_t ID; + }; + struct DeviceRecord { + hiprtScene scene; + InstanceRecord *instanceRecords; + }; + + InstanceGroup(Device *device, + const std::vector &groups, + const std::vector &instanceIDs, + const std::vector &xfms); + ~InstanceGroup(); + void buildAccel() override; + void setTransforms(const std::vector &newXfms) override; + + DeviceRecord *d_deviceRecord = 0; + InstanceRecord *d_instanceRecords = 0; + + hiprtScene scene = nullptr; + void *d_sceneTmp = nullptr; + size_t sceneTmpSize = 0; + void *d_instances = nullptr; + void *d_frames = nullptr; + + const std::vector groups; + const std::vector instanceIDs; + std::vector xfms; + }; + + struct TrianglesGeomGroup : public GeomGroup { + TrianglesGeomGroup(Device *device, const std::vector &geoms); + void buildAccel() override; + DeviceRecord getRecord() override; + }; + + struct UserGeomGroup : public GeomGroup { + UserGeomGroup(Device *device, const std::vector &geoms); + void buildAccel() override; + DeviceRecord getRecord() override; + }; + + } +} diff --git a/rtcore/hiprt/ProgramInterface.h b/rtcore/hiprt/ProgramInterface.h new file mode 100644 index 00000000..7267e7b2 --- /dev/null +++ b/rtcore/hiprt/ProgramInterface.h @@ -0,0 +1,255 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/hiprt/ComputeInterface.h" +#include "rtcore/hiprt/Geom.h" +#include "rtcore/hiprt/Group.h" + +#if RTC_DEVICE_CODE +# include +#endif + +namespace rtc { + namespace hiprt { + + using rtc::cuda_common::ComputeInterface; + + struct Device; + + /*! the HIPRT func table handle for `device`. HIPRT calls intersectFunc/ + filterFunc through this table during traversal; barney registers no + per-(geomType,rayType) data (the SBT carries all per-geom state), so the + table is a single 1x1 slot whose callbacks dispatch via the hit's + instance/prim back into barney's intersect/anyHit programs. */ + void *getFuncTable(Device *device); + + struct TraceInterface { + inline __device__ void ignoreIntersection() + { rejectThisHit = true; } + + inline __device__ const void *getLaunchParamsPointer(); + + inline __device__ void reportIntersection(float t, int i) + { current.tMax = t; } + + inline __device__ void *getPRD() const + { return prd; } + + inline __device__ const void *getProgramData() const + { return geomData; } + + inline __device__ const void *getLPData() const + { return lpData; } + + inline __device__ vec3i getLaunchDims() const + { return vec3i(blockDim.x*gridDim.x, blockDim.y*gridDim.y, blockDim.z*gridDim.z); } + + inline __device__ vec3i getLaunchIndex() const + { return vec3i(threadIdx.x+blockIdx.x*blockDim.x, + threadIdx.y+blockIdx.y*blockDim.y, + threadIdx.z+blockIdx.z*blockDim.z); } + + inline __device__ vec2f getTriangleBarycentrics() const + { return current.triangleBarycentrics; } + + inline __device__ int getGeometryIndex() const + { return current.geomID; } + + inline __device__ int getPrimitiveIndex() const + { return current.primID; } + + inline __device__ int getInstanceID() const + { return currentInstance->ID; } + + inline __device__ int getRTCInstanceIndex() const + { return current.instID; } + + inline __device__ float getRayTmax() const + { return current.tMax; } + + inline __device__ float getRayTmin() const + { return tMin; } + + inline __device__ vec3f getObjectRayDirection() const + { return object.dir; } + + inline __device__ vec3f getObjectRayOrigin() const + { return object.org; } + + inline __device__ vec3f getWorldRayDirection() const + { return world.dir; } + + inline __device__ vec3f getWorldRayOrigin() const + { return world.org; } + + inline __device__ vec3f transformNormalFromObjectToWorldSpace(vec3f v) const; + inline __device__ vec3f transformPointFromObjectToWorldSpace(vec3f v) const; + inline __device__ vec3f transformVectorFromObjectToWorldSpace(vec3f v) const; + inline __device__ vec3f transformNormalFromWorldToObjectSpace(vec3f v) const; + inline __device__ vec3f transformPointFromWorldToObjectSpace(vec3f v) const; + inline __device__ vec3f transformVectorFromWorldToObjectSpace(vec3f v) const; + + inline __device__ void traceRay(rtc::AccelHandle world, + vec3f org, vec3f dir, + float t0, float t1, void *prdPtr); + + inline __device__ bool intersectTriangle(const vec3f v0,const vec3f v1, + const vec3f v2, bool dbg=false); +#if RTC_DEVICE_CODE + // HIPRT func-table hooks: setupHit primes per-hit state; the intersect + // thunk runs barney's intersect program (custom geoms); the filter thunk + // runs anyHit and folds the hit into `accepted` (triangle + custom geoms). + inline __device__ Geom::SBTHeader *setupHit(uint32_t instID, uint32_t localPrimID); + inline __device__ void foldAcceptedHit(Geom::SBTHeader *header); + inline __device__ __attribute__((noinline)) + bool hiprtIntersectThunk(::hiprtHit &hit); + inline __device__ __attribute__((noinline)) + void hiprtFilterThunk(const ::hiprtHit &hit); +#endif + + // launch params + const void *lpData; + + // ray/traversal state (mirrors the software backend's TraceInterface): + void *prd; + const void *geomData; + float tMin; + Geom::SBTHeader *acceptedSBT; + struct { + vec2f triangleBarycentrics; + int primID; + int geomID; + int instID; + float tMax; + } current, accepted; + struct { + vec3f org; + vec3f dir; + } world, object; + const InstanceGroup::InstanceRecord *currentInstance; + bool rejectThisHit; + + // HIPRT-specific: the instance-group device record (carries the + // hiprtScene + instance records) and the func-table handle the traversal + // needs. void* so this header stays usable in non-HIPRT host TUs. + rtc::AccelHandle hiprtWorld; + void *hiprtFuncTableHandle; + }; + + } +} + +#define RTC_DECLARE_GLOBALS(ignore) /* ignore */ + +// ------------------------------------------------------------------ +// User (custom) geometry. As in the cuda backend, the per-geom intersect/anyHit/ +// closestHit programs are stored as device function pointers in the SBT (filled +// by rtc_hiprt_writeAddresses) and called from the trace kernel; HIPRT supplies +// the AABB-BVH traversal. The bounds kernel computes per-prim AABBs that the +// hiprt Group hands to HIPRT as a hiprtAABBListPrimitive. +// ------------------------------------------------------------------ +#if RTC_DEVICE_CODE +# define RTC_HIPRT_USERGEOM_KERNELS(name,type,has_ah,has_ch) \ + __device__ void \ + _rtc_hiprt_boundsFunc__##name(const void *geom, \ + owl::common::box3f &result, \ + int primID) \ + { \ + ::rtc::hiprt::TraceInterface ti; \ + type::bounds(ti,geom,result,primID); \ + } \ + __global__ void \ + _rtc_hiprt_boundsFuncKernel__##name(const void *geom, \ + owl::common::box3f *boundsArray, \ + int numPrims) \ + { \ + uint32_t blockIndex \ + = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y; \ + uint32_t primID \ + = threadIdx.x + blockDim.x*threadIdx.y \ + + blockDim.x*blockDim.y*blockIndex; \ + if (primID < numPrims) \ + _rtc_hiprt_boundsFunc__##name(geom,boundsArray[primID],primID); \ + } \ + __global__ \ + void rtc_hiprt_writeAddresses_##name(rtc::hiprt::Geom::SBTHeader *h) \ + { \ + h->ah = has_ah?type::anyHit:0; \ + h->ch = has_ch?type::closestHit:0; \ + h->user.intersect = type::intersect; \ + } +#else +# define RTC_HIPRT_USERGEOM_KERNELS(name,type,has_ah,has_ch) \ + __global__ void \ + _rtc_hiprt_boundsFuncKernel__##name(const void *geom, \ + owl::common::box3f *boundsArray, \ + int numPrims); \ + __global__ \ + void rtc_hiprt_writeAddresses_##name(rtc::hiprt::Geom::SBTHeader *h); +#endif + +#define RTC_EXPORT_USER_GEOM(name,DD,type,has_ah,has_ch) \ + RTC_HIPRT_USERGEOM_KERNELS(name,type,has_ah,has_ch); \ + extern "C" void \ + _rtc_hiprt_writeBounds__##name(rtc::Device *device, \ + const void *geom, \ + owl::common::box3f *boundsArray, \ + int numPrims) \ + { \ + if (numPrims == 0) return; \ + int bs = 1024; \ + int nb = owl::common::divRoundUp(numPrims,bs); \ + _rtc_hiprt_boundsFuncKernel__##name<<stream>>> \ + (geom,boundsArray,numPrims); \ + } \ + rtc::hiprt::GeomType *createGeomType_##name(rtc::Device *device) \ + { \ + ::rtc::hiprt::SetActiveGPU forDuration(device); \ + rtc::hiprt::Geom::SBTHeader *h; \ + BARNEY_CUDA_CALL(Malloc((void **)&h,sizeof(*h))); \ + rtc_hiprt_writeAddresses_##name<<<1,32>>>(h); \ + device->sync(); \ + rtc::hiprt::Geom::SBTHeader hh; \ + BARNEY_CUDA_CALL(Memcpy(&hh,h,sizeof(hh),cudaMemcpyDefault)); \ + BARNEY_CUDA_CALL(Free(h)); \ + return new rtc::hiprt::UserGeomType \ + (device, sizeof(DD), _rtc_hiprt_writeBounds__##name, \ + hh.user.intersect, hh.ah, hh.ch); \ + } + +// ------------------------------------------------------------------ +// Triangle geometry +// ------------------------------------------------------------------ +#if RTC_DEVICE_CODE +# define RTC_HIPRT_TRIANGLES_WRITEADDR(name,Programs,has_ah,has_ch) \ + __global__ \ + void rtc_hiprt_writeAddresses_##name(rtc::hiprt::Geom::SBTHeader *h) \ + { \ + h->ah = has_ah?Programs::anyHit:0; \ + h->ch = has_ch?Programs::closestHit:0; \ + } +#else +# define RTC_HIPRT_TRIANGLES_WRITEADDR(name,Programs,has_ah,has_ch) \ + __global__ \ + void rtc_hiprt_writeAddresses_##name(rtc::hiprt::Geom::SBTHeader *h); +#endif + +#define RTC_EXPORT_TRIANGLES_GEOM(name,DD,Programs,has_ah,has_ch) \ + RTC_HIPRT_TRIANGLES_WRITEADDR(name,Programs,has_ah,has_ch) \ + rtc::hiprt::GeomType *createGeomType_##name(rtc::Device *device) \ + { \ + ::rtc::hiprt::SetActiveGPU forDuration(device); \ + rtc::hiprt::Geom::SBTHeader *h; \ + BARNEY_CUDA_CALL(Malloc((void **)&h,sizeof(*h))); \ + rtc_hiprt_writeAddresses_##name<<<1,32>>>(h); \ + device->sync(); \ + rtc::hiprt::Geom::SBTHeader hh; \ + BARNEY_CUDA_CALL(Memcpy(&hh,h,sizeof(hh),cudaMemcpyDefault)); \ + return new rtc::hiprt::TrianglesGeomType \ + (device, sizeof(DD), hh.ah, hh.ch); \ + } diff --git a/rtcore/hiprt/TraceInterface.h b/rtcore/hiprt/TraceInterface.h new file mode 100644 index 00000000..1ca829dd --- /dev/null +++ b/rtcore/hiprt/TraceInterface.h @@ -0,0 +1,286 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily +// +// Device-side ray-tracing interface for the HIPRT (hardware-RT) backend. The +// shading surface (getPrimitiveIndex/getInstanceID/getWorldRayOrigin/the +// transform helpers/...) is identical to the software (cuda) backend; the only +// difference is traceRay(), which walks HIPRT's BVH (a hiprtScene built by +// rtcore/hiprt/Group) instead of the cuBQL software two-level walk, then runs +// barney's existing per-geometry function-pointer dispatch (intersect/anyHit/ +// closestHit programs) on the returned hit. This keeps barney's megakernel + +// function-pointer SBT model; HIPRT supplies BVH build + traversal + hit +// enumeration only. +// +// HIPRT device traversal is provided by including +// (the full inline traversal implementation), so the trace kernel links against +// libhiprt without going through HIPRT's runtime JIT. The kernel TU must also +// define the func-table callbacks intersectFunc()/filterFunc() (see the +// RTC_HIPRT_TRACEKERNEL macro) that route HIPRT's custom-primitive intersection +// and any-hit filtering back into barney's intersect/anyHit programs. + +#pragma once + +#include "rtcore/cudaCommon/ComputeInterface.h" +#include "rtcore/hiprt/TraceKernel.h" +#include "rtcore/hiprt/Group.h" +#include "rtcore/hiprt/ProgramInterface.h" +#include "rtcore/cudaCommon/Device.h" + +#if RTC_DEVICE_CODE +# include +#endif + +namespace rtc { + namespace hiprt { + + inline __device__ + bool TraceInterface::intersectTriangle(const vec3f v0, + const vec3f v1, + const vec3f v2, + bool dbg) + { + // barney's triangle test, reused verbatim from the software backend so + // the barycentrics/tMax convention matches the rest of barney exactly. + vec3f e1 = v1-v0; + vec3f e2 = v2-v0; + vec3f N = cross(e1,e2); + if (fabsf(dot(object.dir,N)) < 1e-12f) return false; + float t = -dot(object.org-v0,N)/dot(object.dir,N); + if (t <= 0.f || t >= current.tMax) return false; + vec3f P = object.org - v0 + t*object.dir; + float e1u,e2u,Pu, e1v,e2v,Pv; + if (fabsf(N.x) >= max(fabsf(N.y),fabsf(N.z))) { + e1u = e1.y; e2u = e2.y; Pu = P.y; + e1v = e1.z; e2v = e2.z; Pv = P.z; + } else if (fabsf(N.y) > fabsf(N.z)) { + e1u = e1.x; e2u = e2.x; Pu = P.x; + e1v = e1.z; e2v = e2.z; Pv = P.z; + } else { + e1u = e1.x; e2u = e2.x; Pu = P.x; + e1v = e1.y; e2v = e2.y; Pv = P.y; + } + auto det = [](float a, float b, float c, float d) -> float + { return a*d - c*b; }; + if (det(e1u,e1v,e2u,e2v) == 0.f) return false; + float u = det(Pu,e2u,Pv,e2v)/det(e1u,e2u,e1v,e2v); + float v = det(e1u,Pu,e1v,Pv)/det(e1u,e2u,e1v,e2v); + if ((u < 0.f) || (v < 0.f) || ((u+v) >= 1.f)) return false; + current.triangleBarycentrics = vec2f{ u,v }; + current.tMax = t; + return true; + } + +#if RTC_DEVICE_CODE + /*! set up the per-hit instance/object-space + SBT state for the candidate + (instanceID, geometry-local primID) HIPRT just produced. Mirrors the + cuda backend's intersectPrim preamble. Returns the SBT header. */ + inline __device__ + Geom::SBTHeader *TraceInterface::setupHit(uint32_t instID, uint32_t localPrimID) + { + InstanceGroup::DeviceRecord *model + = (InstanceGroup::DeviceRecord *)hiprtWorld; + currentInstance = model->instanceRecords + instID; + current.instID = instID; + object.org = xfmPoint (currentInstance->worldToObjectXfm, world.org); + object.dir = xfmVector(currentInstance->worldToObjectXfm, world.dir); + + const GeomGroup::DeviceRecord *group = ¤tInstance->group; + GeomGroup::Prim prim = group->prims[localPrimID]; + current.primID = prim.primID; + current.geomID = prim.geomID; + current.tMax = accepted.tMax; + + uint8_t *geomSBT = group->sbt + prim.geomID * group->sbtEntrySize; + Geom::SBTHeader *header = (Geom::SBTHeader *)geomSBT; + this->geomData = (header+1); + return header; + } + + /*! fold the hit currently described by `current`/`header` into `accepted`, + running the geom's anyHit program first (which may reject via + ignoreIntersection). Shared by the custom-geom intersect thunk and the + triangle filter thunk so a hit is committed exactly once, with the same + anyHit/reject/closest-wins logic as the software backend's intersectPrim. */ + inline __device__ + void TraceInterface::foldAcceptedHit(Geom::SBTHeader *header) + { + rejectThisHit = false; + if (header->ah) header->ah(*this); + if (!rejectThisHit) { + accepted.tMax = current.tMax; + accepted.primID = current.primID; + accepted.geomID = current.geomID; + accepted.instID = current.instID; + accepted.triangleBarycentrics = current.triangleBarycentrics; + acceptedSBT = header; + } + } + + /*! HIPRT intersectFunc hook for CUSTOM geometry. Runs barney's intersect + program for this candidate prim EXACTLY ONCE and folds the hit here + (matching the software backend's single intersectPrim per prim). For + has_ch=false geoms (cylinders/cones/capsules) the intersect program + itself writes the hit (material BSDF, hitIDs); running it once per + candidate -- not again in the filter -- is what keeps that hit-write from + corrupting memory as the custom-prim count grows. The closest hit wins + because the intersect tests against current.tMax (= accepted.tMax), so a + farther prim early-returns before writing. + + Returns FALSE unconditionally so HIPRT discards this leaf and keeps + walking the BVH (the record-and-ignore enumerate-all pattern); the filter + is therefore never invoked for custom geoms, so the intersect runs only + once. Kept __noinline__ to bound the trace kernel's register footprint. */ + inline __device__ __attribute__((noinline)) + bool TraceInterface::hiprtIntersectThunk(::hiprtHit &hit) + { + Geom::SBTHeader *header = setupHit(hit.instanceID, hit.primID); + header->user.intersect(*this); + if (current.tMax < accepted.tMax) + foldAcceptedHit(header); + return false; + } + + /*! HIPRT filterFunc hook (any-hit) -- reached only for TRIANGLE geometry + (the custom intersect thunk returns false, so HIPRT never calls the + filter for custom geoms). Recompute the barycentrics with barney's own + triangle test so they match the rest of barney, run anyHit, and fold the + hit. Always returns true to the caller so HIPRT keeps enumerating. */ + inline __device__ __attribute__((noinline)) + void TraceInterface::hiprtFilterThunk(const ::hiprtHit &hit) + { + Geom::SBTHeader *header = setupHit(hit.instanceID, hit.primID); + vec3i indices = header->triangles.indices[current.primID]; + vec3f v0 = header->triangles.vertices[indices.x]; + vec3f v1 = header->triangles.vertices[indices.y]; + vec3f v2 = header->triangles.vertices[indices.z]; + if (!intersectTriangle(v0,v1,v2)) return; + foldAcceptedHit(header); + } +#endif + + inline __device__ + void TraceInterface::traceRay(rtc::AccelHandle _world, + vec3f org, + vec3f dir, + float t0, + float t1, + void *prdPtr) + { +#if RTC_DEVICE_CODE + if (fabsf(dir.x) < 1e-6f) dir.x = 1e-6f; + if (fabsf(dir.y) < 1e-6f) dir.y = 1e-6f; + if (fabsf(dir.z) < 1e-6f) dir.z = 1e-6f; + + prd = prdPtr; + world.org = org; + world.dir = dir; + tMin = t0; + accepted.tMax = t1; + accepted.primID = -1; + accepted.instID = -1; + acceptedSBT = 0; + hiprtWorld = _world; + + InstanceGroup::DeviceRecord *model + = (InstanceGroup::DeviceRecord *)_world; + if (!model || !model->scene) return; + + ::hiprtRay ray; + ray.origin = { org.x, org.y, org.z }; + ray.direction = { dir.x, dir.y, dir.z }; + ray.minT = t0; + ray.maxT = t1; + + // The any-hit traversal enumerates hits; for custom geoms HIPRT calls + // intersectFunc -> hiprtIntersectThunk (runs barney's intersect once and + // folds the hit), for triangles filterFunc -> hiprtFilterThunk (runs the + // triangle test + anyHit and folds). Both fold the closest non-rejected + // hit into this TraceInterface (the payload) and tell HIPRT to discard the + // leaf so it keeps walking. This is the record-and-ignore pattern; + // barney's own anyHit programs (transparent/clip geometry) run during + // traversal. The intersect runs exactly once per candidate prim (as in the + // software backend), so has_ch=false geoms write their hit only once. + ::hiprtSceneTraversalAnyHit tr(model->scene, ray, hiprtFullRayMask, + hiprtTraversalHintDefault, + /*payload*/ this, + (hiprtFuncTable)hiprtFuncTableHandle, + /*rayType*/ 0); + hiprtHit hit = tr.getNextHit(); + // The filter has already folded every hit into `accepted`; the returned + // hit is the final one (or invalid). All shading state lives in this + // TraceInterface (the payload), not in a kernel-stack array. + (void)hit; + + if (acceptedSBT && acceptedSBT->ch) { + current = accepted; + this->geomData = (acceptedSBT+1); + currentInstance = model->instanceRecords + accepted.instID; + object.org = xfmPoint (currentInstance->worldToObjectXfm, world.org); + object.dir = xfmVector(currentInstance->worldToObjectXfm, world.dir); + acceptedSBT->ch(*this); + } +#endif + } + + inline __device__ + vec3f TraceInterface::transformNormalFromObjectToWorldSpace(vec3f v) const + { return xfmVector(currentInstance->objectToWorldXfm.l,(const owl::common::vec3f &)v); } + + inline __device__ + vec3f TraceInterface::transformPointFromObjectToWorldSpace(vec3f v) const + { return xfmPoint(currentInstance->objectToWorldXfm,(const owl::common::vec3f &)v); } + + inline __device__ + vec3f TraceInterface::transformVectorFromObjectToWorldSpace(vec3f v) const + { return xfmVector(currentInstance->objectToWorldXfm.l,(const owl::common::vec3f &)v); } + + inline __device__ + vec3f TraceInterface::transformNormalFromWorldToObjectSpace(vec3f v) const + { return xfmVector(currentInstance->worldToObjectXfm.l,(const owl::common::vec3f &)v); } + + inline __device__ + vec3f TraceInterface::transformPointFromWorldToObjectSpace(vec3f v) const + { return xfmPoint(currentInstance->worldToObjectXfm,(const owl::common::vec3f &)v); } + + inline __device__ + vec3f TraceInterface::transformVectorFromWorldToObjectSpace(vec3f v) const + { return xfmVector(currentInstance->worldToObjectXfm.l,(const owl::common::vec3f &)v); } + + } +} + +// The trace kernel: one launched HIP kernel doing raygen + traversal + shading +// (HIPRT's single-kernel model). __launch_bounds__(256) keeps a small declared +// block size to bound per-thread VGPRs through the HIPRT traversal call +// (barney launches 16x16=256). +#if RTC_DEVICE_CODE +# define RTC_HIPRT_TRACEKERNEL(name,Class) \ + __global__ void __launch_bounds__(256) \ + rtc_hiprt_run_##name(::rtc::hiprt::TraceInterface ti) \ + { \ + Class::run(ti); \ + } +#else +# define RTC_HIPRT_TRACEKERNEL(name,Class) \ + __global__ void rtc_hiprt_run_##name(::rtc::hiprt::TraceInterface ti); +#endif + +#define RTC_EXPORT_TRACE2D(name,Class) \ + RTC_HIPRT_TRACEKERNEL(name,Class) \ + void rtc_hiprt_launch_##name(rtc::Device *device, \ + vec2i dims, \ + const void *lpData) \ + { \ + vec2i bs(16,16); \ + vec2i nb = divRoundUp(dims,bs); \ + ::rtc::hiprt::TraceInterface ti; \ + ti.lpData = lpData; \ + ti.hiprtFuncTableHandle = rtc::hiprt::getFuncTable(device); \ + rtc_hiprt_run_##name \ + <<stream>>>(ti); \ + } diff --git a/rtcore/hiprt/TraceKernel.cpp b/rtcore/hiprt/TraceKernel.cpp new file mode 100644 index 00000000..481446f6 --- /dev/null +++ b/rtcore/hiprt/TraceKernel.cpp @@ -0,0 +1,112 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily +// +// HIPRT func-table callbacks and the host-side trace-kernel launcher. +// +// HIPRT's inline device traversal (pulled in via ) +// references two user-provided device functions, intersectFunc() and filterFunc(). +// They are the hook by which barney's per-geometry intersect/anyHit programs run +// during traversal: HIPRT walks the BVH and, for each custom-geom candidate, calls +// intersectFunc -> the intersect thunk (which runs barney's intersect program once +// and folds the hit), and for each triangle candidate calls filterFunc -> the +// filter thunk (triangle test + anyHit + fold). We route both into barney's +// existing function-pointer SBT dispatch via the TraceInterface payload, so HIPRT +// supplies the BVH/traversal while barney supplies the shading -- the "keep the +// megakernel" first cut. +// +// This TU is compiled -fgpu-rdc and device-linked into barney_hiprt_programs, so +// these device symbols resolve for the trace kernel in the same module. + +#include "rtcore/hiprt/Device.h" +#include "rtcore/hiprt/TraceKernel.h" + +#define BARNEY_DEVICE_PROGRAM 1 +#include "rtcore/hiprt/TraceInterface.h" + +#include + +// ------------------------------------------------------------------ +// device-side func-table callbacks (required by hiprt_device_impl.h) +// ------------------------------------------------------------------ +HIPRT_DEVICE bool intersectFunc(uint32_t /*geomType*/, + uint32_t /*rayType*/, + const hiprtFuncTableHeader& /*tableHeader*/, + const hiprtRay& /*ray*/, + void* payload, + hiprtHit& hit) +{ + // custom (user) geometry: run barney's intersect program for this candidate + // prim exactly once and fold the hit. Returns false unconditionally so HIPRT + // discards this leaf and keeps enumerating (the filter is never called for + // custom geoms, so the hit-writing intersect runs only once per prim). + auto *ti = (::rtc::hiprt::TraceInterface *)payload; + return ti->hiprtIntersectThunk(hit); +} + +HIPRT_DEVICE bool filterFunc(uint32_t /*geomType*/, + uint32_t /*rayType*/, + const hiprtFuncTableHeader& /*tableHeader*/, + const hiprtRay& /*ray*/, + void* payload, + const hiprtHit& hit) +{ + // any-hit (triangle geometry only): recompute barycentrics, run the geom's + // anyHit program, and fold this hit into barney's accepted state, then ALWAYS + // reject (return true) so HIPRT enumerates every hit along the ray and we keep + // the closest non-rejected one ourselves. This is the record-and-ignore + // filter pattern. (Custom geoms fold in the intersect thunk and never reach + // here, so their hit-writing intersect runs only once per prim.) + auto *ti = (::rtc::hiprt::TraceInterface *)payload; + ti->hiprtFilterThunk(hit); + return true; +} + +namespace rtc { + namespace hiprt { + + // ------------------------------------------------------------------ + // a process-wide 1x1 HIPRT func table; barney carries all per-geom state in + // the SBT, so HIPRT needs no per-(geomType,rayType) data -- the callbacks + // above reach everything through the TraceInterface payload. Created lazily + // per device and cached. + void *getFuncTable(Device *device) + { + static thread_local hiprtFuncTable table = nullptr; + static thread_local hiprtContext forCtx = nullptr; + if (table && forCtx == device->hiprtCtx) return table; + hiprtCreateFuncTable(device->hiprtCtx, 1, 1, table); + hiprtFuncDataSet fds{}; + hiprtSetFuncTable(device->hiprtCtx, table, 0, 0, fds); + forCtx = device->hiprtCtx; + return table; + } + + // ------------------------------------------------------------------ + TraceKernel2D::TraceKernel2D(Device *device, + size_t sizeOfLP, + TraceLaunchFct traceLaunchFct) + : device(device), traceLaunchFct(traceLaunchFct), sizeOfLP(sizeOfLP) + { + SetActiveGPU forDuration(device); + BARNEY_CUDA_CALL(Malloc((void **)&d_lpData,sizeOfLP)); + } + + TraceKernel2D::~TraceKernel2D() + { + SetActiveGPU forDuration(device); + BARNEY_CUDA_CALL_NOTHROW(Free(d_lpData)); + } + + void TraceKernel2D::launch(vec2i launchDims, const void *kernelData) + { + SetActiveGPU forDuration(device); + BARNEY_CUDA_CALL(MemcpyAsync(d_lpData,kernelData,sizeOfLP, + cudaMemcpyDefault,device->stream)); + traceLaunchFct(device,launchDims,d_lpData); + } + + } +} diff --git a/rtcore/hiprt/TraceKernel.h b/rtcore/hiprt/TraceKernel.h new file mode 100644 index 00000000..c743f0d1 --- /dev/null +++ b/rtcore/hiprt/TraceKernel.h @@ -0,0 +1,43 @@ +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. +// SPDX-License-Identifier: Apache-2.0 +// +// \author Jeff Daily + +#pragma once + +#include "rtcore/cudaCommon/Device.h" + +namespace rtc { + namespace hiprt { + + struct Device; + + typedef void (*TraceLaunchFct)(Device *device, vec2i dims, const void *lpData); + + struct TraceKernel2D { + TraceKernel2D(Device *device, + size_t sizeOfLP, + TraceLaunchFct traceLaunchFct); + ~TraceKernel2D(); + void launch(vec2i launchDims, const void *kernelData); + + Device *const device; + TraceLaunchFct const traceLaunchFct; + size_t const sizeOfLP; + void *d_lpData = 0; + }; + + } +} + +#define RTC_IMPORT_TRACE2D(fileNameBase,name,sizeOfLP) \ + void rtc_hiprt_launch_##name(rtc::Device *device, \ + vec2i dims, \ + const void *lpData); \ + \ + ::rtc::TraceKernel2D *createTrace_##name(rtc::Device *device) \ + { \ + return new ::rtc::hiprt::TraceKernel2D \ + (device,sizeOfLP,rtc_hiprt_launch_##name); \ + } From acbe97aba6b89d452f46dc6ce03b4b5ef65846f4 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 23 Jun 2026 22:18:06 +0000 Subject: [PATCH 2/2] [ROCm] Rely on HIP arch auto-detect instead of pinning gfx90a The gfx90a pin sat after enable_language(HIP), so its if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) guard was always false and the block was dead -- enable_language(HIP) has already detected the host arch (or errored). Removing it makes intent clear and keeps the build honoring -DCMAKE_HIP_ARCHITECTURES, auto-detecting the host GPU, or erroring on a no-GPU host, rather than risking a silently wrong gfx90a default if file order ever changed. This change was authored with the assistance of the Claude AI assistant. --- CMakeLists.txt | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 47dc973a..8f0d9df5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -55,9 +55,8 @@ if ((NOT BARNEY_RTC_EXT STREQUAL "") AND (EXISTS "${BARNEY_RTC_EXT}")) include(${BARNEY_RTC_EXT}/config.cmake) elseif (USE_HIP) enable_language(HIP) - if (NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "") - set(CMAKE_HIP_ARCHITECTURES "gfx90a") - endif() + # enable_language(HIP) above auto-detects the host GPU arch (and errors on a + # no-GPU build host); pass -DCMAKE_HIP_ARCHITECTURES=... to override. message("#barney: building GPU code with HIP for CMAKE_HIP_ARCHITECTURES=${CMAKE_HIP_ARCHITECTURES}") # the cuda backend's GPU code is reused under HIP; OptiX is NVIDIA-only. set(BARNEY_HAVE_HIP ON)